[llvm] [NVPTX] Respect FTZ flag when lowering atomicrmw fadd. (PR #200732)
Justin Lebar via llvm-commits
llvm-commits at lists.llvm.org
Mon Jun 1 00:07:03 PDT 2026
https://github.com/jlebar updated https://github.com/llvm/llvm-project/pull/200732
>From fea9097bff760419d1b3eaa72870f40820945b5a Mon Sep 17 00:00:00 2001
From: Justin Lebar <justin.lebar at gmail.com>
Date: Sun, 31 May 2026 18:07:20 -0700
Subject: [PATCH] [NVPTX] Respect FTZ flag when lowering atomicrmw fadd.
Previously we unconditionally lowered LLVM atomicrmw fadd to PTX
atom.add. This is incorrect, because it ignores the FTZ behavior of the
LLVM and PTX instructions.
Co-Authored-By: Claude Opus 4.8 (1M context) <noreply at anthropic.com>
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 57 +++--
llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll | 19 +-
llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll | 19 +-
llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll | 247 +++++++++++++++++++-
llvm/test/CodeGen/NVPTX/atomics.ll | 34 ++-
5 files changed, 343 insertions(+), 33 deletions(-)
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index d1d01089aa49e..f45491f66cca8 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -7467,22 +7467,53 @@ NVPTXTargetLowering::AtomicExpansionKind
NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
Type *Ty = AI->getValOperand()->getType();
- if (AI->isFloatingPointOperation()) {
- if (AI->getOperation() == AtomicRMWInst::BinOp::FAdd) {
- if (Ty->isHalfTy() && STI.getSmVersion() >= 70 &&
- STI.getPTXVersion() >= 63)
- return AtomicExpansionKind::None;
- if (Ty->isBFloatTy() && STI.getSmVersion() >= 90 &&
- STI.getPTXVersion() >= 78)
- return AtomicExpansionKind::None;
- if (Ty->isFloatTy())
- return AtomicExpansionKind::None;
- if (Ty->isDoubleTy() && STI.hasAtomAddF64())
- return AtomicExpansionKind::None;
+ // Try to lower LLVM atomicrmw fadd to PTX atomic.add. This is complicated
+ // by the weird FTZ behavior PTX atom.add has:
+ // - atom.add.f32 on global memory flushes denormals
+ // - atom.add.f32 on shared memory does not flush denormals
+ // - atom.add.f16 and atomic.add.bf16 never flush denormals
+ //
+ // We lower to atom.add only if the function's FTZ behavior matches that of
+ // atom.add; otherwise, we lower to a CAS loop. But we always allow
+ // atomic.add.bf16; even though it never flushes denormals, we never flush
+ // bf16 denormals when doing regular arithmetic, even when FTZ is enabled.
+ if (AI->isFloatingPointOperation() &&
+ AI->getOperation() == AtomicRMWInst::BinOp::FAdd) {
+ const bool FTZ =
+ AI->getFunction()->getDenormalMode(APFloat::IEEEsingle()).Output ==
+ DenormalMode::PreserveSign;
+
+ if (Ty->isFloatTy()) {
+ switch (AI->getPointerAddressSpace()) {
+ case llvm::ADDRESS_SPACE_GLOBAL:
+ if (FTZ)
+ return AtomicExpansionKind::None;
+ break;
+ case llvm::ADDRESS_SPACE_SHARED:
+ case llvm::ADDRESS_SPACE_SHARED_CLUSTER:
+ if (!FTZ)
+ return AtomicExpansionKind::None;
+ break;
+ }
}
- return AtomicExpansionKind::CmpXChg;
+
+ if (Ty->isHalfTy() && !FTZ && STI.getSmVersion() >= 70 &&
+ STI.getPTXVersion() >= 63)
+ return AtomicExpansionKind::None;
+
+ if (Ty->isBFloatTy() && STI.getSmVersion() >= 90 &&
+ STI.getPTXVersion() >= 78)
+ return AtomicExpansionKind::None;
+
+ if (Ty->isDoubleTy() && STI.hasAtomAddF64())
+ return AtomicExpansionKind::None;
}
+ // All other atomic fp ops expand to a CAS loop.
+ // TODO: Implement other atomic ops!
+ if (AI->isFloatingPointOperation())
+ return AtomicExpansionKind::CmpXChg;
+
assert(Ty->isIntegerTy() && "Ty should be integer at this point");
const unsigned BitWidth = cast<IntegerType>(Ty)->getBitWidth();
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
index a7b2358a4327f..8d198d619e40b 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm60.ll
@@ -1739,14 +1739,25 @@ define i64 @usub_sat_acq_rel_i64_global_cta(ptr addrspace(1) %addr, i64 %val) {
define float @fadd_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val) {
; SM60-LABEL: fadd_acq_rel_float_global_cta(
; SM60: {
-; SM60-NEXT: .reg .b32 %r<3>;
+; SM60-NEXT: .reg .pred %p<2>;
+; SM60-NEXT: .reg .b32 %r<5>;
; SM60-NEXT: .reg .b64 %rd<2>;
; SM60-EMPTY:
; SM60-NEXT: // %bb.0:
+; SM60-NEXT: ld.param.b32 %r2, [fadd_acq_rel_float_global_cta_param_1];
; SM60-NEXT: ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
-; SM60-NEXT: ld.param.b32 %r1, [fadd_acq_rel_float_global_cta_param_1];
-; SM60-NEXT: atom.cta.global.add.f32 %r2, [%rd1], %r1;
-; SM60-NEXT: st.param.b32 [func_retval0], %r2;
+; SM60-NEXT: membar.cta;
+; SM60-NEXT: ld.volatile.global.b32 %r4, [%rd1];
+; SM60-NEXT: $L__BB60_1: // %atomicrmw.start
+; SM60-NEXT: // =>This Inner Loop Header: Depth=1
+; SM60-NEXT: add.rn.f32 %r3, %r4, %r2;
+; SM60-NEXT: atom.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM60-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; SM60-NEXT: mov.b32 %r4, %r1;
+; SM60-NEXT: @%p1 bra $L__BB60_1;
+; SM60-NEXT: // %bb.2: // %atomicrmw.end
+; SM60-NEXT: membar.cta;
+; SM60-NEXT: st.param.b32 [func_retval0], %r1;
; SM60-NEXT: ret;
%retval = atomicrmw fadd ptr addrspace(1) %addr, float %val syncscope("block") acq_rel
ret float %retval
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
index 18165ba5f7a89..71957b262598a 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
@@ -1739,14 +1739,25 @@ define i64 @usub_sat_acq_rel_i64_global_cta(ptr addrspace(1) %addr, i64 %val) {
define float @fadd_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val) {
; SM70-LABEL: fadd_acq_rel_float_global_cta(
; SM70: {
-; SM70-NEXT: .reg .b32 %r<3>;
+; SM70-NEXT: .reg .pred %p<2>;
+; SM70-NEXT: .reg .b32 %r<5>;
; SM70-NEXT: .reg .b64 %rd<2>;
; SM70-EMPTY:
; SM70-NEXT: // %bb.0:
+; SM70-NEXT: ld.param.b32 %r2, [fadd_acq_rel_float_global_cta_param_1];
; SM70-NEXT: ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
-; SM70-NEXT: ld.param.b32 %r1, [fadd_acq_rel_float_global_cta_param_1];
-; SM70-NEXT: atom.acq_rel.cta.global.add.f32 %r2, [%rd1], %r1;
-; SM70-NEXT: st.param.b32 [func_retval0], %r2;
+; SM70-NEXT: fence.acq_rel.cta;
+; SM70-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
+; SM70-NEXT: $L__BB60_1: // %atomicrmw.start
+; SM70-NEXT: // =>This Inner Loop Header: Depth=1
+; SM70-NEXT: add.rn.f32 %r3, %r4, %r2;
+; SM70-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM70-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; SM70-NEXT: mov.b32 %r4, %r1;
+; SM70-NEXT: @%p1 bra $L__BB60_1;
+; SM70-NEXT: // %bb.2: // %atomicrmw.end
+; SM70-NEXT: fence.acq_rel.cta;
+; SM70-NEXT: st.param.b32 [func_retval0], %r1;
; SM70-NEXT: ret;
%retval = atomicrmw fadd ptr addrspace(1) %addr, float %val syncscope("block") acq_rel
ret float %retval
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
index 52dbb2831e27c..bbb05920bcb4e 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
@@ -1739,14 +1739,25 @@ define i64 @usub_sat_acq_rel_i64_global_cta(ptr addrspace(1) %addr, i64 %val) {
define float @fadd_acq_rel_float_global_cta(ptr addrspace(1) %addr, float %val) {
; SM90-LABEL: fadd_acq_rel_float_global_cta(
; SM90: {
-; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b32 %r<5>;
; SM90-NEXT: .reg .b64 %rd<2>;
; SM90-EMPTY:
; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b32 %r2, [fadd_acq_rel_float_global_cta_param_1];
; SM90-NEXT: ld.param.b64 %rd1, [fadd_acq_rel_float_global_cta_param_0];
-; SM90-NEXT: ld.param.b32 %r1, [fadd_acq_rel_float_global_cta_param_1];
-; SM90-NEXT: atom.acq_rel.cta.global.add.f32 %r2, [%rd1], %r1;
-; SM90-NEXT: st.param.b32 [func_retval0], %r2;
+; SM90-NEXT: fence.release.cta;
+; SM90-NEXT: ld.relaxed.cta.global.b32 %r4, [%rd1];
+; SM90-NEXT: $L__BB60_1: // %atomicrmw.start
+; SM90-NEXT: // =>This Inner Loop Header: Depth=1
+; SM90-NEXT: add.rn.f32 %r3, %r4, %r2;
+; SM90-NEXT: atom.relaxed.cta.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM90-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; SM90-NEXT: mov.b32 %r4, %r1;
+; SM90-NEXT: @%p1 bra $L__BB60_1;
+; SM90-NEXT: // %bb.2: // %atomicrmw.end
+; SM90-NEXT: fence.acquire.cta;
+; SM90-NEXT: st.param.b32 [func_retval0], %r1;
; SM90-NEXT: ret;
%retval = atomicrmw fadd ptr addrspace(1) %addr, float %val syncscope("block") acq_rel
ret float %retval
@@ -2981,3 +2992,231 @@ define i32 @nand_seq_cst_i32_global_cta(ptr addrspace(1) %addr, i32 %val) {
ret i32 %retval
}
+
+; Regression test for bug 033. atom.add.f32's subnormal flushing is
+; address-space dependent: global memory flushes (FTZ), shared memory
+; preserves subnormals (IEEE), and generic memory is indeterminate. We may
+; lower atomicrmw fadd float to native atom.add.f32 only when its flush
+; behavior provably matches the function's f32 denormal output mode;
+; otherwise we expand to a CAS loop using the non-flushing add.rn.f32.
+
+; global + FTZ (preserve-sign): native atom.add.f32 flushes, which matches.
+define float @fadd_f32_global_ftz_native(ptr addrspace(1) %p, float %v) #0 {
+; SM90-LABEL: fadd_f32_global_ftz_native(
+; SM90: {
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ftz_native_param_0];
+; SM90-NEXT: ld.param.b32 %r1, [fadd_f32_global_ftz_native_param_1];
+; SM90-NEXT: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1;
+; SM90-NEXT: st.param.b32 [func_retval0], %r2;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr addrspace(1) %p, float %v monotonic
+ ret float %r
+}
+
+; global + ieee: native flushes but the function wants IEEE, so expand.
+define float @fadd_f32_global_ieee_expand(ptr addrspace(1) %p, float %v) {
+; SM90-LABEL: fadd_f32_global_ieee_expand(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b32 %r<5>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b32 %r2, [fadd_f32_global_ieee_expand_param_1];
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ieee_expand_param_0];
+; SM90-NEXT: ld.relaxed.sys.global.b32 %r4, [%rd1];
+; SM90-NEXT: $L__BB101_1: // %atomicrmw.start
+; SM90-NEXT: // =>This Inner Loop Header: Depth=1
+; SM90-NEXT: add.rn.f32 %r3, %r4, %r2;
+; SM90-NEXT: atom.relaxed.sys.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM90-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; SM90-NEXT: mov.b32 %r4, %r1;
+; SM90-NEXT: @%p1 bra $L__BB101_1;
+; SM90-NEXT: // %bb.2: // %atomicrmw.end
+; SM90-NEXT: st.param.b32 [func_retval0], %r1;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr addrspace(1) %p, float %v monotonic
+ ret float %r
+}
+
+; shared + ieee: native atom.add.f32 preserves subnormals, which matches.
+define float @fadd_f32_shared_ieee_native(ptr addrspace(3) %p, float %v) {
+; SM90-LABEL: fadd_f32_shared_ieee_native(
+; SM90: {
+; SM90-NEXT: .reg .b32 %r<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_f32_shared_ieee_native_param_0];
+; SM90-NEXT: ld.param.b32 %r1, [fadd_f32_shared_ieee_native_param_1];
+; SM90-NEXT: atom.relaxed.sys.shared.add.f32 %r2, [%rd1], %r1;
+; SM90-NEXT: st.param.b32 [func_retval0], %r2;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr addrspace(3) %p, float %v monotonic
+ ret float %r
+}
+
+; shared + FTZ (preserve-sign): native preserves but the function wants
+; flushing, so expand.
+define float @fadd_f32_shared_ftz_expand(ptr addrspace(3) %p, float %v) #0 {
+; SM90-LABEL: fadd_f32_shared_ftz_expand(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b32 %r<5>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b32 %r2, [fadd_f32_shared_ftz_expand_param_1];
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_f32_shared_ftz_expand_param_0];
+; SM90-NEXT: ld.relaxed.sys.shared.b32 %r4, [%rd1];
+; SM90-NEXT: $L__BB103_1: // %atomicrmw.start
+; SM90-NEXT: // =>This Inner Loop Header: Depth=1
+; SM90-NEXT: add.rn.ftz.f32 %r3, %r4, %r2;
+; SM90-NEXT: atom.relaxed.sys.shared.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM90-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; SM90-NEXT: mov.b32 %r4, %r1;
+; SM90-NEXT: @%p1 bra $L__BB103_1;
+; SM90-NEXT: // %bb.2: // %atomicrmw.end
+; SM90-NEXT: st.param.b32 [func_retval0], %r1;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr addrspace(3) %p, float %v monotonic
+ ret float %r
+}
+
+; generic + ieee: flush behavior is indeterminate, so expand.
+define float @fadd_f32_generic_ieee_expand(ptr %p, float %v) {
+; SM90-LABEL: fadd_f32_generic_ieee_expand(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b32 %r<5>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b32 %r2, [fadd_f32_generic_ieee_expand_param_1];
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_f32_generic_ieee_expand_param_0];
+; SM90-NEXT: ld.relaxed.sys.b32 %r4, [%rd1];
+; SM90-NEXT: $L__BB104_1: // %atomicrmw.start
+; SM90-NEXT: // =>This Inner Loop Header: Depth=1
+; SM90-NEXT: add.rn.f32 %r3, %r4, %r2;
+; SM90-NEXT: atom.relaxed.sys.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM90-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; SM90-NEXT: mov.b32 %r4, %r1;
+; SM90-NEXT: @%p1 bra $L__BB104_1;
+; SM90-NEXT: // %bb.2: // %atomicrmw.end
+; SM90-NEXT: st.param.b32 [func_retval0], %r1;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr %p, float %v monotonic
+ ret float %r
+}
+
+; generic + FTZ (preserve-sign): flush behavior is indeterminate, so expand.
+define float @fadd_f32_generic_ftz_expand(ptr %p, float %v) #0 {
+; SM90-LABEL: fadd_f32_generic_ftz_expand(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b32 %r<5>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b32 %r2, [fadd_f32_generic_ftz_expand_param_1];
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_f32_generic_ftz_expand_param_0];
+; SM90-NEXT: ld.relaxed.sys.b32 %r4, [%rd1];
+; SM90-NEXT: $L__BB105_1: // %atomicrmw.start
+; SM90-NEXT: // =>This Inner Loop Header: Depth=1
+; SM90-NEXT: add.rn.ftz.f32 %r3, %r4, %r2;
+; SM90-NEXT: atom.relaxed.sys.cas.b32 %r1, [%rd1], %r4, %r3;
+; SM90-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; SM90-NEXT: mov.b32 %r4, %r1;
+; SM90-NEXT: @%p1 bra $L__BB105_1;
+; SM90-NEXT: // %bb.2: // %atomicrmw.end
+; SM90-NEXT: st.param.b32 [func_retval0], %r1;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr %p, float %v monotonic
+ ret float %r
+}
+
+; atom.add.f16 requires .noftz (always preserves subnormals). f16 flushing
+; follows the f32 denormal mode (useF32FTZ), so under FTZ the native
+; instruction would diverge from the function's other f16 ops; expand instead.
+
+; f16 + FTZ: expand to a CAS loop whose add.rn.ftz.f16 flushes.
+define half @fadd_f16_ftz_expand(ptr addrspace(1) %p, half %v) #0 {
+; SM90-LABEL: fadd_f16_ftz_expand(
+; SM90: {
+; SM90-NEXT: .reg .pred %p<2>;
+; SM90-NEXT: .reg .b16 %rs<4>;
+; SM90-NEXT: .reg .b32 %r<15>;
+; SM90-NEXT: .reg .b64 %rd<3>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b16 %rs1, [fadd_f16_ftz_expand_param_1];
+; SM90-NEXT: ld.param.b64 %rd2, [fadd_f16_ftz_expand_param_0];
+; SM90-NEXT: and.b64 %rd1, %rd2, -4;
+; SM90-NEXT: cvt.u32.u64 %r4, %rd2;
+; SM90-NEXT: and.b32 %r5, %r4, 3;
+; SM90-NEXT: shl.b32 %r1, %r5, 3;
+; SM90-NEXT: mov.b32 %r6, 65535;
+; SM90-NEXT: shl.b32 %r7, %r6, %r1;
+; SM90-NEXT: not.b32 %r2, %r7;
+; SM90-NEXT: ld.relaxed.sys.global.b32 %r14, [%rd1];
+; SM90-NEXT: $L__BB106_1: // %atomicrmw.start
+; SM90-NEXT: // =>This Inner Loop Header: Depth=1
+; SM90-NEXT: shr.u32 %r8, %r14, %r1;
+; SM90-NEXT: cvt.u16.u32 %rs2, %r8;
+; SM90-NEXT: add.rn.ftz.f16 %rs3, %rs2, %rs1;
+; SM90-NEXT: cvt.u32.u16 %r9, %rs3;
+; SM90-NEXT: shl.b32 %r10, %r9, %r1;
+; SM90-NEXT: and.b32 %r11, %r14, %r2;
+; SM90-NEXT: or.b32 %r12, %r11, %r10;
+; SM90-NEXT: atom.relaxed.sys.global.cas.b32 %r3, [%rd1], %r14, %r12;
+; SM90-NEXT: setp.ne.b32 %p1, %r3, %r14;
+; SM90-NEXT: mov.b32 %r14, %r3;
+; SM90-NEXT: @%p1 bra $L__BB106_1;
+; SM90-NEXT: // %bb.2: // %atomicrmw.end
+; SM90-NEXT: shr.u32 %r13, %r3, %r1;
+; SM90-NEXT: st.param.b16 [func_retval0], %r13;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr addrspace(1) %p, half %v monotonic
+ ret half %r
+}
+
+; f16 + ieee: native atom.add.noftz.f16 preserves subnormals, which matches.
+define half @fadd_f16_ieee_native(ptr addrspace(1) %p, half %v) {
+; SM90-LABEL: fadd_f16_ieee_native(
+; SM90: {
+; SM90-NEXT: .reg .b16 %rs<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_f16_ieee_native_param_0];
+; SM90-NEXT: ld.param.b16 %rs1, [fadd_f16_ieee_native_param_1];
+; SM90-NEXT: atom.relaxed.sys.global.add.noftz.f16 %rs2, [%rd1], %rs1;
+; SM90-NEXT: st.param.b16 [func_retval0], %rs2;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr addrspace(1) %p, half %v monotonic
+ ret half %r
+}
+
+; bf16 + FTZ: NVPTX has no flushing bf16 add, so native atom.add.noftz.bf16
+; always matches and stays selected even in FTZ mode.
+define bfloat @fadd_bf16_ftz_native(ptr addrspace(1) %p, bfloat %v) #0 {
+; SM90-LABEL: fadd_bf16_ftz_native(
+; SM90: {
+; SM90-NEXT: .reg .b16 %rs<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param.b64 %rd1, [fadd_bf16_ftz_native_param_0];
+; SM90-NEXT: ld.param.b16 %rs1, [fadd_bf16_ftz_native_param_1];
+; SM90-NEXT: atom.relaxed.sys.global.add.noftz.bf16 %rs2, [%rd1], %rs1;
+; SM90-NEXT: st.param.b16 [func_retval0], %rs2;
+; SM90-NEXT: ret;
+ %r = atomicrmw fadd ptr addrspace(1) %p, bfloat %v monotonic
+ ret bfloat %r
+}
+
+attributes #0 = { denormal_fpenv(float: preservesign) }
diff --git a/llvm/test/CodeGen/NVPTX/atomics.ll b/llvm/test/CodeGen/NVPTX/atomics.ll
index 3820b13f7da1f..ab05bf1b7ea3f 100644
--- a/llvm/test/CodeGen/NVPTX/atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/atomics.ll
@@ -8,14 +8,23 @@ declare float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)
define float @atomic_add_f32_generic(ptr %addr, float %val) {
; CHECK-LABEL: atomic_add_f32_generic(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r2, [atomic_add_f32_generic_param_1];
; CHECK-NEXT: ld.param.b64 %rd1, [atomic_add_f32_generic_param_0];
-; CHECK-NEXT: ld.param.b32 %r1, [atomic_add_f32_generic_param_1];
-; CHECK-NEXT: atom.add.f32 %r2, [%rd1], %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ld.volatile.b32 %r4, [%rd1];
+; CHECK-NEXT: $L__BB0_1: // %atomicrmw.start
+; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
+; CHECK-NEXT: add.rn.f32 %r3, %r4, %r2;
+; CHECK-NEXT: atom.cas.b32 %r1, [%rd1], %r4, %r3;
+; CHECK-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; CHECK-NEXT: mov.b32 %r4, %r1;
+; CHECK-NEXT: @%p1 bra $L__BB0_1;
+; CHECK-NEXT: // %bb.2: // %atomicrmw.end
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%ret = call float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)
ret float %ret
@@ -27,14 +36,23 @@ declare float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %v
define float @atomic_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {
; CHECK-LABEL: atomic_add_f32_addrspace1(
; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b32 %r<5>;
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.b32 %r2, [atomic_add_f32_addrspace1_param_1];
; CHECK-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace1_param_0];
-; CHECK-NEXT: ld.param.b32 %r1, [atomic_add_f32_addrspace1_param_1];
-; CHECK-NEXT: atom.global.add.f32 %r2, [%rd1], %r1;
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
+; CHECK-NEXT: ld.volatile.global.b32 %r4, [%rd1];
+; CHECK-NEXT: $L__BB1_1: // %atomicrmw.start
+; CHECK-NEXT: // =>This Inner Loop Header: Depth=1
+; CHECK-NEXT: add.rn.f32 %r3, %r4, %r2;
+; CHECK-NEXT: atom.global.cas.b32 %r1, [%rd1], %r4, %r3;
+; CHECK-NEXT: setp.ne.b32 %p1, %r1, %r4;
+; CHECK-NEXT: mov.b32 %r4, %r1;
+; CHECK-NEXT: @%p1 bra $L__BB1_1;
+; CHECK-NEXT: // %bb.2: // %atomicrmw.end
+; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%ret = call float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %val)
ret float %ret
More information about the llvm-commits
mailing list