[llvm] [NVPTX] Derive f16 atomic FTZ behavior from the f16 denormal mode (PR #217584)
via llvm-commits
llvm-commits at lists.llvm.org
Thu Aug 20 04:47:28 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Christian Sigg (chsigg)
<details>
<summary>Changes</summary>
shouldExpandAtomicRMWInIR computed the function's FTZ mode once from
APFloat::IEEEsingle() and then consulted it for both the f32 and the f16
atomicrmw fadd cases. The f16 case therefore keyed off the *f32* denormal
mode, so a function built with e.g. -denormal-fp-math-f32=preserve-sign
but IEEE f16 denormals was treated as flushing f16 denormals and expanded
atom.add.f16 into a CAS loop, even though atom.add.f16 never flushes and
so agrees with the function's actual f16 behavior.
Compute the FTZ mode separately for each type from the matching semantics.
The f32 path is unchanged.
---
Full diff: https://github.com/llvm/llvm-project/pull/217584.diff
3 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+12-6)
- (modified) llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll (+11-73)
- (modified) llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll (+11-73)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index e788b0e44041f..928fe94ed3cbd 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -7482,12 +7482,12 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
// 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;
+ const Function *F = AI->getFunction();
// AllowFTZAtomics forces atom.add regardless of the FTZ mismatch.
if (Ty->isFloatTy()) {
+ const bool FTZ = F->getDenormalMode(APFloat::IEEEsingle()).Output ==
+ DenormalMode::PreserveSign;
bool UseNative = AllowFTZAtomics;
switch (AI->getPointerAddressSpace()) {
case llvm::ADDRESS_SPACE_GLOBAL:
@@ -7502,9 +7502,15 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const {
return AtomicExpansionKind::None;
}
- if (Ty->isHalfTy() && (!FTZ || AllowFTZAtomics) &&
- STI.hasFeature(NVPTX::SM70) && STI.hasFeature(NVPTX::PTX63))
- return AtomicExpansionKind::None;
+ if (Ty->isHalfTy()) {
+ // atom.add.f16 never flushes denormals, so it only agrees with a
+ // function that is not in FTZ mode for f16.
+ const bool FTZ = F->getDenormalMode(APFloat::IEEEhalf()).Output ==
+ DenormalMode::PreserveSign;
+ if ((!FTZ || AllowFTZAtomics) && STI.hasFeature(NVPTX::SM70) &&
+ STI.hasFeature(NVPTX::PTX63))
+ return AtomicExpansionKind::None;
+ }
if (Ty->isBFloatTy() && STI.hasFeature(NVPTX::SM90))
return AtomicExpansionKind::None;
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
index 19291c6d0c055..5b9f1a68bceab 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm70.ll
@@ -2495,79 +2495,17 @@ define double @fmaximum_acq_rel_double_global_cta(ptr addrspace(1) %addr, double
}
define half @fadd_acq_rel_half_global_cta(ptr addrspace(1) %addr, half %val) {
-; SM70-NOFTZ-DISALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM70-NOFTZ-DISALLOW: {
-; SM70-NOFTZ-DISALLOW-NEXT: .reg .b16 %rs<3>;
-; SM70-NOFTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
-; SM70-NOFTZ-DISALLOW-EMPTY:
-; SM70-NOFTZ-DISALLOW-NEXT: // %bb.0:
-; SM70-NOFTZ-DISALLOW-NEXT: ld.param.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
-; SM70-NOFTZ-DISALLOW-NEXT: ld.param.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM70-NOFTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
-; SM70-NOFTZ-DISALLOW-NEXT: st.param.b16 [func_retval0], %rs2;
-; SM70-NOFTZ-DISALLOW-NEXT: ret;
-;
-; SM70-NOFTZ-ALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM70-NOFTZ-ALLOW: {
-; SM70-NOFTZ-ALLOW-NEXT: .reg .b16 %rs<3>;
-; SM70-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
-; SM70-NOFTZ-ALLOW-EMPTY:
-; SM70-NOFTZ-ALLOW-NEXT: // %bb.0:
-; SM70-NOFTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
-; SM70-NOFTZ-ALLOW-NEXT: ld.param.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM70-NOFTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
-; SM70-NOFTZ-ALLOW-NEXT: st.param.b16 [func_retval0], %rs2;
-; SM70-NOFTZ-ALLOW-NEXT: ret;
-;
-; SM70-FTZ-DISALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM70-FTZ-DISALLOW: {
-; SM70-FTZ-DISALLOW-NEXT: .reg .pred %p<2>;
-; SM70-FTZ-DISALLOW-NEXT: .reg .b16 %rs<4>;
-; SM70-FTZ-DISALLOW-NEXT: .reg .b32 %r<15>;
-; SM70-FTZ-DISALLOW-NEXT: .reg .b64 %rd<3>;
-; SM70-FTZ-DISALLOW-EMPTY:
-; SM70-FTZ-DISALLOW-NEXT: // %bb.0:
-; SM70-FTZ-DISALLOW-NEXT: ld.param.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM70-FTZ-DISALLOW-NEXT: ld.param.b64 %rd2, [fadd_acq_rel_half_global_cta_param_0];
-; SM70-FTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-DISALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM70-FTZ-DISALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM70-FTZ-DISALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM70-FTZ-DISALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM70-FTZ-DISALLOW-NEXT: mov.b32 %r6, 65535;
-; SM70-FTZ-DISALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM70-FTZ-DISALLOW-NEXT: not.b32 %r2, %r7;
-; SM70-FTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM70-FTZ-DISALLOW-NEXT: $L__BB72_1: // %atomicrmw.start
-; SM70-FTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM70-FTZ-DISALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM70-FTZ-DISALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM70-FTZ-DISALLOW-NEXT: add.rn.ftz.f16 %rs3, %rs2, %rs1;
-; SM70-FTZ-DISALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM70-FTZ-DISALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM70-FTZ-DISALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM70-FTZ-DISALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM70-FTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM70-FTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM70-FTZ-DISALLOW-NEXT: mov.b32 %r14, %r3;
-; SM70-FTZ-DISALLOW-NEXT: @%p1 bra $L__BB72_1;
-; SM70-FTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM70-FTZ-DISALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM70-FTZ-DISALLOW-NEXT: fence.acq_rel.cta;
-; SM70-FTZ-DISALLOW-NEXT: st.param.b16 [func_retval0], %r13;
-; SM70-FTZ-DISALLOW-NEXT: ret;
-;
-; SM70-FTZ-ALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM70-FTZ-ALLOW: {
-; SM70-FTZ-ALLOW-NEXT: .reg .b16 %rs<3>;
-; SM70-FTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
-; SM70-FTZ-ALLOW-EMPTY:
-; SM70-FTZ-ALLOW-NEXT: // %bb.0:
-; SM70-FTZ-ALLOW-NEXT: ld.param.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
-; SM70-FTZ-ALLOW-NEXT: ld.param.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM70-FTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
-; SM70-FTZ-ALLOW-NEXT: st.param.b16 [func_retval0], %rs2;
-; SM70-FTZ-ALLOW-NEXT: ret;
+; SM70-LABEL: fadd_acq_rel_half_global_cta(
+; SM70: {
+; SM70-NEXT: .reg .b16 %rs<3>;
+; SM70-NEXT: .reg .b64 %rd<2>;
+; SM70-EMPTY:
+; SM70-NEXT: // %bb.0:
+; SM70-NEXT: ld.param.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
+; SM70-NEXT: ld.param.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
+; SM70-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
+; SM70-NEXT: st.param.b16 [func_retval0], %rs2;
+; SM70-NEXT: ret;
%retval = atomicrmw fadd ptr addrspace(1) %addr, half %val syncscope("block") acq_rel
ret half %retval
}
diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
index 66071ffa8f45d..a1470453fa427 100644
--- a/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
+++ b/llvm/test/CodeGen/NVPTX/atomicrmw-sm90.ll
@@ -2447,79 +2447,17 @@ define double @fmaximum_acq_rel_double_global_cta(ptr addrspace(1) %addr, double
}
define half @fadd_acq_rel_half_global_cta(ptr addrspace(1) %addr, half %val) {
-; SM90-NOFTZ-DISALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM90-NOFTZ-DISALLOW: {
-; SM90-NOFTZ-DISALLOW-NEXT: .reg .b16 %rs<3>;
-; SM90-NOFTZ-DISALLOW-NEXT: .reg .b64 %rd<2>;
-; SM90-NOFTZ-DISALLOW-EMPTY:
-; SM90-NOFTZ-DISALLOW-NEXT: // %bb.0:
-; SM90-NOFTZ-DISALLOW-NEXT: ld.param::func.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
-; SM90-NOFTZ-DISALLOW-NEXT: ld.param::func.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM90-NOFTZ-DISALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
-; SM90-NOFTZ-DISALLOW-NEXT: st.param::func.b16 [func_retval0], %rs2;
-; SM90-NOFTZ-DISALLOW-NEXT: ret;
-;
-; SM90-NOFTZ-ALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM90-NOFTZ-ALLOW: {
-; SM90-NOFTZ-ALLOW-NEXT: .reg .b16 %rs<3>;
-; SM90-NOFTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
-; SM90-NOFTZ-ALLOW-EMPTY:
-; SM90-NOFTZ-ALLOW-NEXT: // %bb.0:
-; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
-; SM90-NOFTZ-ALLOW-NEXT: ld.param::func.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM90-NOFTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
-; SM90-NOFTZ-ALLOW-NEXT: st.param::func.b16 [func_retval0], %rs2;
-; SM90-NOFTZ-ALLOW-NEXT: ret;
-;
-; SM90-FTZ-DISALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM90-FTZ-DISALLOW: {
-; SM90-FTZ-DISALLOW-NEXT: .reg .pred %p<2>;
-; SM90-FTZ-DISALLOW-NEXT: .reg .b16 %rs<4>;
-; SM90-FTZ-DISALLOW-NEXT: .reg .b32 %r<15>;
-; SM90-FTZ-DISALLOW-NEXT: .reg .b64 %rd<3>;
-; SM90-FTZ-DISALLOW-EMPTY:
-; SM90-FTZ-DISALLOW-NEXT: // %bb.0:
-; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM90-FTZ-DISALLOW-NEXT: ld.param::func.b64 %rd2, [fadd_acq_rel_half_global_cta_param_0];
-; SM90-FTZ-DISALLOW-NEXT: fence.release.cta;
-; SM90-FTZ-DISALLOW-NEXT: and.b64 %rd1, %rd2, -4;
-; SM90-FTZ-DISALLOW-NEXT: cvt.u32.u64 %r4, %rd2;
-; SM90-FTZ-DISALLOW-NEXT: and.b32 %r5, %r4, 3;
-; SM90-FTZ-DISALLOW-NEXT: shl.b32 %r1, %r5, 3;
-; SM90-FTZ-DISALLOW-NEXT: mov.b32 %r6, 65535;
-; SM90-FTZ-DISALLOW-NEXT: shl.b32 %r7, %r6, %r1;
-; SM90-FTZ-DISALLOW-NEXT: not.b32 %r2, %r7;
-; SM90-FTZ-DISALLOW-NEXT: ld.relaxed.cta.global.b32 %r14, [%rd1];
-; SM90-FTZ-DISALLOW-NEXT: $L__BB72_1: // %atomicrmw.start
-; SM90-FTZ-DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; SM90-FTZ-DISALLOW-NEXT: shr.u32 %r8, %r14, %r1;
-; SM90-FTZ-DISALLOW-NEXT: cvt.u16.u32 %rs2, %r8;
-; SM90-FTZ-DISALLOW-NEXT: add.rn.ftz.f16 %rs3, %rs2, %rs1;
-; SM90-FTZ-DISALLOW-NEXT: cvt.u32.u16 %r9, %rs3;
-; SM90-FTZ-DISALLOW-NEXT: shl.b32 %r10, %r9, %r1;
-; SM90-FTZ-DISALLOW-NEXT: and.b32 %r11, %r14, %r2;
-; SM90-FTZ-DISALLOW-NEXT: or.b32 %r12, %r11, %r10;
-; SM90-FTZ-DISALLOW-NEXT: atom.relaxed.cta.global.cas.b32 %r3, [%rd1], %r14, %r12;
-; SM90-FTZ-DISALLOW-NEXT: setp.ne.b32 %p1, %r3, %r14;
-; SM90-FTZ-DISALLOW-NEXT: mov.b32 %r14, %r3;
-; SM90-FTZ-DISALLOW-NEXT: @%p1 bra $L__BB72_1;
-; SM90-FTZ-DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; SM90-FTZ-DISALLOW-NEXT: shr.u32 %r13, %r3, %r1;
-; SM90-FTZ-DISALLOW-NEXT: fence.acquire.cta;
-; SM90-FTZ-DISALLOW-NEXT: st.param::func.b16 [func_retval0], %r13;
-; SM90-FTZ-DISALLOW-NEXT: ret;
-;
-; SM90-FTZ-ALLOW-LABEL: fadd_acq_rel_half_global_cta(
-; SM90-FTZ-ALLOW: {
-; SM90-FTZ-ALLOW-NEXT: .reg .b16 %rs<3>;
-; SM90-FTZ-ALLOW-NEXT: .reg .b64 %rd<2>;
-; SM90-FTZ-ALLOW-EMPTY:
-; SM90-FTZ-ALLOW-NEXT: // %bb.0:
-; SM90-FTZ-ALLOW-NEXT: ld.param::func.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
-; SM90-FTZ-ALLOW-NEXT: ld.param::func.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
-; SM90-FTZ-ALLOW-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
-; SM90-FTZ-ALLOW-NEXT: st.param::func.b16 [func_retval0], %rs2;
-; SM90-FTZ-ALLOW-NEXT: ret;
+; SM90-LABEL: fadd_acq_rel_half_global_cta(
+; SM90: {
+; SM90-NEXT: .reg .b16 %rs<3>;
+; SM90-NEXT: .reg .b64 %rd<2>;
+; SM90-EMPTY:
+; SM90-NEXT: // %bb.0:
+; SM90-NEXT: ld.param::func.b64 %rd1, [fadd_acq_rel_half_global_cta_param_0];
+; SM90-NEXT: ld.param::func.b16 %rs1, [fadd_acq_rel_half_global_cta_param_1];
+; SM90-NEXT: atom.acq_rel.cta.global.add.noftz.f16 %rs2, [%rd1], %rs1;
+; SM90-NEXT: st.param::func.b16 [func_retval0], %rs2;
+; SM90-NEXT: ret;
%retval = atomicrmw fadd ptr addrspace(1) %addr, half %val syncscope("block") acq_rel
ret half %retval
}
``````````
</details>
https://github.com/llvm/llvm-project/pull/217584
More information about the llvm-commits
mailing list