[llvm-branch-commits] [llvm] [NVPTX] Prototype widening scalar fp16 atomics to elementwise operations (PR #233904)
Yonah Goldberg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Tue Oct 20 13:55:32 PDT 2026
https://github.com/YonahGoldberg updated https://github.com/llvm/llvm-project/pull/233904
>From 326cc48bcc5502938311d44022f20fc7ea24c0e4 Mon Sep 17 00:00:00 2001
From: Yonah Goldberg <ygoldberg at nvidia.com>
Date: Tue, 20 Oct 2026 16:38:27 +0000
Subject: [PATCH] [NVPTX] Prototype widening scalar fp16 atomics to elementwise
operations
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 109 ++++
llvm/lib/Target/NVPTX/NVPTXISelLowering.h | 2 +
.../CodeGen/NVPTX/atomicrmw-widen-fp16.ll | 455 +++++++++++++++++
.../AtomicExpand/NVPTX/widen-fp16-atomics.ll | 474 ++++++++++++++++++
4 files changed, 1040 insertions(+)
create mode 100644 llvm/test/CodeGen/NVPTX/atomicrmw-widen-fp16.ll
create mode 100644 llvm/test/Transforms/AtomicExpand/NVPTX/widen-fp16-atomics.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 46818c6282b11..5089eddae18f2 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -122,6 +122,18 @@ static cl::opt<bool> AllowFTZAtomics(
"FTZ behavior does not match the function's denormal mode."),
cl::init(true));
+// Experimental: a floating-point identity does not preserve arbitrary bits in
+// the neighboring lane, in particular NaN encodings. Enabling this option also
+// permits accessing the entire enclosing 32-bit word atomically.
+// PTX half-precision arithmetic returns an unspecified NaN:
+// https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#half-precision-floating-point-instructions
+static cl::opt<bool> WidenFP16Atomics(
+ "nvptx-widen-fp16-atomics", cl::Hidden,
+ cl::desc("Widen scalar fp16/bf16 atomics to packed atomics, assuming the "
+ "neighboring lane is safe to access atomically and contains no "
+ "NaN"),
+ cl::init(false));
+
/// Whereas CUDA's implementation (see libdevice) uses ex2.approx for exp2(), it
/// does NOT use lg2.approx for log2, so this is disabled by default.
static cl::opt<bool> UseApproxLog2F32(
@@ -8239,6 +8251,26 @@ static unsigned getWidestNativeAtomicRMWVectorWidth(
return HasPackedAdd ? 2 : 0;
}
+static bool canWidenFP16AtomicRMW(const AtomicRMWInst *AI,
+ bool IgnoreFTZMismatch,
+ const NVPTXSubtarget &STI) {
+ if (!WidenFP16Atomics || AI->isVolatile())
+ return false;
+
+ Type *Ty = AI->getType();
+ if (auto *VecTy = dyn_cast<FixedVectorType>(Ty)) {
+ if (VecTy->getNumElements() != 1)
+ return false;
+ Ty = VecTy->getElementType();
+ }
+ if (!Ty->isHalfTy() && !Ty->isBFloatTy())
+ return false;
+
+ return getWidestNativeAtomicRMWVectorWidth(
+ AI->getOperation(), Ty, AI->getPointerAddressSpace(),
+ *AI->getFunction(), IgnoreFTZMismatch, STI) >= 2;
+}
+
NVPTXTargetLowering::AtomicExpansionKind
NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
const DataLayout &DL = AI->getDataLayout();
@@ -8249,6 +8281,13 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
AllowFTZAtomics ||
AI->hasMetadata(LLVMContext::MD_atomic_ignore_denormal_mode);
+ if (canWidenFP16AtomicRMW(AI, IgnoreFTZMismatch, STI)) {
+ // Reprocess subtraction as addition before constructing the packed value.
+ if (Op == AtomicRMWInst::FSub)
+ return AtomicExpansionKind::Expand;
+ return AtomicExpansionKind::CustomExpand;
+ }
+
if (AI->isElementwise() &&
cast<FixedVectorType>(AI->getType())->getNumElements() > 1) {
auto *VecTy = cast<FixedVectorType>(AI->getType());
@@ -8318,6 +8357,76 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
STI);
}
+void NVPTXTargetLowering::emitExpandAtomicRMW(AtomicRMWInst *AI) const {
+ assert(shouldExpandAtomicRMWInIR(AI) == AtomicExpansionKind::CustomExpand &&
+ "Expected an eligible scalar fp16 atomic");
+ IRBuilder<> Builder(AI);
+ const DataLayout &DL = AI->getDataLayout();
+ Type *Ty = AI->getType()->getScalarType();
+ auto *VecTy = FixedVectorType::get(Ty, 2);
+ Value *Ptr = AI->getPointerOperand();
+ Type *IntPtrTy = DL.getIntPtrType(Ptr->getType());
+
+ // Select the original halfword in its enclosing aligned 32-bit word.
+ Value *IsHigh = Builder.getFalse();
+ Value *AlignedPtr = Ptr;
+ if (AI->getAlign() < Align(4)) {
+ Value *Addr = Builder.CreatePtrToInt(Ptr, IntPtrTy);
+ Value *Halfword = Builder.CreateAnd(Addr, ConstantInt::get(IntPtrTy, 2));
+ IsHigh = Builder.CreateICmpNE(Halfword, ConstantInt::get(IntPtrTy, 0),
+ "atomic.high");
+ AlignedPtr = Builder.CreateIntrinsic(
+ Intrinsic::ptrmask, {Ptr->getType(), IntPtrTy},
+ {Ptr, ConstantInt::getSigned(IntPtrTy, -4)}, nullptr, "atomic.aligned");
+ }
+
+ // -0 preserves both signs of zero under round-to-nearest addition. Min/max
+ // use +infinity/-infinity. All forms preserve subnormals through .noftz.
+ // These identities require the neighboring lane's no-NaN assumption above.
+ Constant *Identity;
+ switch (AI->getOperation()) {
+ case AtomicRMWInst::FAdd:
+ Identity = ConstantFP::getNegativeZero(Ty);
+ break;
+ case AtomicRMWInst::FMin:
+ case AtomicRMWInst::FMinimumNum:
+ Identity = ConstantFP::getInfinity(Ty);
+ break;
+ case AtomicRMWInst::FMax:
+ case AtomicRMWInst::FMaximumNum:
+ Identity = ConstantFP::getInfinity(Ty, /*Negative=*/true);
+ break;
+ default:
+ llvm_unreachable("Unsupported fp16 atomic widening");
+ }
+ Value *Val = AI->getValOperand();
+ if (AI->getType()->isVectorTy())
+ Val = Builder.CreateExtractElement(Val, Builder.getInt32(0));
+ // Keep vector indices constant: dynamic insertion can spill the pair to
+ // local memory during legalization.
+ Value *Lo = Builder.CreateSelect(IsHigh, Identity, Val);
+ Value *Hi = Builder.CreateSelect(IsHigh, Val, Identity);
+ Value *PackedVal = Builder.CreateInsertElement(PoisonValue::get(VecTy), Lo,
+ Builder.getInt32(0));
+ PackedVal = Builder.CreateInsertElement(PackedVal, Hi, Builder.getInt32(1),
+ "atomic.packed");
+ auto *Wide = Builder.CreateAtomicRMW(AI->getOperation(), AlignedPtr,
+ PackedVal, Align(4), AI->getOrdering(),
+ AI->getSyncScopeID());
+ Wide->setElementwise(true);
+ // Widening may touch a different object, so do not copy alias metadata.
+ Wide->copyMetadata(*AI, {LLVMContext::MD_mem_cache_hint,
+ LLVMContext::MD_atomic_ignore_denormal_mode});
+ Value *OldLo = Builder.CreateExtractElement(Wide, Builder.getInt32(0));
+ Value *OldHi = Builder.CreateExtractElement(Wide, Builder.getInt32(1));
+ Value *Old = Builder.CreateSelect(IsHigh, OldHi, OldLo, "atomic.old");
+ if (AI->getType()->isVectorTy())
+ Old = Builder.CreateInsertElement(PoisonValue::get(AI->getType()), Old,
+ Builder.getInt32(0));
+ AI->replaceAllUsesWith(Old);
+ AI->eraseFromParent();
+}
+
bool NVPTXTargetLowering::shouldInsertFencesForAtomic(
const Instruction *I) const {
// A true result lets AtomicExpand enforce the original ordering with fences
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
index 4942408f6449b..933d0b506564a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.h
@@ -152,6 +152,8 @@ class NVPTXTargetLowering : public TargetLowering {
AtomicExpansionKind
shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const override;
+ void emitExpandAtomicRMW(AtomicRMWInst *AI) const override;
+
bool aggressivelyPreferBuildVectorSources(EVT VecVT) const override {
// There's rarely any point of packing something into a vector type if we
// already have the source data.
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-widen-fp16.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-widen-fp16.ll
new file mode 100644
index 0000000000000..1c62226fea847
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-widen-fp16.ll
@@ -0,0 +1,455 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-widen-fp16-atomics | FileCheck %s --check-prefix=SM90
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx62 -nvptx-widen-fp16-atomics | FileCheck %s --check-prefix=SM60
+; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-widen-fp16-atomics | %ptxas-verify -arch=sm_90 %}
+; RUN: %if ptxas-sm_60 && ptxas-isa-6.2 %{ llc < %s -mtriple=nvptx64 -mcpu=sm_60 -mattr=+ptx62 -nvptx-widen-fp16-atomics | %ptxas-verify -arch=sm_60 %}
+
+; This disabled-by-default experiment assumes the neighboring halfword is safe
+; to access atomically and contains no NaN. Ordinary scalar and one-lane vector
+; operations use the existing elementwise lowering after widening.
+
+define half @fadd_dynamic(ptr addrspace(1) %p, half %v) {
+; SM90-LABEL: fadd_dynamic(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b16 %rs<7>;
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<6>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fadd_dynamic_param_0];
+; SM90-NEXT: ld.param::func.b16 %rs1, [fadd_dynamic_param_1];
+; SM90-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM90-NEXT: and.b64 %rd3, %rd2, 1;
+; SM90-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM90-NEXT: and.b64 %rd4, %rd1, -4;
+; SM90-NEXT: selp.b16 %rs2, 0x8000, %rs1, %p1;
+; SM90-NEXT: selp.b16 %rs3, %rs1, 0x8000, %p1;
+; SM90-NEXT: mov.b32 %r1, {%rs2, %rs3};
+; SM90-NEXT: mov.b64 %rd5, 12345;
+; SM90-NEXT: atom.acquire.sys.global.add.noftz.L2::cache_hint.f16x2 %r2, [%rd4], %r1, %rd5;
+; SM90-NEXT: mov.b32 {%rs4, %rs5}, %r2;
+; SM90-NEXT: selp.b16 %rs6, %rs5, %rs4, %p1;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs6;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fadd_dynamic(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<2>;
+; SM60-NEXT: .reg .b16 %rs<7>;
+; SM60-NEXT: .reg .b32 %r<3>;
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [fadd_dynamic_param_0];
+; SM60-NEXT: ld.param.b16 %rs1, [fadd_dynamic_param_1];
+; SM60-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM60-NEXT: and.b64 %rd3, %rd2, 1;
+; SM60-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM60-NEXT: and.b64 %rd4, %rd1, -4;
+; SM60-NEXT: selp.b16 %rs2, 0x8000, %rs1, %p1;
+; SM60-NEXT: selp.b16 %rs3, %rs1, 0x8000, %p1;
+; SM60-NEXT: mov.b32 %r1, {%rs2, %rs3};
+; SM60-NEXT: atom.sys.global.add.noftz.f16x2 %r2, [%rd4], %r1;
+; SM60-NEXT: mov.b32 {%rs4, %rs5}, %r2;
+; SM60-NEXT: selp.b16 %rs6, %rs5, %rs4, %p1;
+; SM60-NEXT: membar.sys;
+; SM60-NEXT: st.param.b16 [func_retval0], %rs6;
+; SM60-NEXT: ret;
+ %old = atomicrmw fadd ptr addrspace(1) %p, half %v acquire, align 2, !mem.cache_hint !0
+ ret half %old
+}
+
+; -0 must occupy the high lane for low-halfword additions, and vice versa.
+define half @fadd_low(ptr addrspace(1) %p) {
+; SM90-LABEL: fadd_low(
+; SM90: {
+; SM90-NEXT: .reg .b16 %rs<2>;
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fadd_low_param_0];
+; SM90-NEXT: mov.b32 %r1, -2147468288;
+; SM90-NEXT: atom.relaxed.sys.global.add.noftz.f16x2 %r2, [%rd1], %r1;
+; SM90-NEXT: mov.b32 {%rs1, _}, %r2;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fadd_low(
+; SM60: {
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<3>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [fadd_low_param_0];
+; SM60-NEXT: mov.b32 %r1, -2147468288;
+; SM60-NEXT: atom.sys.global.add.noftz.f16x2 %r2, [%rd1], %r1;
+; SM60-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r2; }
+; SM60-NEXT: st.param.b16 [func_retval0], %rs1;
+; SM60-NEXT: ret;
+ %old = atomicrmw fadd ptr addrspace(1) %p, half 1.0 monotonic, align 4
+ ret half %old
+}
+
+define half @fadd_high() {
+; SM90-LABEL: fadd_high(
+; SM90: {
+; SM90-NEXT: .reg .b16 %rs<2>;
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: mov.b32 %r1, 1006665728;
+; SM90-NEXT: mov.b64 %rd1, 4096;
+; SM90-NEXT: atom.relaxed.sys.global.add.noftz.f16x2 %r2, [%rd1], %r1;
+; SM90-NEXT: mov.b32 {_, %rs1}, %r2;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs1;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fadd_high(
+; SM60: {
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<3>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: mov.b32 %r1, 1006665728;
+; SM60-NEXT: mov.b64 %rd1, 4096;
+; SM60-NEXT: atom.sys.global.add.noftz.f16x2 %r2, [%rd1], %r1;
+; SM60-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r2; }
+; SM60-NEXT: st.param.b16 [func_retval0], %rs1;
+; SM60-NEXT: ret;
+ %old = atomicrmw fadd ptr addrspace(1) inttoptr (i64 4098 to ptr addrspace(1)), half 1.0 monotonic, align 2
+ ret half %old
+}
+
+define half @fadd_shared(ptr addrspace(3) %p, half %v) {
+; SM90-LABEL: fadd_shared(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b16 %rs<7>;
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<5>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fadd_shared_param_0];
+; SM90-NEXT: ld.param::func.b16 %rs1, [fadd_shared_param_1];
+; SM90-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM90-NEXT: and.b64 %rd3, %rd2, 1;
+; SM90-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM90-NEXT: and.b64 %rd4, %rd1, -4;
+; SM90-NEXT: selp.b16 %rs2, 0x8000, %rs1, %p1;
+; SM90-NEXT: selp.b16 %rs3, %rs1, 0x8000, %p1;
+; SM90-NEXT: mov.b32 %r1, {%rs2, %rs3};
+; SM90-NEXT: atom.relaxed.sys.shared.add.noftz.f16x2 %r2, [%rd4], %r1;
+; SM90-NEXT: mov.b32 {%rs4, %rs5}, %r2;
+; SM90-NEXT: selp.b16 %rs6, %rs5, %rs4, %p1;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs6;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fadd_shared(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<2>;
+; SM60-NEXT: .reg .b16 %rs<7>;
+; SM60-NEXT: .reg .b32 %r<3>;
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [fadd_shared_param_0];
+; SM60-NEXT: ld.param.b16 %rs1, [fadd_shared_param_1];
+; SM60-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM60-NEXT: and.b64 %rd3, %rd2, 1;
+; SM60-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM60-NEXT: and.b64 %rd4, %rd1, -4;
+; SM60-NEXT: selp.b16 %rs2, 0x8000, %rs1, %p1;
+; SM60-NEXT: selp.b16 %rs3, %rs1, 0x8000, %p1;
+; SM60-NEXT: mov.b32 %r1, {%rs2, %rs3};
+; SM60-NEXT: atom.sys.shared.add.noftz.f16x2 %r2, [%rd4], %r1;
+; SM60-NEXT: mov.b32 {%rs4, %rs5}, %r2;
+; SM60-NEXT: selp.b16 %rs6, %rs5, %rs4, %p1;
+; SM60-NEXT: st.param.b16 [func_retval0], %rs6;
+; SM60-NEXT: ret;
+ %old = atomicrmw fadd ptr addrspace(3) %p, half %v monotonic, align 2
+ ret half %old
+}
+
+; Vector min/max require sm_90 + PTX 8.1; older targets retain CAS loops.
+define half @fmin_global(ptr addrspace(1) %p, half %v) {
+; SM90-LABEL: fmin_global(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b16 %rs<7>;
+; SM90-NEXT: .reg .b64 %rd<5>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fmin_global_param_0];
+; SM90-NEXT: ld.param::func.b16 %rs1, [fmin_global_param_1];
+; SM90-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM90-NEXT: and.b64 %rd3, %rd2, 1;
+; SM90-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM90-NEXT: and.b64 %rd4, %rd1, -4;
+; SM90-NEXT: selp.b16 %rs2, 0x7C00, %rs1, %p1;
+; SM90-NEXT: selp.b16 %rs3, %rs1, 0x7C00, %p1;
+; SM90-NEXT: atom.acq_rel.cta.global.min.noftz.v2.f16 {%rs4, %rs5}, [%rd4], {%rs2, %rs3};
+; SM90-NEXT: selp.b16 %rs6, %rs5, %rs4, %p1;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs6;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fmin_global(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<2>;
+; SM60-NEXT: .reg .b16 %rs<4>;
+; SM60-NEXT: .reg .b32 %r<17>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [fmin_global_param_1];
+; SM60-NEXT: ld.param.b64 %rd2, [fmin_global_param_0];
+; SM60-NEXT: membar.cta;
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r4, %rd2;
+; SM60-NEXT: and.b32 %r5, %r4, 3;
+; SM60-NEXT: shl.b32 %r1, %r5, 3;
+; SM60-NEXT: shl.b32 %r6, 65535, %r1;
+; SM60-NEXT: not.b32 %r2, %r6;
+; SM60-NEXT: ld.volatile.global.b32 %r16, [%rd1];
+; SM60-NEXT: cvt.f32.f16 %r9, %rs1;
+; SM60-NEXT: $L__BB4_1: // %atomicrmw.start
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: shr.u32 %r7, %r16, %r1;
+; SM60-NEXT: cvt.u16.u32 %rs2, %r7;
+; SM60-NEXT: cvt.f32.f16 %r8, %rs2;
+; SM60-NEXT: min.f32 %r10, %r8, %r9;
+; SM60-NEXT: cvt.rn.f16.f32 %rs3, %r10;
+; SM60-NEXT: cvt.u32.u16 %r11, %rs3;
+; SM60-NEXT: shl.b32 %r12, %r11, %r1;
+; SM60-NEXT: and.b32 %r13, %r16, %r2;
+; SM60-NEXT: or.b32 %r14, %r13, %r12;
+; SM60-NEXT: atom.cta.global.cas.b32 %r3, [%rd1], %r16, %r14;
+; SM60-NEXT: setp.ne.b32 %p1, %r3, %r16;
+; SM60-NEXT: mov.b32 %r16, %r3;
+; SM60-NEXT: @%p1 bra $L__BB4_1;
+; SM60-NEXT: // %bb.2: // %atomicrmw.end
+; SM60-NEXT: shr.u32 %r15, %r3, %r1;
+; SM60-NEXT: membar.cta;
+; SM60-NEXT: st.param.b16 [func_retval0], %r15;
+; SM60-NEXT: ret;
+ %old = atomicrmw fmin ptr addrspace(1) %p, half %v syncscope("block") acq_rel, align 2
+ ret half %old
+}
+
+define half @fmaximumnum_global(ptr addrspace(1) %p, half %v) {
+; SM90-LABEL: fmaximumnum_global(
+; SM90: {
+; SM90-NEXT: .reg .b16 %rs<5>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fmaximumnum_global_param_0];
+; SM90-NEXT: ld.param::func.b16 %rs1, [fmaximumnum_global_param_1];
+; SM90-NEXT: mov.b16 %rs2, 0xFC00;
+; SM90-NEXT: atom.relaxed.sys.global.max.noftz.v2.f16 {%rs3, %rs4}, [%rd1], {%rs1, %rs2};
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs3;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fmaximumnum_global(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<2>;
+; SM60-NEXT: .reg .b16 %rs<4>;
+; SM60-NEXT: .reg .b32 %r<9>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [fmaximumnum_global_param_1];
+; SM60-NEXT: ld.param.b64 %rd1, [fmaximumnum_global_param_0];
+; SM60-NEXT: ld.volatile.global.b32 %r8, [%rd1];
+; SM60-NEXT: cvt.f32.f16 %r2, %rs1;
+; SM60-NEXT: $L__BB5_1: // %atomicrmw.start
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: cvt.u16.u32 %rs2, %r8;
+; SM60-NEXT: cvt.f32.f16 %r3, %rs2;
+; SM60-NEXT: max.f32 %r4, %r3, %r2;
+; SM60-NEXT: cvt.rn.f16.f32 %rs3, %r4;
+; SM60-NEXT: cvt.u32.u16 %r5, %rs3;
+; SM60-NEXT: and.b32 %r6, %r8, -65536;
+; SM60-NEXT: or.b32 %r7, %r6, %r5;
+; SM60-NEXT: atom.sys.global.cas.b32 %r1, [%rd1], %r8, %r7;
+; SM60-NEXT: setp.ne.b32 %p1, %r1, %r8;
+; SM60-NEXT: mov.b32 %r8, %r1;
+; SM60-NEXT: @%p1 bra $L__BB5_1;
+; SM60-NEXT: // %bb.2: // %atomicrmw.end
+; SM60-NEXT: st.param.b16 [func_retval0], %r1;
+; SM60-NEXT: ret;
+ %old = atomicrmw fmaximumnum ptr addrspace(1) %p, half %v monotonic, align 4
+ ret half %old
+}
+
+; Packed bf16 add requires sm_90; both scalar and one-lane vector types work.
+define <1 x bfloat> @fadd_bf16_v1(ptr %p, <1 x bfloat> %v) {
+; SM90-LABEL: fadd_bf16_v1(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b16 %rs<7>;
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<5>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fadd_bf16_v1_param_0];
+; SM90-NEXT: ld.param::func.b16 %rs1, [fadd_bf16_v1_param_1];
+; SM90-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM90-NEXT: and.b64 %rd3, %rd2, 1;
+; SM90-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM90-NEXT: and.b64 %rd4, %rd1, -4;
+; SM90-NEXT: selp.b16 %rs2, 0x8000, %rs1, %p1;
+; SM90-NEXT: selp.b16 %rs3, %rs1, 0x8000, %p1;
+; SM90-NEXT: mov.b32 %r1, {%rs2, %rs3};
+; SM90-NEXT: atom.relaxed.sys.add.noftz.bf16x2 %r2, [%rd4], %r1;
+; SM90-NEXT: mov.b32 {%rs4, %rs5}, %r2;
+; SM90-NEXT: selp.b16 %rs6, %rs5, %rs4, %p1;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs6;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fadd_bf16_v1(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<3>;
+; SM60-NEXT: .reg .b16 %rs<2>;
+; SM60-NEXT: .reg .b32 %r<23>;
+; SM60-NEXT: .reg .b64 %rd<3>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b16 %rs1, [fadd_bf16_v1_param_1];
+; SM60-NEXT: ld.param.b64 %rd2, [fadd_bf16_v1_param_0];
+; SM60-NEXT: and.b64 %rd1, %rd2, -4;
+; SM60-NEXT: cvt.u32.u64 %r4, %rd2;
+; SM60-NEXT: and.b32 %r5, %r4, 3;
+; SM60-NEXT: shl.b32 %r1, %r5, 3;
+; SM60-NEXT: shl.b32 %r6, 65535, %r1;
+; SM60-NEXT: not.b32 %r2, %r6;
+; SM60-NEXT: ld.volatile.b32 %r22, [%rd1];
+; SM60-NEXT: cvt.u32.u16 %r9, %rs1;
+; SM60-NEXT: shl.b32 %r10, %r9, 16;
+; SM60-NEXT: $L__BB6_1: // %atomicrmw.start
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: shr.u32 %r7, %r22, %r1;
+; SM60-NEXT: shl.b32 %r8, %r7, 16;
+; SM60-NEXT: add.rn.f32 %r11, %r8, %r10;
+; SM60-NEXT: bfe.u32 %r12, %r11, 16, 1;
+; SM60-NEXT: add.s32 %r13, %r12, %r11;
+; SM60-NEXT: add.s32 %r14, %r13, 32767;
+; SM60-NEXT: setp.nan.f32 %p1, %r11, %r11;
+; SM60-NEXT: or.b32 %r15, %r11, 4194304;
+; SM60-NEXT: selp.b32 %r16, %r15, %r14, %p1;
+; SM60-NEXT: shr.u32 %r17, %r16, 16;
+; SM60-NEXT: shl.b32 %r18, %r17, %r1;
+; SM60-NEXT: and.b32 %r19, %r22, %r2;
+; SM60-NEXT: or.b32 %r20, %r19, %r18;
+; SM60-NEXT: atom.sys.cas.b32 %r3, [%rd1], %r22, %r20;
+; SM60-NEXT: setp.ne.b32 %p2, %r3, %r22;
+; SM60-NEXT: mov.b32 %r22, %r3;
+; SM60-NEXT: @%p2 bra $L__BB6_1;
+; SM60-NEXT: // %bb.2: // %atomicrmw.end
+; SM60-NEXT: shr.u32 %r21, %r3, %r1;
+; SM60-NEXT: st.param.b16 [func_retval0], %r21;
+; SM60-NEXT: ret;
+ %old = atomicrmw elementwise fadd ptr %p, <1 x bfloat> %v monotonic, align 2
+ ret <1 x bfloat> %old
+}
+
+define <1 x half> @fsub_v1(ptr addrspace(1) %p, <1 x half> %v) {
+; SM90-LABEL: fsub_v1(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b16 %rs<8>;
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<5>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fsub_v1_param_0];
+; SM90-NEXT: ld.param::func.b16 %rs1, [fsub_v1_param_1];
+; SM90-NEXT: neg.f16 %rs2, %rs1;
+; SM90-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM90-NEXT: and.b64 %rd3, %rd2, 1;
+; SM90-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM90-NEXT: and.b64 %rd4, %rd1, -4;
+; SM90-NEXT: selp.b16 %rs3, 0x8000, %rs2, %p1;
+; SM90-NEXT: selp.b16 %rs4, %rs2, 0x8000, %p1;
+; SM90-NEXT: mov.b32 %r1, {%rs3, %rs4};
+; SM90-NEXT: atom.acq_rel.gpu.global.add.noftz.f16x2 %r2, [%rd4], %r1;
+; SM90-NEXT: mov.b32 {%rs5, %rs6}, %r2;
+; SM90-NEXT: selp.b16 %rs7, %rs6, %rs5, %p1;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs7;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fsub_v1(
+; SM60: {
+; SM60-NEXT: .reg .pred %p<2>;
+; SM60-NEXT: .reg .b16 %rs<8>;
+; SM60-NEXT: .reg .b32 %r<3>;
+; SM60-NEXT: .reg .b64 %rd<5>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [fsub_v1_param_0];
+; SM60-NEXT: ld.param.b16 %rs1, [fsub_v1_param_1];
+; SM60-NEXT: neg.f16 %rs2, %rs1;
+; SM60-NEXT: membar.gl;
+; SM60-NEXT: shr.u64 %rd2, %rd1, 1;
+; SM60-NEXT: and.b64 %rd3, %rd2, 1;
+; SM60-NEXT: setp.ne.b64 %p1, %rd3, 0;
+; SM60-NEXT: and.b64 %rd4, %rd1, -4;
+; SM60-NEXT: selp.b16 %rs3, 0x8000, %rs2, %p1;
+; SM60-NEXT: selp.b16 %rs4, %rs2, 0x8000, %p1;
+; SM60-NEXT: mov.b32 %r1, {%rs3, %rs4};
+; SM60-NEXT: atom.gpu.global.add.noftz.f16x2 %r2, [%rd4], %r1;
+; SM60-NEXT: mov.b32 {%rs5, %rs6}, %r2;
+; SM60-NEXT: selp.b16 %rs7, %rs6, %rs5, %p1;
+; SM60-NEXT: membar.gl;
+; SM60-NEXT: st.param.b16 [func_retval0], %rs7;
+; SM60-NEXT: ret;
+ %old = atomicrmw elementwise fsub ptr addrspace(1) %p, <1 x half> %v syncscope("device") acq_rel, align 2
+ ret <1 x half> %old
+}
+
+; Scalar seq_cst operations retain the leading fence and acquire atomic.
+define half @fsub_seq_cst(ptr addrspace(1) %p, half %v) {
+; SM90-LABEL: fsub_seq_cst(
+; SM90: {
+; SM90-NEXT: .reg .b16 %rs<4>;
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fsub_seq_cst_param_0];
+; SM90-NEXT: fence.sc.gpu;
+; SM90-NEXT: ld.param::func.b16 %rs1, [fsub_seq_cst_param_1];
+; SM90-NEXT: neg.f16 %rs2, %rs1;
+; SM90-NEXT: mov.b32 %r1, {%rs2, 0x8000};
+; SM90-NEXT: atom.acquire.gpu.global.add.noftz.f16x2 %r2, [%rd1], %r1;
+; SM90-NEXT: mov.b32 {%rs3, _}, %r2;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs3;
+; SM90-NEXT: ret;
+;
+; SM60-LABEL: fsub_seq_cst(
+; SM60: {
+; SM60-NEXT: .reg .b16 %rs<4>;
+; SM60-NEXT: .reg .b32 %r<3>;
+; SM60-NEXT: .reg .b64 %rd<2>;
+; SM60-EMPTY:
+; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b64 %rd1, [fsub_seq_cst_param_0];
+; SM60-NEXT: membar.gl;
+; SM60-NEXT: ld.param.b16 %rs1, [fsub_seq_cst_param_1];
+; SM60-NEXT: neg.f16 %rs2, %rs1;
+; SM60-NEXT: mov.b32 %r1, {%rs2, 0x8000};
+; SM60-NEXT: atom.gpu.global.add.noftz.f16x2 %r2, [%rd1], %r1;
+; SM60-NEXT: { .reg .b16 tmp; mov.b32 {%rs3, tmp}, %r2; }
+; SM60-NEXT: membar.gl;
+; SM60-NEXT: st.param.b16 [func_retval0], %rs3;
+; SM60-NEXT: ret;
+ %old = atomicrmw fsub ptr addrspace(1) %p, half %v syncscope("device") seq_cst, align 4
+ ret half %old
+}
+
+!0 = !{i32 0, !5}
+!5 = !{!"nvvm.l2_cache_hint", i64 12345}
diff --git a/llvm/test/Transforms/AtomicExpand/NVPTX/widen-fp16-atomics.ll b/llvm/test/Transforms/AtomicExpand/NVPTX/widen-fp16-atomics.ll
new file mode 100644
index 0000000000000..89fe543084ca4
--- /dev/null
+++ b/llvm/test/Transforms/AtomicExpand/NVPTX/widen-fp16-atomics.ll
@@ -0,0 +1,474 @@
+; NOTE: Assertions have been autogenerated by utils/update_test_checks.py
+; RUN: opt -S %s -passes='require<libcall-lowering-info>,atomic-expand' -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx87 -nvptx-widen-fp16-atomics -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=CHECK,WIDE
+; RUN: opt -S %s -passes='require<libcall-lowering-info>,atomic-expand' -mtriple=nvptx64-nvidia-cuda -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=CHECK,DEFAULT
+
+target datalayout = "e-p:64:64-p3:32:32"
+target triple = "nvptx64-nvidia-cuda"
+
+; This experiment requires the enclosing 32-bit word to be safe to access
+; atomically and its neighboring lane to contain no NaN. It is off by default.
+; Check dynamic lane selection, identities, ordering, and cache metadata.
+define half @fadd_global(ptr addrspace(1) %p, half %v) {
+; WIDE-LABEL: @fadd_global(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P]], i64 -4)
+; WIDE-NEXT: [[TMP3:%.*]] = select i1 [[ATOMIC_HIGH]], half -0.000000e+00, half [[V:%.*]]
+; WIDE-NEXT: [[TMP4:%.*]] = select i1 [[ATOMIC_HIGH]], half [[V]], half -0.000000e+00
+; WIDE-NEXT: [[TMP5:%.*]] = insertelement <2 x half> poison, half [[TMP3]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP5]], half [[TMP4]], i32 1
+; WIDE-NEXT: [[TMP6:%.*]] = atomicrmw elementwise fadd ptr addrspace(1) [[ATOMIC_ALIGNED]], <2 x half> [[ATOMIC_PACKED]] monotonic, align 4, !mem.cache_hint [[META0:![0-9]+]]
+; WIDE-NEXT: [[TMP7:%.*]] = extractelement <2 x half> [[TMP6]], i32 0
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x half> [[TMP6]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP8]], half [[TMP7]]
+; WIDE-NEXT: ret half [[ATOMIC_OLD]]
+;
+; DEFAULT-LABEL: @fadd_global(
+; DEFAULT-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(1) [[P:%.*]], half [[V:%.*]] monotonic, align 2, !mem.cache_hint [[META0:![0-9]+]]
+; DEFAULT-NEXT: ret half [[OLD]]
+;
+ %old = atomicrmw fadd ptr addrspace(1) %p, half %v monotonic, align 2, !mem.cache_hint !0
+ ret half %old
+}
+
+; A shared pointer is 32 bits, unlike the 64-bit global and generic pointers.
+define half @fadd_shared(ptr addrspace(3) %p, half %v) {
+; WIDE-LABEL: @fadd_shared(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(3) [[P:%.*]] to i32
+; WIDE-NEXT: [[TMP2:%.*]] = and i32 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i32 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(3) @llvm.ptrmask.p3.i32(ptr addrspace(3) [[P]], i32 -4)
+; WIDE-NEXT: [[TMP3:%.*]] = select i1 [[ATOMIC_HIGH]], half -0.000000e+00, half [[V:%.*]]
+; WIDE-NEXT: [[TMP6:%.*]] = select i1 [[ATOMIC_HIGH]], half [[V]], half -0.000000e+00
+; WIDE-NEXT: [[TMP5:%.*]] = insertelement <2 x half> poison, half [[TMP3]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP5]], half [[TMP6]], i32 1
+; WIDE-NEXT: [[TMP4:%.*]] = atomicrmw elementwise fadd ptr addrspace(3) [[ATOMIC_ALIGNED]], <2 x half> [[ATOMIC_PACKED]] acquire, align 4
+; WIDE-NEXT: [[TMP7:%.*]] = extractelement <2 x half> [[TMP4]], i32 0
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x half> [[TMP4]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP8]], half [[TMP7]]
+; WIDE-NEXT: ret half [[ATOMIC_OLD]]
+;
+; DEFAULT-LABEL: @fadd_shared(
+; DEFAULT-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(3) [[P:%.*]], half [[V:%.*]] acquire, align 2
+; DEFAULT-NEXT: ret half [[OLD]]
+;
+ %old = atomicrmw fadd ptr addrspace(3) %p, half %v acquire, align 2
+ ret half %old
+}
+
+; Four-byte alignment makes the target lane statically zero. Widening must
+; discard alias metadata describing only the original halfword.
+define half @fmin_aligned(ptr addrspace(1) %p, half %v) {
+; WIDE-LABEL: @fmin_aligned(
+; WIDE-NEXT: [[TMP4:%.*]] = select i1 false, half +inf, half [[V:%.*]]
+; WIDE-NEXT: [[TMP2:%.*]] = select i1 false, half [[V]], half +inf
+; WIDE-NEXT: [[TMP3:%.*]] = insertelement <2 x half> poison, half [[TMP4]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP3]], half [[TMP2]], i32 1
+; WIDE-NEXT: [[TMP1:%.*]] = atomicrmw elementwise fmin ptr addrspace(1) [[P:%.*]], <2 x half> [[ATOMIC_PACKED]] syncscope("block") acq_rel, align 4
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = extractelement <2 x half> [[TMP1]], i32 0
+; WIDE-NEXT: [[TMP6:%.*]] = extractelement <2 x half> [[TMP1]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD1:%.*]] = select i1 false, half [[TMP6]], half [[ATOMIC_OLD]]
+; WIDE-NEXT: ret half [[ATOMIC_OLD1]]
+;
+; DEFAULT-LABEL: @fmin_aligned(
+; DEFAULT-NEXT: fence syncscope("block") release
+; DEFAULT-NEXT: [[TMP1:%.*]] = load atomic i32, ptr addrspace(1) [[P:%.*]] syncscope("block") monotonic, align 4
+; DEFAULT-NEXT: br label [[ATOMICRMW_START:%.*]]
+; DEFAULT: atomicrmw.start:
+; DEFAULT-NEXT: [[LOADED:%.*]] = phi i32 [ [[TMP1]], [[TMP0:%.*]] ], [ [[NEWLOADED:%.*]], [[ATOMICRMW_START]] ]
+; DEFAULT-NEXT: [[EXTRACTED:%.*]] = trunc i32 [[LOADED]] to i16
+; DEFAULT-NEXT: [[TMP2:%.*]] = bitcast i16 [[EXTRACTED]] to half
+; DEFAULT-NEXT: [[TMP3:%.*]] = call half @llvm.minnum.f16(half [[TMP2]], half [[V:%.*]])
+; DEFAULT-NEXT: [[TMP4:%.*]] = bitcast half [[TMP3]] to i16
+; DEFAULT-NEXT: [[EXTENDED:%.*]] = zext i16 [[TMP4]] to i32
+; DEFAULT-NEXT: [[UNMASKED:%.*]] = and i32 [[LOADED]], -65536
+; DEFAULT-NEXT: [[INSERTED:%.*]] = or i32 [[UNMASKED]], [[EXTENDED]]
+; DEFAULT-NEXT: [[TMP5:%.*]] = cmpxchg ptr addrspace(1) [[P]], i32 [[LOADED]], i32 [[INSERTED]] syncscope("block") monotonic monotonic, align 4, !alias.scope [[META2:![0-9]+]]
+; DEFAULT-NEXT: [[SUCCESS:%.*]] = extractvalue { i32, i1 } [[TMP5]], 1
+; DEFAULT-NEXT: [[NEWLOADED]] = extractvalue { i32, i1 } [[TMP5]], 0
+; DEFAULT-NEXT: br i1 [[SUCCESS]], label [[ATOMICRMW_END:%.*]], label [[ATOMICRMW_START]]
+; DEFAULT: atomicrmw.end:
+; DEFAULT-NEXT: [[EXTRACTED1:%.*]] = trunc i32 [[NEWLOADED]] to i16
+; DEFAULT-NEXT: [[TMP6:%.*]] = bitcast i16 [[EXTRACTED1]] to half
+; DEFAULT-NEXT: fence syncscope("block") acquire
+; DEFAULT-NEXT: ret half [[TMP6]]
+;
+ %old = atomicrmw fmin ptr addrspace(1) %p, half %v syncscope("block") acq_rel, align 4, !alias.scope !1
+ ret half %old
+}
+
+define half @fmax_global(ptr addrspace(1) %p, half %v) {
+; WIDE-LABEL: @fmax_global(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P]], i64 -4)
+; WIDE-NEXT: [[TMP3:%.*]] = select i1 [[ATOMIC_HIGH]], half -inf, half [[V:%.*]]
+; WIDE-NEXT: [[TMP6:%.*]] = select i1 [[ATOMIC_HIGH]], half [[V]], half -inf
+; WIDE-NEXT: [[TMP5:%.*]] = insertelement <2 x half> poison, half [[TMP3]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP5]], half [[TMP6]], i32 1
+; WIDE-NEXT: [[TMP4:%.*]] = atomicrmw elementwise fmax ptr addrspace(1) [[ATOMIC_ALIGNED]], <2 x half> [[ATOMIC_PACKED]] monotonic, align 4
+; WIDE-NEXT: [[TMP7:%.*]] = extractelement <2 x half> [[TMP4]], i32 0
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x half> [[TMP4]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP8]], half [[TMP7]]
+; WIDE-NEXT: ret half [[ATOMIC_OLD]]
+;
+; DEFAULT-LABEL: @fmax_global(
+; DEFAULT-NEXT: [[ALIGNEDADDR:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P:%.*]], i64 -4)
+; DEFAULT-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64
+; DEFAULT-NEXT: [[PTRLSB:%.*]] = and i64 [[TMP1]], 3
+; DEFAULT-NEXT: [[TMP2:%.*]] = shl i64 [[PTRLSB]], 3
+; DEFAULT-NEXT: [[SHIFTAMT:%.*]] = trunc i64 [[TMP2]] to i32
+; DEFAULT-NEXT: [[MASK:%.*]] = shl i32 65535, [[SHIFTAMT]]
+; DEFAULT-NEXT: [[INV_MASK:%.*]] = xor i32 [[MASK]], -1
+; DEFAULT-NEXT: [[TMP3:%.*]] = load atomic i32, ptr addrspace(1) [[ALIGNEDADDR]] monotonic, align 4
+; DEFAULT-NEXT: br label [[ATOMICRMW_START:%.*]]
+; DEFAULT: atomicrmw.start:
+; DEFAULT-NEXT: [[LOADED:%.*]] = phi i32 [ [[TMP3]], [[TMP0:%.*]] ], [ [[NEWLOADED:%.*]], [[ATOMICRMW_START]] ]
+; DEFAULT-NEXT: [[SHIFTED:%.*]] = lshr i32 [[LOADED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[EXTRACTED:%.*]] = trunc i32 [[SHIFTED]] to i16
+; DEFAULT-NEXT: [[TMP4:%.*]] = bitcast i16 [[EXTRACTED]] to half
+; DEFAULT-NEXT: [[TMP5:%.*]] = call half @llvm.maxnum.f16(half [[TMP4]], half [[V:%.*]])
+; DEFAULT-NEXT: [[TMP6:%.*]] = bitcast half [[TMP5]] to i16
+; DEFAULT-NEXT: [[EXTENDED:%.*]] = zext i16 [[TMP6]] to i32
+; DEFAULT-NEXT: [[SHIFTED1:%.*]] = shl nuw i32 [[EXTENDED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[UNMASKED:%.*]] = and i32 [[LOADED]], [[INV_MASK]]
+; DEFAULT-NEXT: [[INSERTED:%.*]] = or i32 [[UNMASKED]], [[SHIFTED1]]
+; DEFAULT-NEXT: [[TMP7:%.*]] = cmpxchg ptr addrspace(1) [[ALIGNEDADDR]], i32 [[LOADED]], i32 [[INSERTED]] monotonic monotonic, align 4
+; DEFAULT-NEXT: [[SUCCESS:%.*]] = extractvalue { i32, i1 } [[TMP7]], 1
+; DEFAULT-NEXT: [[NEWLOADED]] = extractvalue { i32, i1 } [[TMP7]], 0
+; DEFAULT-NEXT: br i1 [[SUCCESS]], label [[ATOMICRMW_END:%.*]], label [[ATOMICRMW_START]]
+; DEFAULT: atomicrmw.end:
+; DEFAULT-NEXT: [[SHIFTED2:%.*]] = lshr i32 [[NEWLOADED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[EXTRACTED3:%.*]] = trunc i32 [[SHIFTED2]] to i16
+; DEFAULT-NEXT: [[TMP8:%.*]] = bitcast i16 [[EXTRACTED3]] to half
+; DEFAULT-NEXT: ret half [[TMP8]]
+;
+ %old = atomicrmw fmax ptr addrspace(1) %p, half %v monotonic, align 2
+ ret half %old
+}
+
+; Subtraction is negated once, then reprocessed as addition.
+define <1 x half> @fsub_v1(ptr addrspace(1) %p, <1 x half> %v) {
+; WIDE-LABEL: @fsub_v1(
+; WIDE-NEXT: [[FNEG:%.*]] = fneg <1 x half> [[V:%.*]]
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P]], i64 -4)
+; WIDE-NEXT: [[TMP4:%.*]] = extractelement <1 x half> [[FNEG]], i32 0
+; WIDE-NEXT: [[TMP7:%.*]] = select i1 [[ATOMIC_HIGH]], half -0.000000e+00, half [[TMP4]]
+; WIDE-NEXT: [[TMP10:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP4]], half -0.000000e+00
+; WIDE-NEXT: [[TMP11:%.*]] = insertelement <2 x half> poison, half [[TMP7]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP11]], half [[TMP10]], i32 1
+; WIDE-NEXT: [[TMP5:%.*]] = atomicrmw elementwise fadd ptr addrspace(1) [[ATOMIC_ALIGNED]], <2 x half> [[ATOMIC_PACKED]] syncscope("device") acq_rel, align 4
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x half> [[TMP5]], i32 0
+; WIDE-NEXT: [[TMP9:%.*]] = extractelement <2 x half> [[TMP5]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP9]], half [[TMP8]]
+; WIDE-NEXT: [[TMP6:%.*]] = insertelement <1 x half> poison, half [[ATOMIC_OLD]], i32 0
+; WIDE-NEXT: ret <1 x half> [[TMP6]]
+;
+; DEFAULT-LABEL: @fsub_v1(
+; DEFAULT-NEXT: [[FNEG:%.*]] = fneg <1 x half> [[V:%.*]]
+; DEFAULT-NEXT: [[OLD:%.*]] = atomicrmw elementwise fadd ptr addrspace(1) [[P:%.*]], <1 x half> [[FNEG]] syncscope("device") acq_rel, align 2
+; DEFAULT-NEXT: ret <1 x half> [[OLD]]
+;
+ %old = atomicrmw elementwise fsub ptr addrspace(1) %p, <1 x half> %v syncscope("device") acq_rel, align 2
+ ret <1 x half> %old
+}
+
+; One-lane whole-value vectors have the same atomicity as a scalar.
+define <1 x half> @fminimumnum_v1(ptr addrspace(1) %p, <1 x half> %v) {
+; WIDE-LABEL: @fminimumnum_v1(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P]], i64 -4)
+; WIDE-NEXT: [[TMP4:%.*]] = extractelement <1 x half> [[V:%.*]], i32 0
+; WIDE-NEXT: [[TMP7:%.*]] = select i1 [[ATOMIC_HIGH]], half +inf, half [[TMP4]]
+; WIDE-NEXT: [[TMP10:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP4]], half +inf
+; WIDE-NEXT: [[TMP11:%.*]] = insertelement <2 x half> poison, half [[TMP7]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP11]], half [[TMP10]], i32 1
+; WIDE-NEXT: [[TMP5:%.*]] = atomicrmw elementwise fminimumnum ptr addrspace(1) [[ATOMIC_ALIGNED]], <2 x half> [[ATOMIC_PACKED]] monotonic, align 4
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x half> [[TMP5]], i32 0
+; WIDE-NEXT: [[TMP9:%.*]] = extractelement <2 x half> [[TMP5]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP9]], half [[TMP8]]
+; WIDE-NEXT: [[TMP6:%.*]] = insertelement <1 x half> poison, half [[ATOMIC_OLD]], i32 0
+; WIDE-NEXT: ret <1 x half> [[TMP6]]
+;
+; DEFAULT-LABEL: @fminimumnum_v1(
+; DEFAULT-NEXT: [[ALIGNEDADDR:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P:%.*]], i64 -4)
+; DEFAULT-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64
+; DEFAULT-NEXT: [[PTRLSB:%.*]] = and i64 [[TMP1]], 3
+; DEFAULT-NEXT: [[TMP2:%.*]] = shl i64 [[PTRLSB]], 3
+; DEFAULT-NEXT: [[SHIFTAMT:%.*]] = trunc i64 [[TMP2]] to i32
+; DEFAULT-NEXT: [[MASK:%.*]] = shl i32 65535, [[SHIFTAMT]]
+; DEFAULT-NEXT: [[INV_MASK:%.*]] = xor i32 [[MASK]], -1
+; DEFAULT-NEXT: [[TMP3:%.*]] = load atomic i32, ptr addrspace(1) [[ALIGNEDADDR]] monotonic, align 4
+; DEFAULT-NEXT: br label [[ATOMICRMW_START:%.*]]
+; DEFAULT: atomicrmw.start:
+; DEFAULT-NEXT: [[LOADED:%.*]] = phi i32 [ [[TMP3]], [[TMP0:%.*]] ], [ [[NEWLOADED:%.*]], [[ATOMICRMW_START]] ]
+; DEFAULT-NEXT: [[SHIFTED:%.*]] = lshr i32 [[LOADED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[EXTRACTED:%.*]] = trunc i32 [[SHIFTED]] to i16
+; DEFAULT-NEXT: [[TMP4:%.*]] = bitcast i16 [[EXTRACTED]] to <1 x half>
+; DEFAULT-NEXT: [[TMP5:%.*]] = call <1 x half> @llvm.minimumnum.v1f16(<1 x half> [[TMP4]], <1 x half> [[V:%.*]])
+; DEFAULT-NEXT: [[TMP6:%.*]] = bitcast <1 x half> [[TMP5]] to i16
+; DEFAULT-NEXT: [[EXTENDED:%.*]] = zext i16 [[TMP6]] to i32
+; DEFAULT-NEXT: [[SHIFTED1:%.*]] = shl nuw i32 [[EXTENDED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[UNMASKED:%.*]] = and i32 [[LOADED]], [[INV_MASK]]
+; DEFAULT-NEXT: [[INSERTED:%.*]] = or i32 [[UNMASKED]], [[SHIFTED1]]
+; DEFAULT-NEXT: [[TMP7:%.*]] = cmpxchg ptr addrspace(1) [[ALIGNEDADDR]], i32 [[LOADED]], i32 [[INSERTED]] monotonic monotonic, align 4
+; DEFAULT-NEXT: [[SUCCESS:%.*]] = extractvalue { i32, i1 } [[TMP7]], 1
+; DEFAULT-NEXT: [[NEWLOADED]] = extractvalue { i32, i1 } [[TMP7]], 0
+; DEFAULT-NEXT: br i1 [[SUCCESS]], label [[ATOMICRMW_END:%.*]], label [[ATOMICRMW_START]]
+; DEFAULT: atomicrmw.end:
+; DEFAULT-NEXT: [[SHIFTED2:%.*]] = lshr i32 [[NEWLOADED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[EXTRACTED3:%.*]] = trunc i32 [[SHIFTED2]] to i16
+; DEFAULT-NEXT: [[TMP8:%.*]] = bitcast i16 [[EXTRACTED3]] to <1 x half>
+; DEFAULT-NEXT: ret <1 x half> [[TMP8]]
+;
+ %old = atomicrmw fminimumnum ptr addrspace(1) %p, <1 x half> %v monotonic, align 2
+ ret <1 x half> %old
+}
+
+define bfloat @fadd_bf16_generic(ptr %p, bfloat %v) {
+; WIDE-LABEL: @fadd_bf16_generic(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr @llvm.ptrmask.p0.i64(ptr [[P]], i64 -4)
+; WIDE-NEXT: [[TMP3:%.*]] = select i1 [[ATOMIC_HIGH]], bfloat -0.000000e+00, bfloat [[V:%.*]]
+; WIDE-NEXT: [[TMP6:%.*]] = select i1 [[ATOMIC_HIGH]], bfloat [[V]], bfloat -0.000000e+00
+; WIDE-NEXT: [[TMP5:%.*]] = insertelement <2 x bfloat> poison, bfloat [[TMP3]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x bfloat> [[TMP5]], bfloat [[TMP6]], i32 1
+; WIDE-NEXT: [[TMP4:%.*]] = atomicrmw elementwise fadd ptr [[ATOMIC_ALIGNED]], <2 x bfloat> [[ATOMIC_PACKED]] monotonic, align 4
+; WIDE-NEXT: [[TMP7:%.*]] = extractelement <2 x bfloat> [[TMP4]], i32 0
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x bfloat> [[TMP4]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], bfloat [[TMP8]], bfloat [[TMP7]]
+; WIDE-NEXT: ret bfloat [[ATOMIC_OLD]]
+;
+; DEFAULT-LABEL: @fadd_bf16_generic(
+; DEFAULT-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr [[P:%.*]], bfloat [[V:%.*]] monotonic, align 2
+; DEFAULT-NEXT: ret bfloat [[OLD]]
+;
+ %old = atomicrmw fadd ptr %p, bfloat %v monotonic, align 2
+ ret bfloat %old
+}
+
+define bfloat @fmaximumnum_bf16(ptr addrspace(1) %p, bfloat %v) {
+; WIDE-LABEL: @fmaximumnum_bf16(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P]], i64 -4)
+; WIDE-NEXT: [[TMP3:%.*]] = select i1 [[ATOMIC_HIGH]], bfloat -inf, bfloat [[V:%.*]]
+; WIDE-NEXT: [[TMP6:%.*]] = select i1 [[ATOMIC_HIGH]], bfloat [[V]], bfloat -inf
+; WIDE-NEXT: [[TMP5:%.*]] = insertelement <2 x bfloat> poison, bfloat [[TMP3]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x bfloat> [[TMP5]], bfloat [[TMP6]], i32 1
+; WIDE-NEXT: [[TMP4:%.*]] = atomicrmw elementwise fmaximumnum ptr addrspace(1) [[ATOMIC_ALIGNED]], <2 x bfloat> [[ATOMIC_PACKED]] monotonic, align 4
+; WIDE-NEXT: [[TMP7:%.*]] = extractelement <2 x bfloat> [[TMP4]], i32 0
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x bfloat> [[TMP4]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], bfloat [[TMP8]], bfloat [[TMP7]]
+; WIDE-NEXT: ret bfloat [[ATOMIC_OLD]]
+;
+; DEFAULT-LABEL: @fmaximumnum_bf16(
+; DEFAULT-NEXT: [[ALIGNEDADDR:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P:%.*]], i64 -4)
+; DEFAULT-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64
+; DEFAULT-NEXT: [[PTRLSB:%.*]] = and i64 [[TMP1]], 3
+; DEFAULT-NEXT: [[TMP2:%.*]] = shl i64 [[PTRLSB]], 3
+; DEFAULT-NEXT: [[SHIFTAMT:%.*]] = trunc i64 [[TMP2]] to i32
+; DEFAULT-NEXT: [[MASK:%.*]] = shl i32 65535, [[SHIFTAMT]]
+; DEFAULT-NEXT: [[INV_MASK:%.*]] = xor i32 [[MASK]], -1
+; DEFAULT-NEXT: [[TMP3:%.*]] = load atomic i32, ptr addrspace(1) [[ALIGNEDADDR]] monotonic, align 4
+; DEFAULT-NEXT: br label [[ATOMICRMW_START:%.*]]
+; DEFAULT: atomicrmw.start:
+; DEFAULT-NEXT: [[LOADED:%.*]] = phi i32 [ [[TMP3]], [[TMP0:%.*]] ], [ [[NEWLOADED:%.*]], [[ATOMICRMW_START]] ]
+; DEFAULT-NEXT: [[SHIFTED:%.*]] = lshr i32 [[LOADED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[EXTRACTED:%.*]] = trunc i32 [[SHIFTED]] to i16
+; DEFAULT-NEXT: [[TMP4:%.*]] = bitcast i16 [[EXTRACTED]] to bfloat
+; DEFAULT-NEXT: [[TMP5:%.*]] = call bfloat @llvm.maximumnum.bf16(bfloat [[TMP4]], bfloat [[V:%.*]])
+; DEFAULT-NEXT: [[TMP6:%.*]] = bitcast bfloat [[TMP5]] to i16
+; DEFAULT-NEXT: [[EXTENDED:%.*]] = zext i16 [[TMP6]] to i32
+; DEFAULT-NEXT: [[SHIFTED1:%.*]] = shl nuw i32 [[EXTENDED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[UNMASKED:%.*]] = and i32 [[LOADED]], [[INV_MASK]]
+; DEFAULT-NEXT: [[INSERTED:%.*]] = or i32 [[UNMASKED]], [[SHIFTED1]]
+; DEFAULT-NEXT: [[TMP7:%.*]] = cmpxchg ptr addrspace(1) [[ALIGNEDADDR]], i32 [[LOADED]], i32 [[INSERTED]] monotonic monotonic, align 4
+; DEFAULT-NEXT: [[SUCCESS:%.*]] = extractvalue { i32, i1 } [[TMP7]], 1
+; DEFAULT-NEXT: [[NEWLOADED]] = extractvalue { i32, i1 } [[TMP7]], 0
+; DEFAULT-NEXT: br i1 [[SUCCESS]], label [[ATOMICRMW_END:%.*]], label [[ATOMICRMW_START]]
+; DEFAULT: atomicrmw.end:
+; DEFAULT-NEXT: [[SHIFTED2:%.*]] = lshr i32 [[NEWLOADED]], [[SHIFTAMT]]
+; DEFAULT-NEXT: [[EXTRACTED3:%.*]] = trunc i32 [[SHIFTED2]] to i16
+; DEFAULT-NEXT: [[TMP8:%.*]] = bitcast i16 [[EXTRACTED3]] to bfloat
+; DEFAULT-NEXT: ret bfloat [[TMP8]]
+;
+ %old = atomicrmw fmaximumnum ptr addrspace(1) %p, bfloat %v monotonic, align 2
+ ret bfloat %old
+}
+
+; Volatile accesses cannot introduce extra memory accesses under this option.
+define half @fadd_volatile(ptr addrspace(1) %p, half %v) {
+; CHECK-LABEL: @fadd_volatile(
+; CHECK-NEXT: [[OLD:%.*]] = atomicrmw volatile fadd ptr addrspace(1) [[P:%.*]], half [[V:%.*]] monotonic, align 2
+; CHECK-NEXT: ret half [[OLD]]
+;
+ %old = atomicrmw volatile fadd ptr addrspace(1) %p, half %v monotonic, align 2
+ ret half %old
+}
+
+; PTX has no shared-memory vector min/max.
+define half @fmin_shared(ptr addrspace(3) %p, half %v) {
+; CHECK-LABEL: @fmin_shared(
+; CHECK-NEXT: [[ALIGNEDADDR:%.*]] = call ptr addrspace(3) @llvm.ptrmask.p3.i32(ptr addrspace(3) [[P:%.*]], i32 -4)
+; CHECK-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(3) [[P]] to i32
+; CHECK-NEXT: [[PTRLSB:%.*]] = and i32 [[TMP1]], 3
+; CHECK-NEXT: [[SHIFTAMT:%.*]] = shl i32 [[PTRLSB]], 3
+; CHECK-NEXT: [[MASK:%.*]] = shl i32 65535, [[SHIFTAMT]]
+; CHECK-NEXT: [[INV_MASK:%.*]] = xor i32 [[MASK]], -1
+; CHECK-NEXT: [[TMP3:%.*]] = load atomic i32, ptr addrspace(3) [[ALIGNEDADDR]] monotonic, align 4
+; CHECK-NEXT: br label [[ATOMICRMW_START:%.*]]
+; CHECK: atomicrmw.start:
+; CHECK-NEXT: [[LOADED:%.*]] = phi i32 [ [[TMP3]], [[TMP0:%.*]] ], [ [[NEWLOADED:%.*]], [[ATOMICRMW_START]] ]
+; CHECK-NEXT: [[SHIFTED:%.*]] = lshr i32 [[LOADED]], [[SHIFTAMT]]
+; CHECK-NEXT: [[EXTRACTED:%.*]] = trunc i32 [[SHIFTED]] to i16
+; CHECK-NEXT: [[TMP4:%.*]] = bitcast i16 [[EXTRACTED]] to half
+; CHECK-NEXT: [[TMP5:%.*]] = call half @llvm.minnum.f16(half [[TMP4]], half [[V:%.*]])
+; CHECK-NEXT: [[TMP6:%.*]] = bitcast half [[TMP5]] to i16
+; CHECK-NEXT: [[EXTENDED:%.*]] = zext i16 [[TMP6]] to i32
+; CHECK-NEXT: [[SHIFTED1:%.*]] = shl nuw i32 [[EXTENDED]], [[SHIFTAMT]]
+; CHECK-NEXT: [[UNMASKED:%.*]] = and i32 [[LOADED]], [[INV_MASK]]
+; CHECK-NEXT: [[INSERTED:%.*]] = or i32 [[UNMASKED]], [[SHIFTED1]]
+; CHECK-NEXT: [[TMP7:%.*]] = cmpxchg ptr addrspace(3) [[ALIGNEDADDR]], i32 [[LOADED]], i32 [[INSERTED]] monotonic monotonic, align 4
+; CHECK-NEXT: [[SUCCESS:%.*]] = extractvalue { i32, i1 } [[TMP7]], 1
+; CHECK-NEXT: [[NEWLOADED]] = extractvalue { i32, i1 } [[TMP7]], 0
+; CHECK-NEXT: br i1 [[SUCCESS]], label [[ATOMICRMW_END:%.*]], label [[ATOMICRMW_START]]
+; CHECK: atomicrmw.end:
+; CHECK-NEXT: [[SHIFTED2:%.*]] = lshr i32 [[NEWLOADED]], [[SHIFTAMT]]
+; CHECK-NEXT: [[EXTRACTED3:%.*]] = trunc i32 [[SHIFTED2]] to i16
+; CHECK-NEXT: [[TMP8:%.*]] = bitcast i16 [[EXTRACTED3]] to half
+; CHECK-NEXT: ret half [[TMP8]]
+;
+ %old = atomicrmw fmin ptr addrspace(3) %p, half %v monotonic, align 2
+ ret half %old
+}
+
+; NaN-propagating fminimum/fmaximum cannot use PTX's number-preferring min/max.
+define half @fminimum_global(ptr addrspace(1) %p, half %v) {
+; CHECK-LABEL: @fminimum_global(
+; CHECK-NEXT: [[ALIGNEDADDR:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P:%.*]], i64 -4)
+; CHECK-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P]] to i64
+; CHECK-NEXT: [[PTRLSB:%.*]] = and i64 [[TMP1]], 3
+; CHECK-NEXT: [[TMP2:%.*]] = shl i64 [[PTRLSB]], 3
+; CHECK-NEXT: [[SHIFTAMT:%.*]] = trunc i64 [[TMP2]] to i32
+; CHECK-NEXT: [[MASK:%.*]] = shl i32 65535, [[SHIFTAMT]]
+; CHECK-NEXT: [[INV_MASK:%.*]] = xor i32 [[MASK]], -1
+; CHECK-NEXT: [[TMP3:%.*]] = load atomic i32, ptr addrspace(1) [[ALIGNEDADDR]] monotonic, align 4
+; CHECK-NEXT: br label [[ATOMICRMW_START:%.*]]
+; CHECK: atomicrmw.start:
+; CHECK-NEXT: [[LOADED:%.*]] = phi i32 [ [[TMP3]], [[TMP0:%.*]] ], [ [[NEWLOADED:%.*]], [[ATOMICRMW_START]] ]
+; CHECK-NEXT: [[SHIFTED:%.*]] = lshr i32 [[LOADED]], [[SHIFTAMT]]
+; CHECK-NEXT: [[EXTRACTED:%.*]] = trunc i32 [[SHIFTED]] to i16
+; CHECK-NEXT: [[TMP4:%.*]] = bitcast i16 [[EXTRACTED]] to half
+; CHECK-NEXT: [[TMP5:%.*]] = call half @llvm.minimum.f16(half [[TMP4]], half [[V:%.*]])
+; CHECK-NEXT: [[TMP6:%.*]] = bitcast half [[TMP5]] to i16
+; CHECK-NEXT: [[EXTENDED:%.*]] = zext i16 [[TMP6]] to i32
+; CHECK-NEXT: [[SHIFTED1:%.*]] = shl nuw i32 [[EXTENDED]], [[SHIFTAMT]]
+; CHECK-NEXT: [[UNMASKED:%.*]] = and i32 [[LOADED]], [[INV_MASK]]
+; CHECK-NEXT: [[INSERTED:%.*]] = or i32 [[UNMASKED]], [[SHIFTED1]]
+; CHECK-NEXT: [[TMP7:%.*]] = cmpxchg ptr addrspace(1) [[ALIGNEDADDR]], i32 [[LOADED]], i32 [[INSERTED]] monotonic monotonic, align 4
+; CHECK-NEXT: [[SUCCESS:%.*]] = extractvalue { i32, i1 } [[TMP7]], 1
+; CHECK-NEXT: [[NEWLOADED]] = extractvalue { i32, i1 } [[TMP7]], 0
+; CHECK-NEXT: br i1 [[SUCCESS]], label [[ATOMICRMW_END:%.*]], label [[ATOMICRMW_START]]
+; CHECK: atomicrmw.end:
+; CHECK-NEXT: [[SHIFTED2:%.*]] = lshr i32 [[NEWLOADED]], [[SHIFTAMT]]
+; CHECK-NEXT: [[EXTRACTED3:%.*]] = trunc i32 [[SHIFTED2]] to i16
+; CHECK-NEXT: [[TMP8:%.*]] = bitcast i16 [[EXTRACTED3]] to half
+; CHECK-NEXT: ret half [[TMP8]]
+;
+ %old = atomicrmw fminimum ptr addrspace(1) %p, half %v monotonic, align 2
+ ret half %old
+}
+
+; Widening still respects the active lane's denormal mode requirements.
+define half @fadd_ftz(ptr addrspace(1) %p, half %v) #0 {
+; WIDE-LABEL: @fadd_ftz(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P]], i64 -4)
+; WIDE-NEXT: [[TMP3:%.*]] = select i1 [[ATOMIC_HIGH]], half -0.000000e+00, half [[V:%.*]]
+; WIDE-NEXT: [[TMP6:%.*]] = select i1 [[ATOMIC_HIGH]], half [[V]], half -0.000000e+00
+; WIDE-NEXT: [[TMP5:%.*]] = insertelement <2 x half> poison, half [[TMP3]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP5]], half [[TMP6]], i32 1
+; WIDE-NEXT: [[TMP4:%.*]] = atomicrmw elementwise fadd ptr addrspace(1) [[ATOMIC_ALIGNED]], <2 x half> [[ATOMIC_PACKED]] monotonic, align 4
+; WIDE-NEXT: [[TMP7:%.*]] = extractelement <2 x half> [[TMP4]], i32 0
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x half> [[TMP4]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP8]], half [[TMP7]]
+; WIDE-NEXT: ret half [[ATOMIC_OLD]]
+;
+; DEFAULT-LABEL: @fadd_ftz(
+; DEFAULT-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(1) [[P:%.*]], half [[V:%.*]] monotonic, align 2
+; DEFAULT-NEXT: ret half [[OLD]]
+;
+ %old = atomicrmw fadd ptr addrspace(1) %p, half %v monotonic, align 2
+ ret half %old
+}
+
+define half @fadd_ftz_ignore(ptr addrspace(1) %p, half %v) #0 {
+; WIDE-LABEL: @fadd_ftz_ignore(
+; WIDE-NEXT: [[TMP1:%.*]] = ptrtoint ptr addrspace(1) [[P:%.*]] to i64
+; WIDE-NEXT: [[TMP2:%.*]] = and i64 [[TMP1]], 2
+; WIDE-NEXT: [[ATOMIC_HIGH:%.*]] = icmp ne i64 [[TMP2]], 0
+; WIDE-NEXT: [[ATOMIC_ALIGNED:%.*]] = call ptr addrspace(1) @llvm.ptrmask.p1.i64(ptr addrspace(1) [[P]], i64 -4)
+; WIDE-NEXT: [[TMP3:%.*]] = select i1 [[ATOMIC_HIGH]], half -0.000000e+00, half [[V:%.*]]
+; WIDE-NEXT: [[TMP4:%.*]] = select i1 [[ATOMIC_HIGH]], half [[V]], half -0.000000e+00
+; WIDE-NEXT: [[TMP5:%.*]] = insertelement <2 x half> poison, half [[TMP3]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP5]], half [[TMP4]], i32 1
+; WIDE-NEXT: [[TMP6:%.*]] = atomicrmw elementwise fadd ptr addrspace(1) [[ATOMIC_ALIGNED]], <2 x half> [[ATOMIC_PACKED]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2:![0-9]+]]
+; WIDE-NEXT: [[TMP7:%.*]] = extractelement <2 x half> [[TMP6]], i32 0
+; WIDE-NEXT: [[TMP8:%.*]] = extractelement <2 x half> [[TMP6]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = select i1 [[ATOMIC_HIGH]], half [[TMP8]], half [[TMP7]]
+; WIDE-NEXT: ret half [[ATOMIC_OLD]]
+;
+; DEFAULT-LABEL: @fadd_ftz_ignore(
+; DEFAULT-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(1) [[P:%.*]], half [[V:%.*]] monotonic, align 2, !atomic.ignore.denormal.mode [[META5:![0-9]+]]
+; DEFAULT-NEXT: ret half [[OLD]]
+;
+ %old = atomicrmw fadd ptr addrspace(1) %p, half %v monotonic, align 2, !atomic.ignore.denormal.mode !4
+ ret half %old
+}
+
+attributes #0 = { "denormal-fp-math"="preserve-sign,preserve-sign" }
+
+; Scalar seq_cst operations retain the leading fence and acquire atomic.
+define half @fsub_seq_cst(ptr addrspace(1) %p, half %v) {
+; WIDE-LABEL: @fsub_seq_cst(
+; WIDE-NEXT: fence syncscope("device") seq_cst
+; WIDE-NEXT: [[FNEG:%.*]] = fneg half [[V:%.*]]
+; WIDE-NEXT: [[TMP4:%.*]] = select i1 false, half -0.000000e+00, half [[FNEG]]
+; WIDE-NEXT: [[TMP2:%.*]] = select i1 false, half [[FNEG]], half -0.000000e+00
+; WIDE-NEXT: [[TMP3:%.*]] = insertelement <2 x half> poison, half [[TMP4]], i32 0
+; WIDE-NEXT: [[ATOMIC_PACKED:%.*]] = insertelement <2 x half> [[TMP3]], half [[TMP2]], i32 1
+; WIDE-NEXT: [[TMP1:%.*]] = atomicrmw elementwise fadd ptr addrspace(1) [[P:%.*]], <2 x half> [[ATOMIC_PACKED]] syncscope("device") acquire, align 4
+; WIDE-NEXT: [[ATOMIC_OLD:%.*]] = extractelement <2 x half> [[TMP1]], i32 0
+; WIDE-NEXT: [[TMP6:%.*]] = extractelement <2 x half> [[TMP1]], i32 1
+; WIDE-NEXT: [[ATOMIC_OLD1:%.*]] = select i1 false, half [[TMP6]], half [[ATOMIC_OLD]]
+; WIDE-NEXT: ret half [[ATOMIC_OLD1]]
+;
+; DEFAULT-LABEL: @fsub_seq_cst(
+; DEFAULT-NEXT: fence syncscope("device") seq_cst
+; DEFAULT-NEXT: [[FNEG:%.*]] = fneg half [[V:%.*]]
+; DEFAULT-NEXT: [[OLD:%.*]] = atomicrmw fadd ptr addrspace(1) [[P:%.*]], half [[FNEG]] syncscope("device") acquire, align 4
+; DEFAULT-NEXT: ret half [[OLD]]
+;
+ %old = atomicrmw fsub ptr addrspace(1) %p, half %v syncscope("device") seq_cst, align 4
+ ret half %old
+}
+
+!0 = !{i32 0, !5}
+!5 = !{!"nvvm.l2_cache_hint", i64 12345}
+!1 = !{!2}
+!2 = distinct !{!2, !3, !"halfword scope"}
+!3 = distinct !{!3, i1 false, !"halfword domain"}
+!4 = !{}
More information about the llvm-branch-commits
mailing list