https://github.com/chsigg created https://github.com/llvm/llvm-project/pull/217586
PTX atom.add has a fixed denormal behavior that the program cannot control: atom.add.f32 flushes denormals on global memory but not on shared, and atom.add.f16 never flushes. When that disagrees with the function's denormal mode, the backend expands the atomic into a CAS loop so the denormal behavior is preserved. !atomic.ignore.denormal.mode says the denormal behavior of this particular atomic does not matter, so use the native instruction even when it disagrees. This is the same thing -nvptx-allow-ftz-atomics does, except per-instruction instead of per-compilation, which lets a frontend opt in only the operations it knows about -- notably CUDA's atomicAdd(), which is defined in terms of atom.add. Note that -nvptx-allow-ftz-atomics defaults to true, so the new behavior is only observable with -nvptx-allow-ftz-atomics=false. >From 2364e440efaecbf0f330dd7b472ebae0ba8128a7 Mon Sep 17 00:00:00 2001 From: Christian Sigg <[email protected]> Date: Wed, 19 Aug 2026 14:50:24 +0200 Subject: [PATCH] [NVPTX] Honor !atomic.ignore.denormal.mode on atomicrmw fadd PTX atom.add has a fixed denormal behavior that the program cannot control: atom.add.f32 flushes denormals on global memory but not on shared, and atom.add.f16 never flushes. When that disagrees with the function's denormal mode, the backend expands the atomic into a CAS loop so the denormal behavior is preserved. !atomic.ignore.denormal.mode says the denormal behavior of this particular atomic does not matter, so use the native instruction even when it disagrees. This is the same thing -nvptx-allow-ftz-atomics does, except per-instruction instead of per-compilation, which lets a frontend opt in only the operations it knows about -- notably CUDA's atomicAdd(), which is defined in terms of atom.add. Note that -nvptx-allow-ftz-atomics defaults to true, so the new behavior is only observable with -nvptx-allow-ftz-atomics=false. --- llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 14 +- .../NVPTX/atomicrmw-ignore-denormal-mode.ll | 258 ++++++++++++++++++ 2 files changed, 269 insertions(+), 3 deletions(-) create mode 100644 llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp index 928fe94ed3cbd..7f1b7e28f013f 100644 --- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp +++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp @@ -52,6 +52,7 @@ #include "llvm/IR/Instruction.h" #include "llvm/IR/Instructions.h" #include "llvm/IR/IntrinsicsNVPTX.h" +#include "llvm/IR/LLVMContext.h" #include "llvm/IR/Module.h" #include "llvm/IR/Type.h" #include "llvm/IR/Value.h" @@ -7484,11 +7485,18 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const { AI->getOperation() == AtomicRMWInst::BinOp::FAdd) { const Function *F = AI->getFunction(); - // AllowFTZAtomics forces atom.add regardless of the FTZ mismatch. + // Both the -nvptx-allow-ftz-atomics option and per-instruction + // !atomic.ignore.denormal.mode say that denormal handling is insignificant + // here, so atom.add may be used even when its FTZ behavior disagrees with + // the function's. + const bool IgnoreFTZMismatch = + AllowFTZAtomics || + AI->hasMetadata(LLVMContext::MD_atomic_ignore_denormal_mode); + if (Ty->isFloatTy()) { const bool FTZ = F->getDenormalMode(APFloat::IEEEsingle()).Output == DenormalMode::PreserveSign; - bool UseNative = AllowFTZAtomics; + bool UseNative = IgnoreFTZMismatch; switch (AI->getPointerAddressSpace()) { case llvm::ADDRESS_SPACE_GLOBAL: UseNative |= FTZ; @@ -7507,7 +7515,7 @@ NVPTXTargetLowering::shouldExpandAtomicRMWInIR(const AtomicRMWInst *AI) const { // 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) && + if ((!FTZ || IgnoreFTZMismatch) && STI.hasFeature(NVPTX::SM70) && STI.hasFeature(NVPTX::PTX63)) return AtomicExpansionKind::None; } diff --git a/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll b/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll new file mode 100644 index 0000000000000..f9d90c50b5886 --- /dev/null +++ b/llvm/test/CodeGen/NVPTX/atomicrmw-ignore-denormal-mode.ll @@ -0,0 +1,258 @@ +; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 6 +; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | FileCheck %s --check-prefixes=CHECK,STRICT +; RUN: llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 | FileCheck %s --check-prefixes=CHECK,DEFAULT +; RUN: %if ptxas-sm_90 && ptxas-isa-8.7 %{ llc < %s -march=nvptx64 -mcpu=sm_90 -mattr=+ptx87 -nvptx-allow-ftz-atomics=false | %ptxas-verify -arch=sm_90 %} + +; PTX atom.add has a fixed FTZ behavior that need not match the function's +; denormal mode: +; - atom.add.f32 on global memory flushes denormals +; - atom.add.f32 on shared memory does not +; - atom.add.f16 never flushes denormals +; When the two disagree the atomic is expanded into a CAS loop, unless the +; mismatch is explicitly declared insignificant -- either globally via +; -nvptx-allow-ftz-atomics (on by default, hence the STRICT run lines above, +; without which the metadata would never be the deciding factor) or per +; instruction via !atomic.ignore.denormal.mode. + +; f32, global: the native instruction flushes, so an IEEE function needs the +; metadata to avoid the CAS loop. + +define float @fadd_f32_global_ieee_no_md(ptr addrspace(1) %addr, float %val) { +; STRICT-LABEL: fadd_f32_global_ieee_no_md( +; STRICT: { +; STRICT-NEXT: .reg .pred %p<2>; +; STRICT-NEXT: .reg .b32 %r<5>; +; STRICT-NEXT: .reg .b64 %rd<2>; +; STRICT-EMPTY: +; STRICT-NEXT: // %bb.0: +; STRICT-NEXT: ld.param.b32 %r2, [fadd_f32_global_ieee_no_md_param_1]; +; STRICT-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ieee_no_md_param_0]; +; STRICT-NEXT: ld.relaxed.sys.global.b32 %r4, [%rd1]; +; STRICT-NEXT: $L__BB0_1: // %atomicrmw.start +; STRICT-NEXT: // =>This Inner Loop Header: Depth=1 +; STRICT-NEXT: add.rn.f32 %r3, %r4, %r2; +; STRICT-NEXT: atom.relaxed.sys.global.cas.b32 %r1, [%rd1], %r4, %r3; +; STRICT-NEXT: setp.ne.b32 %p1, %r1, %r4; +; STRICT-NEXT: mov.b32 %r4, %r1; +; STRICT-NEXT: @%p1 bra $L__BB0_1; +; STRICT-NEXT: // %bb.2: // %atomicrmw.end +; STRICT-NEXT: st.param.b32 [func_retval0], %r1; +; STRICT-NEXT: ret; +; +; DEFAULT-LABEL: fadd_f32_global_ieee_no_md( +; DEFAULT: { +; DEFAULT-NEXT: .reg .b32 %r<3>; +; DEFAULT-NEXT: .reg .b64 %rd<2>; +; DEFAULT-EMPTY: +; DEFAULT-NEXT: // %bb.0: +; DEFAULT-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ieee_no_md_param_0]; +; DEFAULT-NEXT: ld.param.b32 %r1, [fadd_f32_global_ieee_no_md_param_1]; +; DEFAULT-NEXT: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1; +; DEFAULT-NEXT: st.param.b32 [func_retval0], %r2; +; DEFAULT-NEXT: ret; + %r = atomicrmw fadd ptr addrspace(1) %addr, float %val monotonic + ret float %r +} + +define float @fadd_f32_global_ieee_md(ptr addrspace(1) %addr, float %val) { +; CHECK-LABEL: fadd_f32_global_ieee_md( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ieee_md_param_0]; +; CHECK-NEXT: ld.param.b32 %r1, [fadd_f32_global_ieee_md_param_1]; +; CHECK-NEXT: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %r = atomicrmw fadd ptr addrspace(1) %addr, float %val monotonic, !atomic.ignore.denormal.mode !0 + ret float %r +} + +; An FTZ function already agrees with global atom.add, so no metadata needed. + +define float @fadd_f32_global_ftz_no_md(ptr addrspace(1) %addr, float %val) #0 { +; CHECK-LABEL: fadd_f32_global_ftz_no_md( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_global_ftz_no_md_param_0]; +; CHECK-NEXT: ld.param.b32 %r1, [fadd_f32_global_ftz_no_md_param_1]; +; CHECK-NEXT: atom.relaxed.sys.global.add.f32 %r2, [%rd1], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %r = atomicrmw fadd ptr addrspace(1) %addr, float %val monotonic + ret float %r +} + +; f32, generic: no address space guarantee at all, so only the metadata (or the +; option) can authorize the native instruction. + +define float @fadd_f32_generic_ieee_no_md(ptr %addr, float %val) { +; STRICT-LABEL: fadd_f32_generic_ieee_no_md( +; STRICT: { +; STRICT-NEXT: .reg .pred %p<2>; +; STRICT-NEXT: .reg .b32 %r<5>; +; STRICT-NEXT: .reg .b64 %rd<2>; +; STRICT-EMPTY: +; STRICT-NEXT: // %bb.0: +; STRICT-NEXT: ld.param.b32 %r2, [fadd_f32_generic_ieee_no_md_param_1]; +; STRICT-NEXT: ld.param.b64 %rd1, [fadd_f32_generic_ieee_no_md_param_0]; +; STRICT-NEXT: ld.relaxed.sys.b32 %r4, [%rd1]; +; STRICT-NEXT: $L__BB3_1: // %atomicrmw.start +; STRICT-NEXT: // =>This Inner Loop Header: Depth=1 +; STRICT-NEXT: add.rn.f32 %r3, %r4, %r2; +; STRICT-NEXT: atom.relaxed.sys.cas.b32 %r1, [%rd1], %r4, %r3; +; STRICT-NEXT: setp.ne.b32 %p1, %r1, %r4; +; STRICT-NEXT: mov.b32 %r4, %r1; +; STRICT-NEXT: @%p1 bra $L__BB3_1; +; STRICT-NEXT: // %bb.2: // %atomicrmw.end +; STRICT-NEXT: st.param.b32 [func_retval0], %r1; +; STRICT-NEXT: ret; +; +; DEFAULT-LABEL: fadd_f32_generic_ieee_no_md( +; DEFAULT: { +; DEFAULT-NEXT: .reg .b32 %r<3>; +; DEFAULT-NEXT: .reg .b64 %rd<2>; +; DEFAULT-EMPTY: +; DEFAULT-NEXT: // %bb.0: +; DEFAULT-NEXT: ld.param.b64 %rd1, [fadd_f32_generic_ieee_no_md_param_0]; +; DEFAULT-NEXT: ld.param.b32 %r1, [fadd_f32_generic_ieee_no_md_param_1]; +; DEFAULT-NEXT: atom.relaxed.sys.add.f32 %r2, [%rd1], %r1; +; DEFAULT-NEXT: st.param.b32 [func_retval0], %r2; +; DEFAULT-NEXT: ret; + %r = atomicrmw fadd ptr %addr, float %val monotonic + ret float %r +} + +define float @fadd_f32_generic_ieee_md(ptr %addr, float %val) { +; CHECK-LABEL: fadd_f32_generic_ieee_md( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_generic_ieee_md_param_0]; +; CHECK-NEXT: ld.param.b32 %r1, [fadd_f32_generic_ieee_md_param_1]; +; CHECK-NEXT: atom.relaxed.sys.add.f32 %r2, [%rd1], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %r = atomicrmw fadd ptr %addr, float %val monotonic, !atomic.ignore.denormal.mode !0 + ret float %r +} + +; f32, shared: the native instruction does not flush, matching an IEEE function. + +define float @fadd_f32_shared_ieee_no_md(ptr addrspace(3) %addr, float %val) { +; CHECK-LABEL: fadd_f32_shared_ieee_no_md( +; CHECK: { +; CHECK-NEXT: .reg .b32 %r<3>; +; CHECK-NEXT: .reg .b64 %rd<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f32_shared_ieee_no_md_param_0]; +; CHECK-NEXT: ld.param.b32 %r1, [fadd_f32_shared_ieee_no_md_param_1]; +; CHECK-NEXT: atom.relaxed.sys.shared.add.f32 %r2, [%rd1], %r1; +; CHECK-NEXT: st.param.b32 [func_retval0], %r2; +; CHECK-NEXT: ret; + %r = atomicrmw fadd ptr addrspace(3) %addr, float %val monotonic + ret float %r +} + +; f16: atom.add.f16 never flushes, so it disagrees with an FTZ function. + +define half @fadd_f16_global_ftz_no_md(ptr addrspace(1) %addr, half %val) #1 { +; STRICT-LABEL: fadd_f16_global_ftz_no_md( +; STRICT: { +; STRICT-NEXT: .reg .pred %p<2>; +; STRICT-NEXT: .reg .b16 %rs<4>; +; STRICT-NEXT: .reg .b32 %r<15>; +; STRICT-NEXT: .reg .b64 %rd<3>; +; STRICT-EMPTY: +; STRICT-NEXT: // %bb.0: +; STRICT-NEXT: ld.param.b16 %rs1, [fadd_f16_global_ftz_no_md_param_1]; +; STRICT-NEXT: ld.param.b64 %rd2, [fadd_f16_global_ftz_no_md_param_0]; +; STRICT-NEXT: and.b64 %rd1, %rd2, -4; +; STRICT-NEXT: cvt.u32.u64 %r4, %rd2; +; STRICT-NEXT: and.b32 %r5, %r4, 3; +; STRICT-NEXT: shl.b32 %r1, %r5, 3; +; STRICT-NEXT: mov.b32 %r6, 65535; +; STRICT-NEXT: shl.b32 %r7, %r6, %r1; +; STRICT-NEXT: not.b32 %r2, %r7; +; STRICT-NEXT: ld.relaxed.sys.global.b32 %r14, [%rd1]; +; STRICT-NEXT: $L__BB6_1: // %atomicrmw.start +; STRICT-NEXT: // =>This Inner Loop Header: Depth=1 +; STRICT-NEXT: shr.u32 %r8, %r14, %r1; +; STRICT-NEXT: cvt.u16.u32 %rs2, %r8; +; STRICT-NEXT: add.rn.ftz.f16 %rs3, %rs2, %rs1; +; STRICT-NEXT: cvt.u32.u16 %r9, %rs3; +; STRICT-NEXT: shl.b32 %r10, %r9, %r1; +; STRICT-NEXT: and.b32 %r11, %r14, %r2; +; STRICT-NEXT: or.b32 %r12, %r11, %r10; +; STRICT-NEXT: atom.relaxed.sys.global.cas.b32 %r3, [%rd1], %r14, %r12; +; STRICT-NEXT: setp.ne.b32 %p1, %r3, %r14; +; STRICT-NEXT: mov.b32 %r14, %r3; +; STRICT-NEXT: @%p1 bra $L__BB6_1; +; STRICT-NEXT: // %bb.2: // %atomicrmw.end +; STRICT-NEXT: shr.u32 %r13, %r3, %r1; +; STRICT-NEXT: st.param.b16 [func_retval0], %r13; +; STRICT-NEXT: ret; +; +; DEFAULT-LABEL: fadd_f16_global_ftz_no_md( +; DEFAULT: { +; DEFAULT-NEXT: .reg .b16 %rs<3>; +; DEFAULT-NEXT: .reg .b64 %rd<2>; +; DEFAULT-EMPTY: +; DEFAULT-NEXT: // %bb.0: +; DEFAULT-NEXT: ld.param.b64 %rd1, [fadd_f16_global_ftz_no_md_param_0]; +; DEFAULT-NEXT: ld.param.b16 %rs1, [fadd_f16_global_ftz_no_md_param_1]; +; DEFAULT-NEXT: atom.relaxed.sys.global.add.noftz.f16 %rs2, [%rd1], %rs1; +; DEFAULT-NEXT: st.param.b16 [func_retval0], %rs2; +; DEFAULT-NEXT: ret; + %r = atomicrmw fadd ptr addrspace(1) %addr, half %val monotonic + ret half %r +} + +define half @fadd_f16_global_ftz_md(ptr addrspace(1) %addr, half %val) #1 { +; CHECK-LABEL: fadd_f16_global_ftz_md( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b64 %rd<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f16_global_ftz_md_param_0]; +; CHECK-NEXT: ld.param.b16 %rs1, [fadd_f16_global_ftz_md_param_1]; +; CHECK-NEXT: atom.relaxed.sys.global.add.noftz.f16 %rs2, [%rd1], %rs1; +; CHECK-NEXT: st.param.b16 [func_retval0], %rs2; +; CHECK-NEXT: ret; + %r = atomicrmw fadd ptr addrspace(1) %addr, half %val monotonic, !atomic.ignore.denormal.mode !0 + ret half %r +} + +; The f16 decision must come from the f16 denormal mode, not the f32 one. Here +; only f32 is flushed, so f16 is still IEEE and agrees with atom.add.f16 -- no +; metadata and no CAS loop. + +define half @fadd_f16_global_f32_ftz_only_no_md(ptr addrspace(1) %addr, half %val) #0 { +; CHECK-LABEL: fadd_f16_global_f32_ftz_only_no_md( +; CHECK: { +; CHECK-NEXT: .reg .b16 %rs<3>; +; CHECK-NEXT: .reg .b64 %rd<2>; +; CHECK-EMPTY: +; CHECK-NEXT: // %bb.0: +; CHECK-NEXT: ld.param.b64 %rd1, [fadd_f16_global_f32_ftz_only_no_md_param_0]; +; CHECK-NEXT: ld.param.b16 %rs1, [fadd_f16_global_f32_ftz_only_no_md_param_1]; +; CHECK-NEXT: atom.relaxed.sys.global.add.noftz.f16 %rs2, [%rd1], %rs1; +; CHECK-NEXT: st.param.b16 [func_retval0], %rs2; +; CHECK-NEXT: ret; + %r = atomicrmw fadd ptr addrspace(1) %addr, half %val monotonic + ret half %r +} + +attributes #0 = { denormal_fpenv(float: preservesign) } +attributes #1 = { denormal_fpenv(preservesign) } + +!0 = !{} _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
