Repository navigation
[NVPTX] Attach !atomic.ignore.denormal.mode when upgrading legacy NVVM atomic add intrinsics - #232725
Merged
chsigg merged 1 commit intoOct 17, 2026
Conversation
Contributor
Author
This stack of pull requests is managed by Graphite. Learn more about stacking. |
|
@llvm/pr-subscribers-llvm-ir @llvm/pr-subscribers-backend-nvptx Author: Christian Sigg (chsigg) ChangesFull diff: https://github.com/llvm/llvm-project/pull/232725.diff 3 Files Affected:
diff --git a/llvm/lib/IR/AutoUpgrade.cpp b/llvm/lib/IR/AutoUpgrade.cpp
index afb1cd1d48960..4e754e34fcc64 100644
--- a/llvm/lib/IR/AutoUpgrade.cpp
+++ b/llvm/lib/IR/AutoUpgrade.cpp
@@ -3317,6 +3317,9 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
Rep = Builder.CreateAtomicRMW(
AtomicRMWInst::FAdd, Ptr, Val, MaybeAlign(), AtomicOrdering::Monotonic,
CI->getContext().getOrInsertSyncScopeID("device"));
+ cast<Instruction>(Rep)->setMetadata(
+ LLVMContext::MD_atomic_ignore_denormal_mode,
+ MDNode::get(CI->getContext(), {}));
// The default scope for atomic.load.* intrinsics is device
// (= gpu scope in ptx), but the default LLVM atomic scope is
// "system"
@@ -3363,6 +3366,11 @@ static Value *upgradeNVVMIntrinsicCall(StringRef Name, CallBase *CI,
"unexpected nvvm scoped atomic intrinsic");
Rep = Builder.CreateAtomicRMW(BinOp, Ptr, Val, MaybeAlign(),
AtomicOrdering::Monotonic, SSID);
+ if (BinOp == AtomicRMWInst::FAdd) {
+ cast<Instruction>(Rep)->setMetadata(
+ LLVMContext::MD_atomic_ignore_denormal_mode,
+ MDNode::get(CI->getContext(), {}));
+ }
}
} else if (Name == "clz.ll") {
// llvm.nvvm.clz.ll returns an i32, but llvm.ctlz.i64 returns an i64.
diff --git a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
index a2174877ac091..1e201281cecfd 100644
--- a/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
+++ b/llvm/test/Assembler/auto_upgrade_nvvm_intrinsics.ll
@@ -399,8 +399,8 @@ define void @ldg(ptr %p0, ptr addrspace(1) %p1) {
define i32 @atomics(ptr %p0, i32 %a, float %b, double %c) {
; CHECK: %1 = atomicrmw uinc_wrap ptr %p0, i32 %a syncscope("device") monotonic
; CHECK: %2 = atomicrmw udec_wrap ptr %p0, i32 %a syncscope("device") monotonic
-; CHECK: %3 = atomicrmw fadd ptr %p0, float %b syncscope("device") monotonic
-; CHECK: %4 = atomicrmw fadd ptr %p0, double %c syncscope("device") monotonic
+; CHECK: %3 = atomicrmw fadd ptr %p0, float %b syncscope("device") monotonic, align 4, !atomic.ignore.denormal.mode !0
+; CHECK: %4 = atomicrmw fadd ptr %p0, double %c syncscope("device") monotonic, align 8, !atomic.ignore.denormal.mode !0
%r1 = call i32 @llvm.nvvm.atomic.load.inc.32.p0(ptr %p0, i32 %a)
%r2 = call i32 @llvm.nvvm.atomic.load.dec.32.p0(ptr %p0, i32 %a)
@@ -501,13 +501,13 @@ define void @atomics_scoped(ptr %p, i32 %a, i64 %b, float %c, double %d) {
; CHECK: [[CAS_sys_i64:%.*]] = cmpxchg ptr %p, i64 %b, i64 %b monotonic monotonic
; CHECK: extractvalue { i64, i1 } [[CAS_sys_i64]], 0
%cas.i.sys.i64 = call i64 @llvm.nvvm.atomic.cas.gen.i.sys.i64.p0(ptr %p, i64 %b, i64 %b)
-; CHECK: atomicrmw fadd ptr %p, float %c syncscope("block") monotonic
+; CHECK: atomicrmw fadd ptr %p, float %c syncscope("block") monotonic, align 4, !atomic.ignore.denormal.mode !0
%add.f.cta.f32 = call float @llvm.nvvm.atomic.add.gen.f.cta.f32.p0(ptr %p, float %c)
-; CHECK: atomicrmw fadd ptr %p, double %d syncscope("block") monotonic
+; CHECK: atomicrmw fadd ptr %p, double %d syncscope("block") monotonic, align 8, !atomic.ignore.denormal.mode !0
%add.f.cta.f64 = call double @llvm.nvvm.atomic.add.gen.f.cta.f64.p0(ptr %p, double %d)
-; CHECK: atomicrmw fadd ptr %p, float %c monotonic
+; CHECK: atomicrmw fadd ptr %p, float %c monotonic, align 4, !atomic.ignore.denormal.mode !0
%add.f.sys.f32 = call float @llvm.nvvm.atomic.add.gen.f.sys.f32.p0(ptr %p, float %c)
-; CHECK: atomicrmw fadd ptr %p, double %d monotonic
+; CHECK: atomicrmw fadd ptr %p, double %d monotonic, align 8, !atomic.ignore.denormal.mode !0
%add.f.sys.f64 = call double @llvm.nvvm.atomic.add.gen.f.sys.f64.p0(ptr %p, double %d)
ret void
}
diff --git a/llvm/test/CodeGen/NVPTX/atomics.ll b/llvm/test/CodeGen/NVPTX/atomics.ll
index e1853f1b00cf3..cb708fd855c77 100644
--- a/llvm/test/CodeGen/NVPTX/atomics.ll
+++ b/llvm/test/CodeGen/NVPTX/atomics.ll
@@ -11,38 +11,17 @@ declare float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)
; CHECK-LABEL: atomic_add_f32_generic
define float @atomic_add_f32_generic(ptr %addr, float %val) {
;
-; DISALLOW-LABEL: atomic_add_f32_generic(
-; DISALLOW: {
-; DISALLOW-NEXT: .reg .pred %p<2>;
-; DISALLOW-NEXT: .reg .b32 %r<5>;
-; DISALLOW-NEXT: .reg .b64 %rd<2>;
-; DISALLOW-EMPTY:
-; DISALLOW-NEXT: // %bb.0:
-; DISALLOW-NEXT: ld.param.b32 %r2, [atomic_add_f32_generic_param_1];
-; DISALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_generic_param_0];
-; DISALLOW-NEXT: ld.volatile.b32 %r4, [%rd1];
-; DISALLOW-NEXT: $L__BB0_1: // %atomicrmw.start
-; DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; DISALLOW-NEXT: add.rn.f32 %r3, %r4, %r2;
-; DISALLOW-NEXT: atom.cas.b32 %r1, [%rd1], %r4, %r3;
-; DISALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; DISALLOW-NEXT: mov.b32 %r4, %r1;
-; DISALLOW-NEXT: @%p1 bra $L__BB0_1;
-; DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; DISALLOW-NEXT: st.param.b32 [func_retval0], %r1;
-; DISALLOW-NEXT: ret;
-;
-; ALLOW-LABEL: atomic_add_f32_generic(
-; ALLOW: {
-; ALLOW-NEXT: .reg .b32 %r<3>;
-; ALLOW-NEXT: .reg .b64 %rd<2>;
-; ALLOW-EMPTY:
-; ALLOW-NEXT: // %bb.0:
-; ALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_generic_param_0];
-; ALLOW-NEXT: ld.param.b32 %r1, [atomic_add_f32_generic_param_1];
-; ALLOW-NEXT: atom.add.f32 %r2, [%rd1], %r1;
-; ALLOW-NEXT: st.param.b32 [func_retval0], %r2;
-; ALLOW-NEXT: ret;
+; CHECK-LABEL: atomic_add_f32_generic(
+; 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, [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: ret;
%ret = call float @llvm.nvvm.atomic.load.add.f32.p0(ptr %addr, float %val)
ret float %ret
}
@@ -52,38 +31,17 @@ declare float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %v
; CHECK-LABEL: atomic_add_f32_addrspace1
define float @atomic_add_f32_addrspace1(ptr addrspace(1) %addr, float %val) {
;
-; DISALLOW-LABEL: atomic_add_f32_addrspace1(
-; DISALLOW: {
-; DISALLOW-NEXT: .reg .pred %p<2>;
-; DISALLOW-NEXT: .reg .b32 %r<5>;
-; DISALLOW-NEXT: .reg .b64 %rd<2>;
-; DISALLOW-EMPTY:
-; DISALLOW-NEXT: // %bb.0:
-; DISALLOW-NEXT: ld.param.b32 %r2, [atomic_add_f32_addrspace1_param_1];
-; DISALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace1_param_0];
-; DISALLOW-NEXT: ld.volatile.global.b32 %r4, [%rd1];
-; DISALLOW-NEXT: $L__BB1_1: // %atomicrmw.start
-; DISALLOW-NEXT: // =>This Inner Loop Header: Depth=1
-; DISALLOW-NEXT: add.rn.f32 %r3, %r4, %r2;
-; DISALLOW-NEXT: atom.global.cas.b32 %r1, [%rd1], %r4, %r3;
-; DISALLOW-NEXT: setp.ne.b32 %p1, %r1, %r4;
-; DISALLOW-NEXT: mov.b32 %r4, %r1;
-; DISALLOW-NEXT: @%p1 bra $L__BB1_1;
-; DISALLOW-NEXT: // %bb.2: // %atomicrmw.end
-; DISALLOW-NEXT: st.param.b32 [func_retval0], %r1;
-; DISALLOW-NEXT: ret;
-;
-; ALLOW-LABEL: atomic_add_f32_addrspace1(
-; ALLOW: {
-; ALLOW-NEXT: .reg .b32 %r<3>;
-; ALLOW-NEXT: .reg .b64 %rd<2>;
-; ALLOW-EMPTY:
-; ALLOW-NEXT: // %bb.0:
-; ALLOW-NEXT: ld.param.b64 %rd1, [atomic_add_f32_addrspace1_param_0];
-; ALLOW-NEXT: ld.param.b32 %r1, [atomic_add_f32_addrspace1_param_1];
-; ALLOW-NEXT: atom.global.add.f32 %r2, [%rd1], %r1;
-; ALLOW-NEXT: st.param.b32 [func_retval0], %r2;
-; ALLOW-NEXT: ret;
+; CHECK-LABEL: atomic_add_f32_addrspace1(
+; 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, [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: ret;
%ret = call float @llvm.nvvm.atomic.load.add.f32.p1(ptr addrspace(1) %addr, float %val)
ret float %ret
}
@@ -106,3 +64,6 @@ define float @atomic_add_f32_addrspace3(ptr addrspace(3) %addr, float %val) {
%ret = call float @llvm.nvvm.atomic.load.add.f32.p3(ptr addrspace(3) %addr, float %val)
ret float %ret
}
+;; NOTE: These prefixes are unused and the list is autogenerated. Do not add tests below this line:
+; ALLOW: {{.*}}
+; DISALLOW: {{.*}}
|
Contributor
Author
Merge activity
|
chsigg
force-pushed
the
users/chsigg/nvvm-autoupgrade-atomic-ignore-denormal-mode
branch
6 times, most recently
from
October 17, 2026 07:29
1366bc8 to
6b8a837
Compare
…M atomic add intrinsics
chsigg
force-pushed
the
users/chsigg/nvvm-autoupgrade-atomic-ignore-denormal-mode
branch
from
October 17, 2026 15:55
6b8a837 to
0c314a6
Compare
chsigg
enabled auto-merge (squash)
October 17, 2026 16:01
chsigg
deleted the
users/chsigg/nvvm-autoupgrade-atomic-ignore-denormal-mode
branch
October 17, 2026 16:44
chsigg
added a commit
that referenced
this pull request
Oct 18, 2026
This restores the strict FTZ default originally introduced in #200732 and temporarily relaxed in #206154. Frontends can now explicitly opt in to native FTZ-mismatched atomics on a per-instruction basis using `!atomic.ignore.denormal.mode` metadata (#217585, #217586, #217587, #217588, #232725, #232726), which is supported in Clang, Flang, MLIR, and XLA. Users who need the previous behavior globally can still pass `-nvptx-allow-ftz-atomics=true`. Update test cases to verify the new default disallow behavior.
alanzhao1
pushed a commit
to alanzhao1/llvm-project
that referenced
this pull request
Oct 25, 2026
…31429) This restores the strict FTZ default originally introduced in llvm#200732 and temporarily relaxed in llvm#206154. Frontends can now explicitly opt in to native FTZ-mismatched atomics on a per-instruction basis using `!atomic.ignore.denormal.mode` metadata (llvm#217585, llvm#217586, llvm#217587, llvm#217588, llvm#232725, llvm#232726), which is supported in Clang, Flang, MLIR, and XLA. Users who need the previous behavior globally can still pass `-nvptx-allow-ftz-atomics=true`. Update test cases to verify the new default disallow behavior.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Legacy NVVM floating-point atomic add intrinsics (
@llvm.nvvm.atomic.load.add.f32.p*,@llvm.nvvm.atomic.load.add.f64.p*, and@llvm.nvvm.atomic.add.gen.f.*) specifically targeted the PTXatom.addinstruction, which flushes f32 denormals to zero on global memory.Attach
!atomic.ignore.denormal.modemetadata when auto-upgrading these intrinsics toatomicrmw faddso that they continue to lower to nativeatom.addinstructions when-nvptx-allow-ftz-atomicsdefaults tofalse(#231429), matching the Clang lowering of__nvvm_atom_add_gen_f(#217587).