Repository navigation
[clang][NVPTX] Emit !atomic.ignore.denormal.mode for CUDA atomics - #217587
Merged
Merged
Conversation
This was referenced Aug 20, 2026
Contributor
Author
This was referenced Aug 20, 2026
🐧 Linux x64 Test Results
✅ The build succeeded and all tests passed. |
🪟 Windows x64 Test Results
✅ The build succeeded and all tests passed. |
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
from
August 20, 2026 16:38
2364e44 to
f654917
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
2 times, most recently
from
August 20, 2026 16:47
6494a10 to
cbc913a
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
from
August 20, 2026 16:47
f654917 to
6a5536a
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
August 20, 2026 17:18
cbc913a to
bafeed9
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
2 times, most recently
from
August 20, 2026 19:44
636a3c6 to
32accb8
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
August 20, 2026 19:44
bafeed9 to
5d2cb45
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
from
August 21, 2026 07:02
32accb8 to
4fad0ff
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
August 21, 2026 07:02
5d2cb45 to
bc3d073
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
from
August 24, 2026 10:18
4fad0ff to
1774a54
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
2 times, most recently
from
August 24, 2026 12:15
7c7705b to
e4aaa27
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
2 times, most recently
from
August 24, 2026 12:27
7d9fc10 to
5a0ac93
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
August 24, 2026 12:27
e4aaa27 to
1c1276b
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
September 15, 2026 15:17
1c1276b to
5a2f006
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
from
September 15, 2026 15:17
5a0ac93 to
5ab17f9
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
September 15, 2026 19:21
5a2f006 to
7100407
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
2 times, most recently
from
September 16, 2026 07:08
ba5bdf2 to
9f1578b
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
September 16, 2026 07:08
7100407 to
36bae08
Compare
chsigg
force-pushed
the
users/chsigg/nvptx-honor-atomic-ignore-denormal
branch
from
October 7, 2026 13:10
9f1578b to
507f6bf
Compare
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
October 7, 2026 13:10
36bae08 to
3e509d8
Compare
Base automatically changed from
users/chsigg/nvptx-honor-atomic-ignore-denormal
to
main
October 7, 2026 14:05
CUDA's atomicAdd() family is defined in terms of PTX atom.add, whose denormal behavior is fixed by the hardware. Without any annotation the backend has to assume the function's denormal mode must be honored and expands these into CAS loops whenever the two disagree. Mark them with !atomic.ignore.denormal.mode so the native instruction is used. That covers the __nvvm_atom_*_add_gen_f builtins that atomicAdd(), atomicAdd_block() and atomicAdd_system() are written in terms of, plus C11/C++11 atomics under -fatomic-ignore-denormal-mode and the [[clang::atomic(ignore_denormal_mode)]] attribute, which requires teaching the NVPTX target about AtomicOptions. The condition for when the metadata is meaningful is now shared with the AMDGPU and SPIR-V targets in addAtomicIgnoreDenormalModeMetadata(). It takes an AllowHalf flag because whether f16 denormals are observable is target specific: PTX exposes no FTZ control for f16 operations, so atom.add.f16 never flushes and the opt-in is meaningful there, whereas the AMDGPU backend only consults the metadata for f32. Co-authored-by: Artem Belevich <tra@google.com>
chsigg
marked this pull request as ready for review
October 7, 2026 15:05
chsigg
force-pushed
the
users/chsigg/clang-emit-atomic-ignore-denormal
branch
from
October 7, 2026 15:05
3e509d8 to
081fbc3
Compare
llvmorg-github-actions
Bot
added
clang
|
@llvm/pr-subscribers-clang @llvm/pr-subscribers-backend-amdgpu Author: Christian Sigg (chsigg) ChangesCUDA's atomicAdd() family is defined in terms of PTX atom.add, whose That covers the _nvvm_atom*_add_gen_f builtins that atomicAdd(), The condition for when the metadata is meaningful is now shared with the Patch is 28.40 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/217587.diff 11 Files Affected:
diff --git a/clang/lib/Basic/Targets/NVPTX.cpp b/clang/lib/Basic/Targets/NVPTX.cpp
index de99a6718f26c..61efae2ccae87 100644
--- a/clang/lib/Basic/Targets/NVPTX.cpp
+++ b/clang/lib/Basic/Targets/NVPTX.cpp
@@ -206,6 +206,12 @@ void NVPTXTargetInfo::getTargetDefines(const LangOptions &Opts,
}
}
+void NVPTXTargetInfo::adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
+ const TargetInfo *Aux) {
+ TargetInfo::adjust(Diags, Opts, Aux);
+ AtomicOpts = AtomicOptions(Opts);
+}
+
llvm::SmallVector<Builtin::InfosShard>
NVPTXTargetInfo::getTargetBuiltins() const {
return {{&BuiltinStrings, BuiltinInfos}};
diff --git a/clang/lib/Basic/Targets/NVPTX.h b/clang/lib/Basic/Targets/NVPTX.h
index b6d2bfd320103..e55fe89712385 100644
--- a/clang/lib/Basic/Targets/NVPTX.h
+++ b/clang/lib/Basic/Targets/NVPTX.h
@@ -147,6 +147,8 @@ class LLVM_LIBRARY_VISIBILITY NVPTXTargetInfo : public TargetInfo {
GPU = StringToOffloadArch(Name);
return !GPU.isUnknown();
}
+ void adjust(DiagnosticsEngine &Diags, LangOptions &Opts,
+ const TargetInfo *Aux) override;
void setSupportedOpenCLOpts() override {
auto &Opts = getSupportedOpenCLOpts();
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index f3bf71dddc341..2f29e9557ba56 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -21,7 +21,6 @@
#include "llvm/IR/IntrinsicsAMDGPU.h"
#include "llvm/IR/IntrinsicsR600.h"
#include "llvm/IR/IntrinsicsSPIRV.h"
-#include "llvm/IR/LLVMContext.h"
#include "llvm/IR/MemoryModelRelaxationAnnotations.h"
#include "llvm/Support/AMDGPUAddrSpace.h"
#include "llvm/Support/AtomicOrdering.h"
@@ -2047,9 +2046,7 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
// Most targets require "atomic.ignore.denormal.mode" to emit the native
// instruction, but this only matters for float fadd.
- if (BinOp == llvm::AtomicRMWInst::FAdd && Val->getType()->isFloatTy())
- RMW->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode,
- EmptyMD);
+ addAtomicIgnoreDenormalModeMetadata(*this, *RMW);
}
return Builder.CreateBitCast(RMW, OrigTy);
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index e304eed8a1aab..2781d6b1dbd9b 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -11,6 +11,7 @@
//===----------------------------------------------------------------------===//
#include "CGBuiltin.h"
+#include "TargetInfo.h"
#include "clang/Basic/TargetBuiltins.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
#include "llvm/TargetParser/AtomicScope.h"
@@ -362,8 +363,15 @@ static Value *MakeScopedAtomicRMW(CodeGenFunction &CGF, const CallExpr *E,
Value *Val = CGF.EmitScalarExpr(E->getArg(1));
llvm::SyncScope::ID SSID = CGF.getLLVMContext().getOrInsertSyncScopeID(
*llvm::getAtomicScopeIRString(CGF.getTarget().getTriple(), Scope));
- return CGF.Builder.CreateAtomicRMW(Kind, Ptr, Val,
- llvm::AtomicOrdering::Monotonic, SSID);
+ llvm::AtomicRMWInst *RMW = CGF.Builder.CreateAtomicRMW(
+ Kind, Ptr, Val, llvm::AtomicOrdering::Monotonic, SSID);
+
+ // CUDA's atomicAdd_block()/atomicAdd_system() reach this through the scoped
+ // add builtins, and are defined in terms of the native atom.add just as the
+ // unscoped atomicAdd() is. The helper ignores everything that is not a
+ // floating-point add, so the other scoped operations are unaffected.
+ addAtomicIgnoreDenormalModeMetadata(CGF, *RMW, /*AllowHalf=*/true);
+ return RMW;
}
// `Scope` is AtomicScope::Workgroup for _cta builtins and AtomicScope::System
@@ -529,8 +537,15 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
Address DestAddr = EmitPointerWithAlignment(E->getArg(0));
Value *Val = EmitScalarExpr(E->getArg(1));
- return Builder.CreateAtomicRMW(llvm::AtomicRMWInst::FAdd, DestAddr, Val,
- AtomicOrdering::Monotonic);
+ llvm::AtomicRMWInst *RMW = Builder.CreateAtomicRMW(
+ llvm::AtomicRMWInst::FAdd, DestAddr, Val, AtomicOrdering::Monotonic);
+
+ // CUDA's atomicAdd() is defined in terms of the native atom.add
+ // instruction, whose denormal behavior is fixed by the hardware and need
+ // not agree with the function's denormal mode. Without this the atomic is
+ // expanded into a CAS loop whenever the two disagree.
+ addAtomicIgnoreDenormalModeMetadata(*this, *RMW, /*AllowHalf=*/true);
+ return RMW;
}
case NVPTX::BI__nvvm_atom_inc_gen_ui:
diff --git a/clang/lib/CodeGen/TargetInfo.cpp b/clang/lib/CodeGen/TargetInfo.cpp
index 2f89154c2d2d9..9786077f9062c 100644
--- a/clang/lib/CodeGen/TargetInfo.cpp
+++ b/clang/lib/CodeGen/TargetInfo.cpp
@@ -160,6 +160,21 @@ TargetCodeGenInfo::getLLVMSyncScopeID(const LangOptions &LangOpts,
getLLVMSyncScopeStr(LangOpts, Scope, Ordering));
}
+void CodeGen::addAtomicIgnoreDenormalModeMetadata(CodeGenFunction &CGF,
+ llvm::Instruction &AtomicInst,
+ bool AllowHalf) {
+ auto *RMW = dyn_cast<llvm::AtomicRMWInst>(&AtomicInst);
+ if (!RMW || RMW->getOperation() != llvm::AtomicRMWInst::FAdd)
+ return;
+
+ llvm::Type *Ty = RMW->getType();
+ if (!Ty->isFloatTy() && !(AllowHalf && Ty->isHalfTy()))
+ return;
+
+ RMW->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode,
+ llvm::MDNode::get(CGF.getLLVMContext(), {}));
+}
+
void TargetCodeGenInfo::addStackProbeTargetAttributes(
const Decl *D, llvm::GlobalValue *GV, CodeGen::CodeGenModule &CGM) const {
if (llvm::Function *Fn = dyn_cast_or_null<llvm::Function>(GV)) {
diff --git a/clang/lib/CodeGen/TargetInfo.h b/clang/lib/CodeGen/TargetInfo.h
index fae32068b1570..3b1dad75f88df 100644
--- a/clang/lib/CodeGen/TargetInfo.h
+++ b/clang/lib/CodeGen/TargetInfo.h
@@ -499,6 +499,24 @@ class TargetCodeGenInfo {
CodeGen::CodeGenModule &CGM) const;
};
+/// Attach `!atomic.ignore.denormal.mode` to \p AtomicInst, if it is an
+/// `atomicrmw` the metadata is meaningful for, i.e. an `fadd` on a type whose
+/// native atomic has a fixed denormal behavior that may disagree with the
+/// function's denormal mode.
+///
+/// `float` always qualifies. `half` only does when \p AllowHalf is set, since
+/// whether f16 denormals are observable is target specific: NVPTX exposes no
+/// FTZ control for f16 operations, so atom.add.f16 never flushes and the
+/// per-instruction opt-in is meaningful, whereas the AMDGPU backend only
+/// consults the metadata for f32.
+///
+/// This is shared by the targets whose backends honor the metadata; it does
+/// not itself check whether the user asked for it (see
+/// AtomicOptionKind::IgnoreDenormalMode).
+void addAtomicIgnoreDenormalModeMetadata(CodeGenFunction &CGF,
+ llvm::Instruction &AtomicInst,
+ bool AllowHalf = false);
+
std::unique_ptr<TargetCodeGenInfo>
createDefaultTargetCodeGenInfo(CodeGenModule &CGM);
diff --git a/clang/lib/CodeGen/Targets/AMDGPU.cpp b/clang/lib/CodeGen/Targets/AMDGPU.cpp
index 0b5ed1898f138..a74f0cf22d8e5 100644
--- a/clang/lib/CodeGen/Targets/AMDGPU.cpp
+++ b/clang/lib/CodeGen/Targets/AMDGPU.cpp
@@ -581,10 +581,8 @@ void AMDGPUTargetCodeGenInfo::setTargetAtomicMetadata(
RMW->setMetadata("amdgpu.no.fine.grained.memory", Empty);
if (!AO.getOption(clang::AtomicOptionKind::RemoteMemory))
RMW->setMetadata("amdgpu.no.remote.memory", Empty);
- if (AO.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) &&
- RMW->getOperation() == llvm::AtomicRMWInst::FAdd &&
- RMW->getType()->isFloatTy())
- RMW->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode, Empty);
+ if (AO.getOption(clang::AtomicOptionKind::IgnoreDenormalMode))
+ addAtomicIgnoreDenormalModeMetadata(CGF, *RMW);
}
bool AMDGPUTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
diff --git a/clang/lib/CodeGen/Targets/NVPTX.cpp b/clang/lib/CodeGen/Targets/NVPTX.cpp
index 00fe271c85b0b..0229389fe1716 100644
--- a/clang/lib/CodeGen/Targets/NVPTX.cpp
+++ b/clang/lib/CodeGen/Targets/NVPTX.cpp
@@ -13,6 +13,7 @@
#include "llvm/ADT/StringExtras.h"
#include "llvm/IR/CallingConv.h"
#include "llvm/IR/IntrinsicsNVPTX.h"
+#include "llvm/IR/LLVMContext.h"
#include "llvm/Support/NVVMAttributes.h"
using namespace clang;
@@ -50,6 +51,9 @@ class NVPTXTargetCodeGenInfo : public TargetCodeGenInfo {
void setTargetAttributes(const Decl *D, llvm::GlobalValue *GV,
CodeGen::CodeGenModule &M) const override;
+ void setTargetAtomicMetadata(CodeGenFunction &CGF,
+ llvm::Instruction &AtomicInst,
+ const AtomicExpr *AE) const override;
bool shouldEmitStaticExternCAliases() const override;
StringRef getLLVMSyncScopeStr(const LangOptions &LangOpts, SyncScope Scope,
@@ -304,6 +308,14 @@ bool NVPTXTargetCodeGenInfo::shouldEmitStaticExternCAliases() const {
return false;
}
+void NVPTXTargetCodeGenInfo::setTargetAtomicMetadata(
+ CodeGenFunction &CGF, llvm::Instruction &AtomicInst,
+ const AtomicExpr *AE) const {
+ AtomicOptions AO = CGF.CGM.getAtomicOpts();
+ if (AO.getOption(clang::AtomicOptionKind::IgnoreDenormalMode))
+ addAtomicIgnoreDenormalModeMetadata(CGF, AtomicInst, /*AllowHalf=*/true);
+}
+
StringRef NVPTXTargetCodeGenInfo::getLLVMSyncScopeStr(
const LangOptions &LangOpts, SyncScope Scope,
llvm::AtomicOrdering Ordering) const {
diff --git a/clang/lib/CodeGen/Targets/SPIR.cpp b/clang/lib/CodeGen/Targets/SPIR.cpp
index 244dd10d1f7ca..de7880c641db3 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -588,10 +588,8 @@ void SPIRVTargetCodeGenInfo::setTargetAtomicMetadata(
RMW->setMetadata("amdgpu.no.fine.grained.memory", Empty);
if (!AO.getOption(clang::AtomicOptionKind::RemoteMemory))
RMW->setMetadata("amdgpu.no.remote.memory", Empty);
- if (AO.getOption(clang::AtomicOptionKind::IgnoreDenormalMode) &&
- RMW->getOperation() == llvm::AtomicRMWInst::FAdd &&
- RMW->getType()->isFloatTy())
- RMW->setMetadata(llvm::LLVMContext::MD_atomic_ignore_denormal_mode, Empty);
+ if (AO.getOption(clang::AtomicOptionKind::IgnoreDenormalMode))
+ addAtomicIgnoreDenormalModeMetadata(CGF, *RMW);
}
/// Construct a SPIR-V target extension type for the given OpenCL image type.
diff --git a/clang/test/CodeGen/builtins-nvptx.c b/clang/test/CodeGen/builtins-nvptx.c
index acd7d40277bfa..25dc7e15bf455 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -417,7 +417,10 @@ __device__ void nvvm_atom(float *fp, float f, double *dfp, double df,
// CHECK-NEXT: extractvalue { i64, i1 } {{%[0-9]+}}, 0
__nvvm_atom_cas_gen_ll(&sll, 0, ll);
- // CHECK: atomicrmw fadd ptr {{.*}} monotonic, align 4
+ // CUDA's atomicAdd() lowers to this builtin. The native atom.add has a fixed
+ // denormal behavior, so the atomic is marked to keep it from being expanded
+ // into a CAS loop when the function's denormal mode disagrees.
+ // CHECK: atomicrmw fadd ptr {{.*}} monotonic, align 4, !atomic.ignore.denormal.mode
__nvvm_atom_add_gen_f(fp, f);
// CHECK: atomicrmw uinc_wrap ptr {{.*}} monotonic, align 4
diff --git a/clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu b/clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu
new file mode 100644
index 0000000000000..e3ae34322a83d
--- /dev/null
+++ b/clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu
@@ -0,0 +1,321 @@
+// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --version 6
+// REQUIRES: nvptx-registered-target
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90 \
+// RUN: -target-feature +ptx87 -fcuda-is-device \
+// RUN: -emit-llvm -o - %s | FileCheck --check-prefix=DEV %s
+// RUN: %clang_cc1 -triple nvptx64-nvidia-cuda -target-cpu sm_90 \
+// RUN: -target-feature +ptx87 -fcuda-is-device \
+// RUN: -fatomic-ignore-denormal-mode \
+// RUN: -emit-llvm -o - %s | FileCheck --check-prefix=OPT %s
+
+#include "Inputs/cuda.h"
+
+// -fatomic-ignore-denormal-mode reaches NVPTX, and only marks float fadd --
+// the one case where the native instruction's fixed denormal behavior is
+// observable.
+
+// DEV-LABEL: define dso_local void @_Z10test_floatPf(
+// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4
+// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
+// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
+// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
+// DEV-NEXT: ret void
+//
+// OPT-LABEL: define dso_local void @_Z10test_floatPf(
+// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0:[0-9]+]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4
+// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
+// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2:![0-9]+]]
+// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
+// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
+// OPT-NEXT: ret void
+//
+__device__ void test_float(float *a) {
+ __scoped_atomic_fetch_add(a, 1.0f, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
+}
+
+// DEV-LABEL: define dso_local void @_Z11test_doublePd(
+// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca double, align 8
+// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca double, align 8
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: store double 1.000000e+00, ptr [[DOTATOMICTMP]], align 8
+// DEV-NEXT: [[TMP1:%.*]] = load double, ptr [[DOTATOMICTMP]], align 8
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], double [[TMP1]] monotonic, align 8
+// DEV-NEXT: store double [[TMP2]], ptr [[ATOMIC_TEMP]], align 8
+// DEV-NEXT: [[TMP3:%.*]] = load double, ptr [[ATOMIC_TEMP]], align 8
+// DEV-NEXT: ret void
+//
+// OPT-LABEL: define dso_local void @_Z11test_doublePd(
+// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca double, align 8
+// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca double, align 8
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: store double 1.000000e+00, ptr [[DOTATOMICTMP]], align 8
+// OPT-NEXT: [[TMP1:%.*]] = load double, ptr [[DOTATOMICTMP]], align 8
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], double [[TMP1]] monotonic, align 8
+// OPT-NEXT: store double [[TMP2]], ptr [[ATOMIC_TEMP]], align 8
+// OPT-NEXT: [[TMP3:%.*]] = load double, ptr [[ATOMIC_TEMP]], align 8
+// OPT-NEXT: ret void
+//
+__device__ void test_double(double *a) {
+ __scoped_atomic_fetch_add(a, 1.0, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
+}
+
+// The scoped attribute overrides the command line in both directions.
+
+// DEV-LABEL: define dso_local void @_Z12test_attr_onPf(
+// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4
+// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
+// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2:![0-9]+]]
+// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
+// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
+// DEV-NEXT: ret void
+//
+// OPT-LABEL: define dso_local void @_Z12test_attr_onPf(
+// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4
+// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
+// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
+// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
+// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
+// OPT-NEXT: ret void
+//
+__device__ void test_attr_on(float *a) {
+ [[clang::atomic(ignore_denormal_mode)]] {
+ __scoped_atomic_fetch_add(a, 1.0f, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
+ }
+}
+
+// DEV-LABEL: define dso_local void @_Z13test_attr_offPf(
+// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4
+// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
+// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// DEV-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
+// DEV-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
+// DEV-NEXT: ret void
+//
+// OPT-LABEL: define dso_local void @_Z13test_attr_offPf(
+// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca float, align 4
+// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca float, align 4
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: store float 1.000000e+00, ptr [[DOTATOMICTMP]], align 4
+// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[DOTATOMICTMP]], align 4
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4
+// OPT-NEXT: store float [[TMP2]], ptr [[ATOMIC_TEMP]], align 4
+// OPT-NEXT: [[TMP3:%.*]] = load float, ptr [[ATOMIC_TEMP]], align 4
+// OPT-NEXT: ret void
+//
+__device__ void test_attr_off(float *a) {
+ [[clang::atomic(no_ignore_denormal_mode)]] {
+ __scoped_atomic_fetch_add(a, 1.0f, __ATOMIC_RELAXED, __MEMORY_SCOPE_SYSTEM);
+ }
+}
+
+// CUDA's atomicAdd() lowers to __nvvm_atom_add_gen_f rather than going through
+// the C11 atomic path, so it is marked unconditionally: the native atom.add is
+// wha...
[truncated]
|
Artem-B
approved these changes
Oct 10, 2026
|
LLVM Buildbot has detected a new failure on builder Full details are available at: https://lab.llvm.org/buildbot/#/builders/72/builds/778 Here is the relevant piece of the build log for the reference |
ram-NK
pushed a commit
to ram-NK/llvm-project
that referenced
this pull request
Oct 14, 2026
…vm#217587) CUDA's atomicAdd() family is defined in terms of PTX atom.add, whose denormal behavior is fixed by the hardware. Without any annotation the backend has to assume the function's denormal mode must be honored and expands these into CAS loops whenever the two disagree. Mark them with !atomic.ignore.denormal.mode so the native instruction is used. That covers the __nvvm_atom_*_add_gen_f builtins that atomicAdd(), atomicAdd_block() and atomicAdd_system() are written in terms of, plus C11/C++11 atomics under -fatomic-ignore-denormal-mode and the [[clang::atomic(ignore_denormal_mode)]] attribute, which requires teaching the NVPTX target about AtomicOptions. The condition for when the metadata is meaningful is now shared with the AMDGPU and SPIR-V targets in addAtomicIgnoreDenormalModeMetadata(). It takes an AllowHalf flag because whether f16 denormals are observable is target specific: PTX exposes no FTZ control for f16 operations, so atom.add.f16 never flushes and the opt-in is meaningful there, whereas the AMDGPU backend only consults the metadata for f32. --------- Co-authored-by: Artem Belevich <tra@google.com>
MazenEwiss
pushed a commit
to MazenEwiss/llvm-project
that referenced
this pull request
Oct 14, 2026
…vm#217587) CUDA's atomicAdd() family is defined in terms of PTX atom.add, whose denormal behavior is fixed by the hardware. Without any annotation the backend has to assume the function's denormal mode must be honored and expands these into CAS loops whenever the two disagree. Mark them with !atomic.ignore.denormal.mode so the native instruction is used. That covers the __nvvm_atom_*_add_gen_f builtins that atomicAdd(), atomicAdd_block() and atomicAdd_system() are written in terms of, plus C11/C++11 atomics under -fatomic-ignore-denormal-mode and the [[clang::atomic(ignore_denormal_mode)]] attribute, which requires teaching the NVPTX target about AtomicOptions. The condition for when the metadata is meaningful is now shared with the AMDGPU and SPIR-V targets in addAtomicIgnoreDenormalModeMetadata(). It takes an AllowHalf flag because whether f16 denormals are observable is target specific: PTX exposes no FTZ control for f16 operations, so atom.add.f16 never flushes and the opt-in is meaningful there, whereas the AMDGPU backend only consults the metadata for f32. --------- Co-authored-by: Artem Belevich <tra@google.com>
StarOne01
pushed a commit
to StarOne01/llvm-project
that referenced
this pull request
Oct 15, 2026
…vm#217587) CUDA's atomicAdd() family is defined in terms of PTX atom.add, whose denormal behavior is fixed by the hardware. Without any annotation the backend has to assume the function's denormal mode must be honored and expands these into CAS loops whenever the two disagree. Mark them with !atomic.ignore.denormal.mode so the native instruction is used. That covers the __nvvm_atom_*_add_gen_f builtins that atomicAdd(), atomicAdd_block() and atomicAdd_system() are written in terms of, plus C11/C++11 atomics under -fatomic-ignore-denormal-mode and the [[clang::atomic(ignore_denormal_mode)]] attribute, which requires teaching the NVPTX target about AtomicOptions. The condition for when the metadata is meaningful is now shared with the AMDGPU and SPIR-V targets in addAtomicIgnoreDenormalModeMetadata(). It takes an AllowHalf flag because whether f16 denormals are observable is target specific: PTX exposes no FTZ control for f16 operations, so atom.add.f16 never flushes and the opt-in is meaningful there, whereas the AMDGPU backend only consults the metadata for f32. --------- Co-authored-by: Artem Belevich <tra@google.com>
chsigg
added a commit
that referenced
this pull request
Oct 17, 2026
…M atomic add intrinsics (#232725) 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 PTX `atom.add` instruction, which flushes f32 denormals to zero on global memory. Attach `!atomic.ignore.denormal.mode` metadata when auto-upgrading these intrinsics to `atomicrmw fadd` so that they continue to lower to native `atom.add` instructions when `-nvptx-allow-ftz-atomics` defaults to `false` (#231429), matching the Clang lowering of `__nvvm_atom_add_gen_f` (#217587).
chsigg
added a commit
that referenced
this pull request
Oct 18, 2026
…232726) CUDA Fortran `atomicadd` has the same device semantics as CUDA C++ `atomicAdd()`, targeting PTX `atom.add` (which flushes f32 denormals to zero on global memory). Set `ignore_denormal_mode` on `mlir::LLVM::AtomicRMWOp` when lowering CUDA Fortran `atomicadd` for floating-point types in `genAtomBinOp`, matching Clang's CUDA/HIP `atomicAdd` lowering (#217587) so it continues to emit native `atom.add.f32` when `-nvptx-allow-ftz-atomics` defaults to `false` (#231429).
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.
CUDA's atomicAdd() family is defined in terms of PTX atom.add, whose
denormal behavior is fixed by the hardware. Without any annotation the
backend has to assume the function's denormal mode must be honored and
expands these into CAS loops whenever the two disagree. Mark them with
!atomic.ignore.denormal.mode so the native instruction is used.
That covers the _nvvm_atom*_add_gen_f builtins that atomicAdd(),
atomicAdd_block() and atomicAdd_system() are written in terms of, plus
C11/C++11 atomics under -fatomic-ignore-denormal-mode and the
[[clang::atomic(ignore_denormal_mode)]] attribute, which requires
teaching the NVPTX target about AtomicOptions.
The condition for when the metadata is meaningful is now shared with the
AMDGPU and SPIR-V targets in addAtomicIgnoreDenormalModeMetadata(). It
takes an AllowHalf flag because whether f16 denormals are observable is
target specific: PTX exposes no FTZ control for f16 operations, so
atom.add.f16 never flushes and the opt-in is meaningful there, whereas
the AMDGPU backend only consults the metadata for f32.