[llvm-branch-commits] [llvm] [NVPTX] Honor !atomic.ignore.denormal.mode on atomicrmw fadd (PR #217586)
Christian Sigg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Aug 20 10:18:38 PDT 2026
https://github.com/chsigg updated https://github.com/llvm/llvm-project/pull/217586
>From 636a3c663b6b346248e0d03a2d1f37cb40b47029 Mon Sep 17 00:00:00 2001
From: Christian Sigg <csigg at google.com>
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.
Co-authored-by: Artem Belevich <tra at google.com>
---
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 6798875f5fec5..da70d61d2af2c 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"
@@ -7494,11 +7495,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;
@@ -7517,7 +7525,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 = !{}
More information about the llvm-branch-commits
mailing list