llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-clang
Author: Chinmay Ingle (chinmayingle)
<details>
<summary>Changes</summary>
[Clang][NVPTX] Emit device scope for generic atomic builtins
The unscoped __nvvm_atom_*_gen_* builtins (CUDA's atomicAdd() etc.)
were emitted without a syncscope, which LLVM treats as system scope.
Since #<!-- -->222458, the NVPTX backend rejects system scope atomics below
sm_60, so these builtins failed to compile for sm_50-53.
Emit syncscope("device") for them instead. On sm_60 and newer, plain
atomicAdd now lowers to a device-scope atom rather than .sys, which
matches nvcc.
- Tests: regenerated atomic-ignore-denormal-mode-nvptx.cu and added
syncscope("device") checks to builtins-nvptx.c.
Fixes #<!-- -->224420
---
Patch is 20.20 KiB, truncated to 20.00 KiB below, full version:
https://github.com/llvm/llvm-project/pull/233280.diff
5 Files Affected:
- (modified) clang/lib/CodeGen/CGBuiltin.cpp (+8-6)
- (modified) clang/lib/CodeGen/CGBuiltin.h (+12-9)
- (modified) clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp (+39-16)
- (modified) clang/test/CodeGen/builtins-nvptx.c (+36-36)
- (modified) clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu (+2-2)
``````````diff
diff --git a/clang/lib/CodeGen/CGBuiltin.cpp b/clang/lib/CodeGen/CGBuiltin.cpp
index c93375ff3a18ef..22d89df39ef7ad 100644
--- a/clang/lib/CodeGen/CGBuiltin.cpp
+++ b/clang/lib/CodeGen/CGBuiltin.cpp
@@ -272,9 +272,10 @@ Address CheckAtomicAlignment(CodeGenFunction &CGF, const
CallExpr *E) {
/// Utility to insert an atomic instruction based on Intrinsic::ID
/// and the expression node.
-Value *MakeBinaryAtomicValue(
- CodeGenFunction &CGF, llvm::AtomicRMWInst::BinOp Kind, const CallExpr *E,
- AtomicOrdering Ordering) {
+Value *MakeBinaryAtomicValue(CodeGenFunction &CGF,
+ llvm::AtomicRMWInst::BinOp Kind, const CallExpr
*E,
+ AtomicOrdering Ordering,
+ llvm::SyncScope::ID SSID) {
QualType T = E->getType();
assert(E->getArg(0)->getType()->isPointerType());
@@ -292,7 +293,7 @@ Value *MakeBinaryAtomicValue(
Val = EmitToInt(CGF, Val, T, IntType);
llvm::Value *Result =
- CGF.Builder.CreateAtomicRMW(Kind, DestAddr, Val, Ordering);
+ CGF.Builder.CreateAtomicRMW(Kind, DestAddr, Val, Ordering, SSID);
// Consider atomics to be volatile in MS kernel mode.
if (CGF.CGM.getLangOpts().Kernel)
cast<llvm::AtomicRMWInst>(Result)->setVolatile(true);
@@ -375,7 +376,8 @@ static RValue EmitBinaryAtomicPost(CodeGenFunction &CGF,
Value *MakeAtomicCmpXchgValue(CodeGenFunction &CGF, const CallExpr *E,
bool ReturnBool,
llvm::AtomicOrdering SuccessOrdering,
- llvm::AtomicOrdering FailureOrdering) {
+ llvm::AtomicOrdering FailureOrdering,
+ llvm::SyncScope::ID SSID) {
QualType T = ReturnBool ? E->getArg(1)->getType() : E->getType();
Address DestAddr = CheckAtomicAlignment(CGF, E);
@@ -388,7 +390,7 @@ Value *MakeAtomicCmpXchgValue(CodeGenFunction &CGF, const
CallExpr *E,
Value *New = EmitToInt(CGF, CGF.EmitScalarExpr(E->getArg(2)), T, IntType);
Value *Pair = CGF.Builder.CreateAtomicCmpXchg(
- DestAddr, Cmp, New, SuccessOrdering, FailureOrdering);
+ DestAddr, Cmp, New, SuccessOrdering, FailureOrdering, SSID);
if (ReturnBool)
// Extract boolean success flag and zext it to int.
return CGF.Builder.CreateZExt(CGF.Builder.CreateExtractValue(Pair, 1),
diff --git a/clang/lib/CodeGen/CGBuiltin.h b/clang/lib/CodeGen/CGBuiltin.h
index 394b397fde6a9f..e651dd9cc03a9b 100644
--- a/clang/lib/CodeGen/CGBuiltin.h
+++ b/clang/lib/CodeGen/CGBuiltin.h
@@ -91,11 +91,12 @@ llvm::Value *EmitFromInt(clang::CodeGen::CodeGenFunction
&CGF, llvm::Value *V,
clang::CodeGen::Address CheckAtomicAlignment(clang::CodeGen::CodeGenFunction
&CGF,
const clang::CallExpr *E);
-llvm::Value *MakeBinaryAtomicValue(clang::CodeGen::CodeGenFunction &CGF,
- llvm::AtomicRMWInst::BinOp Kind,
- const clang::CallExpr *E,
- llvm::AtomicOrdering Ordering =
-
llvm::AtomicOrdering::SequentiallyConsistent);
+llvm::Value *
+MakeBinaryAtomicValue(clang::CodeGen::CodeGenFunction &CGF,
+ llvm::AtomicRMWInst::BinOp Kind, const clang::CallExpr
*E,
+ llvm::AtomicOrdering Ordering =
+ llvm::AtomicOrdering::SequentiallyConsistent,
+ llvm::SyncScope::ID SSID = llvm::SyncScope::System);
llvm::Value *EmitOverflowIntrinsic(clang::CodeGen::CodeGenFunction &CGF,
const llvm::Intrinsic::ID IntrinsicID,
@@ -103,9 +104,11 @@ llvm::Value
*EmitOverflowIntrinsic(clang::CodeGen::CodeGenFunction &CGF,
llvm::Value *Y,
llvm::Value *&Carry);
-llvm::Value *MakeAtomicCmpXchgValue(clang::CodeGen::CodeGenFunction &CGF,
- const clang::CallExpr *E, bool ReturnBool,
- llvm::AtomicOrdering SuccessOrdering,
- llvm::AtomicOrdering FailureOrdering);
+llvm::Value *
+MakeAtomicCmpXchgValue(clang::CodeGen::CodeGenFunction &CGF,
+ const clang::CallExpr *E, bool ReturnBool,
+ llvm::AtomicOrdering SuccessOrdering,
+ llvm::AtomicOrdering FailureOrdering,
+ llvm::SyncScope::ID SSID = llvm::SyncScope::System);
#endif
diff --git a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
index a513fcc12953ba..cc230f451e8125 100644
--- a/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/NVPTX.cpp
@@ -451,6 +451,12 @@ static Value *MakeFPArith(unsigned IntrinsicID,
APFloat::roundingMode RM,
} // namespace
+static llvm::SyncScope::ID getDeviceSyncScope(CodeGenFunction &CGF) {
+ return CGF.getLLVMContext().getOrInsertSyncScopeID(
+ *llvm::getAtomicScopeIRString(CGF.getTarget().getTriple(),
+ llvm::AtomicScope::Device));
+}
+
Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned BuiltinID,
const CallExpr *E) {
switch (BuiltinID) {
@@ -462,65 +468,79 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned
BuiltinID,
E->getType()->hasSignedIntegerRepresentation() ? Intrinsic::smulh
: Intrinsic::umulh,
EmitScalarExpr(E->getArg(0)), EmitScalarExpr(E->getArg(1)));
+
+ // Unscoped __nvvm_atom_*_gen_* builtins (CUDA's atomicAdd() etc.) have
device
+ // scope, but LLVM's default sync scope is system, so request device
+ // explicitly.
case NVPTX::BI__nvvm_atom_add_gen_i:
case NVPTX::BI__nvvm_atom_add_gen_l:
case NVPTX::BI__nvvm_atom_add_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Add, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_sub_gen_i:
case NVPTX::BI__nvvm_atom_sub_gen_l:
case NVPTX::BI__nvvm_atom_sub_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Sub, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_and_gen_i:
case NVPTX::BI__nvvm_atom_and_gen_l:
case NVPTX::BI__nvvm_atom_and_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::And, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_or_gen_i:
case NVPTX::BI__nvvm_atom_or_gen_l:
case NVPTX::BI__nvvm_atom_or_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Or, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_xor_gen_i:
case NVPTX::BI__nvvm_atom_xor_gen_l:
case NVPTX::BI__nvvm_atom_xor_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Xor, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_xchg_gen_i:
case NVPTX::BI__nvvm_atom_xchg_gen_l:
case NVPTX::BI__nvvm_atom_xchg_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Xchg, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_max_gen_i:
case NVPTX::BI__nvvm_atom_max_gen_l:
case NVPTX::BI__nvvm_atom_max_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Max, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_max_gen_ui:
case NVPTX::BI__nvvm_atom_max_gen_ul:
case NVPTX::BI__nvvm_atom_max_gen_ull:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::UMax, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_min_gen_i:
case NVPTX::BI__nvvm_atom_min_gen_l:
case NVPTX::BI__nvvm_atom_min_gen_ll:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::Min, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_min_gen_ui:
case NVPTX::BI__nvvm_atom_min_gen_ul:
case NVPTX::BI__nvvm_atom_min_gen_ull:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::UMin, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_cas_gen_us:
case NVPTX::BI__nvvm_atom_cas_gen_i:
@@ -528,9 +548,9 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned
BuiltinID,
case NVPTX::BI__nvvm_atom_cas_gen_ll:
// __nvvm_atom_cas_gen_* should return the old value rather than the
// success flag.
- return MakeAtomicCmpXchgValue(*this, E, /*ReturnBool=*/false,
- AtomicOrdering::Monotonic,
- AtomicOrdering::Monotonic);
+ return MakeAtomicCmpXchgValue(
+ *this, E, /*ReturnBool=*/false, AtomicOrdering::Monotonic,
+ AtomicOrdering::Monotonic, getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_add_gen_f:
case NVPTX::BI__nvvm_atom_add_gen_d: {
@@ -538,7 +558,8 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned
BuiltinID,
Value *Val = EmitScalarExpr(E->getArg(1));
llvm::AtomicRMWInst *RMW = Builder.CreateAtomicRMW(
- llvm::AtomicRMWInst::FAdd, DestAddr, Val, AtomicOrdering::Monotonic);
+ llvm::AtomicRMWInst::FAdd, DestAddr, Val, AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
// CUDA's atomicAdd() is defined in terms of the native atom.add
// instruction, whose denormal behavior is fixed by the hardware and need
@@ -550,11 +571,13 @@ Value *CodeGenFunction::EmitNVPTXBuiltinExpr(unsigned
BuiltinID,
case NVPTX::BI__nvvm_atom_inc_gen_ui:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::UIncWrap, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_atom_dec_gen_ui:
return MakeBinaryAtomicValue(*this, llvm::AtomicRMWInst::UDecWrap, E,
- AtomicOrdering::Monotonic);
+ AtomicOrdering::Monotonic,
+ getDeviceSyncScope(*this));
case NVPTX::BI__nvvm_ldg_c:
case NVPTX::BI__nvvm_ldg_sc:
diff --git a/clang/test/CodeGen/builtins-nvptx.c
b/clang/test/CodeGen/builtins-nvptx.c
index b7126a82b8064b..aad71845aa0790 100644
--- a/clang/test/CodeGen/builtins-nvptx.c
+++ b/clang/test/CodeGen/builtins-nvptx.c
@@ -339,94 +339,94 @@ __device__ void nvvm_atom(float *fp, float f, double
*dfp, double df,
unsigned short *usp, unsigned short us, int *ip,
int i, unsigned int *uip, unsigned ui, long *lp,
long l, long long *llp, long long ll) {
- // CHECK: atomicrmw add ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw add ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_add_gen_i(ip, i);
- // CHECK: atomicrmw add ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw add ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_add_gen_l(&dl, l);
- // CHECK: atomicrmw add ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw add ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_add_gen_ll(&sll, ll);
- // CHECK: atomicrmw sub ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw sub ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_sub_gen_i(ip, i);
- // CHECK: atomicrmw sub ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw sub ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_sub_gen_l(&dl, l);
- // CHECK: atomicrmw sub ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw sub ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_sub_gen_ll(&sll, ll);
- // CHECK: atomicrmw and ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw and ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_and_gen_i(ip, i);
- // CHECK: atomicrmw and ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw and ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_and_gen_l(&dl, l);
- // CHECK: atomicrmw and ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw and ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_and_gen_ll(&sll, ll);
- // CHECK: atomicrmw or ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw or ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_or_gen_i(ip, i);
- // CHECK: atomicrmw or ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw or ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_or_gen_l(&dl, l);
- // CHECK: atomicrmw or ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw or ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_or_gen_ll(&sll, ll);
- // CHECK: atomicrmw xor ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw xor ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_xor_gen_i(ip, i);
- // CHECK: atomicrmw xor ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw xor ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_xor_gen_l(&dl, l);
- // CHECK: atomicrmw xor ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw xor ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_xor_gen_ll(&sll, ll);
- // CHECK: atomicrmw xchg ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw xchg ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_xchg_gen_i(ip, i);
- // CHECK: atomicrmw xchg ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw xchg ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_xchg_gen_l(&dl, l);
- // CHECK: atomicrmw xchg ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw xchg ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_xchg_gen_ll(&sll, ll);
- // CHECK: atomicrmw max ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw max ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_max_gen_i(ip, i);
- // CHECK: atomicrmw umax ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw umax ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_max_gen_ui((unsigned int *)ip, i);
- // CHECK: atomicrmw max ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw max ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_max_gen_l(&dl, l);
- // CHECK: atomicrmw umax ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw umax ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_max_gen_ul((unsigned long *)&dl, l);
- // CHECK: atomicrmw max ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw max ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_max_gen_ll(&sll, ll);
- // CHECK: atomicrmw umax ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw umax ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_max_gen_ull((unsigned long long *)&sll, ll);
- // CHECK: atomicrmw min ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw min ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_min_gen_i(ip, i);
- // CHECK: atomicrmw umin ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw umin ptr {{.*}} syncscope("device") monotonic, align 4
__nvvm_atom_min_gen_ui((unsigned int *)ip, i);
- // CHECK: atomicrmw min ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw min ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_min_gen_l(&dl, l);
- // CHECK: atomicrmw umin ptr {{.*}} monotonic, align {{4|8}}
+ // CHECK: atomicrmw umin ptr {{.*}} syncscope("device") monotonic, align
{{4|8}}
__nvvm_atom_min_gen_ul((unsigned long *)&dl, l);
- // CHECK: atomicrmw min ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw min ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_min_gen_ll(&sll, ll);
- // CHECK: atomicrmw umin ptr {{.*}} monotonic, align 8
+ // CHECK: atomicrmw umin ptr {{.*}} syncscope("device") monotonic, align 8
__nvvm_atom_min_gen_ull((unsigned long long *)&sll, ll);
- // CHECK: cmpxchg ptr {{.*}} monotonic monotonic, align 4
+ // CHECK: cmpxchg ptr {{.*}} syncscope("device") monotonic monotonic, align 4
// CHECK-NEXT: extractvalue { i32, i1 } {{%[0-9]+}}, 0
__nvvm_atom_cas_gen_i(ip, 0, i);
- // CHECK: cmpxchg ptr {{.*}} monotonic monotonic, align {{4|8}}
+ // CHECK: cmpxchg ptr {{.*}} syncscope("device") monotonic monotonic, align
{{4|8}}
// CHECK-NEXT: extractvalue { {{i32|i64}}, i1 } {{%[0-9]+}}, 0
__nvvm_atom_cas_gen_l(&dl, 0, l);
- // CHECK: cmpxchg ptr {{.*}} monotonic monotonic, align 8
+ // CHECK: cmpxchg ptr {{.*}} syncscope("device") monotonic monotonic, align 8
// CHECK-NEXT: extractvalue { i64, i1 } {{%[0-9]+}}, 0
__nvvm_atom_cas_gen_ll(&sll, 0, ll);
// 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
+ // CHECK: atomicrmw fadd ptr {{.*}} syncscope("device") monotonic, align 4,
!atomic.ignore.denormal.mode
__nvvm_atom_add_gen_f(fp, f);
- // CHECK: atomicrmw uinc_wrap ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw uinc_wrap ptr {{.*}} syncscope("device") monotonic,
align 4
__nvvm_atom_inc_gen_ui(uip, ui);
- // CHECK: atomicrmw udec_wrap ptr {{.*}} monotonic, align 4
+ // CHECK: atomicrmw udec_wrap ptr {{.*}} syncscope("device") monotonic,
align 4
__nvvm_atom_dec_gen_ui(uip, ui);
diff --git a/clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu
b/clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu
index e3ae34322a83d9..ba5eecfd37580e 100644
--- a/clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu
+++ b/clang/test/CodeGenCUDA/atomic-ignore-denormal-mode-nvptx.cu
@@ -169,7 +169,7 @@ __device__ void test_attr_off(float *a) {
// 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: [[TMP2:%.*]] = atomicrmw fadd ptr [[TMP0]], float [[TMP1]]
syncscope("device") monotonic, align 4, !atomic.ignore.denormal.mode [[META2]]
// DEV-NEXT: ret float [[TMP2]]
//
// OPT-LABEL: define dso_local noundef float @_Z23test_atomic_add_builtinPff(
@@ -181,7 +181,7 @@ __device__ void test_attr_off(float *a) {
// 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: [[TMP2:%.*]] = atomicrmw fadd ptr [...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/233280
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
