[llvm-branch-commits] [clang] [clang][NVPTX] Emit !atomic.ignore.denormal.mode for CUDA atomics (PR #217587)
Christian Sigg via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Mon Aug 24 05:27:36 PDT 2026
https://github.com/chsigg updated https://github.com/llvm/llvm-project/pull/217587
>From 1c1276b561fb36140557f032e8c574132c6da76c Mon Sep 17 00:00:00 2001
From: Christian Sigg <csigg at google.com>
Date: Wed, 19 Aug 2026 14:50:42 +0200
Subject: [PATCH] [clang][NVPTX] Emit !atomic.ignore.denormal.mode for CUDA
atomics
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 at google.com>
---
clang/lib/Basic/Targets/NVPTX.cpp | 6 +
clang/lib/Basic/Targets/NVPTX.h | 2 +
clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 5 +-
clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp | 23 +-
clang/lib/CodeGen/TargetInfo.cpp | 15 +
clang/lib/CodeGen/TargetInfo.h | 18 +
clang/lib/CodeGen/Targets/AMDGPU.cpp | 6 +-
clang/lib/CodeGen/Targets/NVPTX.cpp | 12 +
clang/lib/CodeGen/Targets/SPIR.cpp | 6 +-
clang/test/CodeGen/builtins-nvptx.c | 5 +-
.../atomic-ignore-denormal-mode-nvptx.cu | 321 ++++++++++++++++++
11 files changed, 402 insertions(+), 17 deletions(-)
create mode 100644 clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu
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 9a951eee44f14..0f0aa9f913c7e 100644
--- a/clang/lib/Basic/Targets/NVPTX.h
+++ b/clang/lib/Basic/Targets/NVPTX.h
@@ -144,6 +144,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 a4eb7ec126583..d6f53c6fdf10c 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"
@@ -2053,9 +2052,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 64fdae9d8934d..28394655537d1 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
@@ -507,8 +515,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 0b3bd5a4aab7f..8e09102c542d5 100644
--- a/clang/lib/CodeGen/Targets/SPIR.cpp
+++ b/clang/lib/CodeGen/Targets/SPIR.cpp
@@ -587,10 +587,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 87be7b46aad8e..f4e7b18bca292 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -387,7 +387,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
+// what atomicAdd() is defined to be, independent of the command line.
+
+// DEV-LABEL: define dso_local noundef float @_Z23test_atomic_add_builtinPff(
+// DEV-SAME: ptr noundef [[A:%.*]], float noundef [[V:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[V_ADDR:%.*]] = alloca float, align 4
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: store float [[V]], ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
+// DEV-NEXT: ret float [[TMP2]]
+//
+// OPT-LABEL: define dso_local noundef float @_Z23test_atomic_add_builtinPff(
+// OPT-SAME: ptr noundef [[A:%.*]], float noundef [[V:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[V_ADDR:%.*]] = alloca float, align 4
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: store float [[V]], ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
+// OPT-NEXT: ret float [[TMP2]]
+//
+__device__ float test_atomic_add_builtin(float *a, float v) {
+ return __nvvm_atom_add_gen_f(a, v);
+}
+
+// CUDA's atomicAdd_block()/atomicAdd_system() go through the scoped add
+// builtins, which are the same native atom.add and are marked the same way.
+
+// DEV-LABEL: define dso_local noundef float @_Z29test_atomic_add_block_builtinPff(
+// DEV-SAME: ptr noundef [[A:%.*]], float noundef [[V:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[V_ADDR:%.*]] = alloca float, align 4
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: store float [[V]], ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] syncscope("block") monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
+// DEV-NEXT: ret float [[TMP2]]
+//
+// OPT-LABEL: define dso_local noundef float @_Z29test_atomic_add_block_builtinPff(
+// OPT-SAME: ptr noundef [[A:%.*]], float noundef [[V:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[V_ADDR:%.*]] = alloca float, align 4
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: store float [[V]], ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] syncscope("block") monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
+// OPT-NEXT: ret float [[TMP2]]
+//
+__device__ float test_atomic_add_block_builtin(float *a, float v) {
+ return __nvvm_atom_cta_add_gen_f(a, v);
+}
+
+// DEV-LABEL: define dso_local noundef float @_Z30test_atomic_add_system_builtinPff(
+// DEV-SAME: ptr noundef [[A:%.*]], float noundef [[V:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[V_ADDR:%.*]] = alloca float, align 4
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: store float [[V]], ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP1:%.*]] = load float, ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
+// DEV-NEXT: ret float [[TMP2]]
+//
+// OPT-LABEL: define dso_local noundef float @_Z30test_atomic_add_system_builtinPff(
+// OPT-SAME: ptr noundef [[A:%.*]], float noundef [[V:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[V_ADDR:%.*]] = alloca float, align 4
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: store float [[V]], ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP1:%.*]] = load float, ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]] monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
+// OPT-NEXT: ret float [[TMP2]]
+//
+__device__ float test_atomic_add_system_builtin(float *a, float v) {
+ return __nvvm_atom_sys_add_gen_f(a, v);
+}
+
+// Scoped operations that are not a floating-point add stay unmarked.
+
+// DEV-LABEL: define dso_local noundef i32 @_Z25test_atomic_add_block_intPii(
+// DEV-SAME: ptr noundef [[A:%.*]], i32 noundef [[V:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[V_ADDR:%.*]] = alloca i32, align 4
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: store i32 [[V]], ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP1:%.*]] = load i32, ptr [[V_ADDR]], align 4
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("block") monotonic, align 4
+// DEV-NEXT: ret i32 [[TMP2]]
+//
+// OPT-LABEL: define dso_local noundef i32 @_Z25test_atomic_add_block_intPii(
+// OPT-SAME: ptr noundef [[A:%.*]], i32 noundef [[V:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[V_ADDR:%.*]] = alloca i32, align 4
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: store i32 [[V]], ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP1:%.*]] = load i32, ptr [[V_ADDR]], align 4
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw add ptr [[TMP0]], i32 [[TMP1]] syncscope("block") monotonic, align 4
+// OPT-NEXT: ret i32 [[TMP2]]
+//
+__device__ int test_atomic_add_block_int(int *a, int v) {
+ return __nvvm_atom_cta_add_gen_i(a, v);
+}
+
+// PTX has no FTZ control for f16 operations, so atom.add.f16 never flushes and
+// the per-instruction opt-in is meaningful for _Float16 as well.
+
+// DEV-LABEL: define dso_local void @_Z9test_halfPDF16_(
+// DEV-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// DEV-NEXT: [[ENTRY:.*:]]
+// DEV-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// DEV-NEXT: [[DOTATOMICTMP:%.*]] = alloca half, align 2
+// DEV-NEXT: [[ATOMIC_TEMP:%.*]] = alloca half, align 2
+// DEV-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// DEV-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// DEV-NEXT: store half 1.000000e+00, ptr [[DOTATOMICTMP]], align 2
+// DEV-NEXT: [[TMP1:%.*]] = load half, ptr [[DOTATOMICTMP]], align 2
+// DEV-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], half [[TMP1]] monotonic, align 2
+// DEV-NEXT: store half [[TMP2]], ptr [[ATOMIC_TEMP]], align 2
+// DEV-NEXT: [[TMP3:%.*]] = load half, ptr [[ATOMIC_TEMP]], align 2
+// DEV-NEXT: ret void
+//
+// OPT-LABEL: define dso_local void @_Z9test_halfPDF16_(
+// OPT-SAME: ptr noundef [[A:%.*]]) #[[ATTR0]] {
+// OPT-NEXT: [[ENTRY:.*:]]
+// OPT-NEXT: [[A_ADDR:%.*]] = alloca ptr, align 8
+// OPT-NEXT: [[DOTATOMICTMP:%.*]] = alloca half, align 2
+// OPT-NEXT: [[ATOMIC_TEMP:%.*]] = alloca half, align 2
+// OPT-NEXT: store ptr [[A]], ptr [[A_ADDR]], align 8
+// OPT-NEXT: [[TMP0:%.*]] = load ptr, ptr [[A_ADDR]], align 8
+// OPT-NEXT: store half 1.000000e+00, ptr [[DOTATOMICTMP]], align 2
+// OPT-NEXT: [[TMP1:%.*]] = load half, ptr [[DOTATOMICTMP]], align 2
+// OPT-NEXT: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], half [[TMP1]] monotonic, align 2, !atomic.ignore.denormal.mode [[META2]]
+// OPT-NEXT: store half [[TMP2]], ptr [[ATOMIC_TEMP]], align 2
+// OPT-NEXT: [[TMP3:%.*]] = load half, ptr [[ATOMIC_TEMP]], align 2
+// OPT-NEXT: ret void
+//
+__device__ void test_half(_Float16 *a) {
+ __scoped_atomic_fetch_add(a, (_Float16)1.0, __ATOMIC_RELAXED,
+ __MEMORY_SCOPE_SYSTEM);
+}
+//.
+// DEV: [[META2]] = !{}
+//.
+// OPT: [[META2]] = !{}
+//.
More information about the llvm-branch-commits
mailing list