Skip to content

Commit 0027fd3

Browse files
jhuber6jhuber6
committed
[Clang][AMDGPU] Stop defaulting to one-as for all atomic scopes
Summary: The documentation at https://llvm.org/docs/AMDGPUUsage.html#memory-scopes states that these 'one-as' modifiers are more specific versions of the scopes that only apply to a specific address space. This doesn't make sense for fences which have no associated address space to use, and it's a more restrictive version the normal scope. This should not tbe the default behavior, but it is currently emitted in all cases except for sequentially consistent.
1 parent 2067e60 commit 0027fd3

File tree

8 files changed

+2519
-376
lines changed

8 files changed

+2519
-376
lines changed

‎clang/lib/CodeGen/Targets/AMDGPU.cpp‎

Lines changed: 0 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -537,13 +537,6 @@ AMDGPUTargetCodeGenInfo::getLLVMSyncScopeID(const LangOptions &LangOpts,
537537
break;
538538
}
539539

540-
if (Ordering != llvm::AtomicOrdering::SequentiallyConsistent) {
541-
if (!Name.empty())
542-
Name = Twine(Twine(Name) + Twine("-")).str();
543-
544-
Name = Twine(Twine(Name) + Twine("one-as")).str();
545-
}
546-
547540
return Ctx.getOrInsertSyncScopeID(Name);
548541
}
549542

‎clang/test/CodeGen/scoped-atomic-ops.c‎

Lines changed: 2311 additions & 161 deletions
Large diffs are not rendered by default.

‎clang/test/CodeGen/scoped-fence-ops.c‎

Lines changed: 10 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@
99
// AMDGCN-LABEL: define hidden void @fe1a(
1010
// AMDGCN-SAME: ) #[[ATTR0:[0-9]+]] {
1111
// AMDGCN-NEXT: [[ENTRY:.*:]]
12-
// AMDGCN-NEXT: fence syncscope("workgroup-one-as") release
12+
// AMDGCN-NEXT: fence syncscope("workgroup") release
1313
// AMDGCN-NEXT: ret void
1414
//
1515
// SPIRV-LABEL: define hidden spir_func void @fe1a(
@@ -45,13 +45,13 @@ void fe1a() {
4545
// AMDGCN: [[ATOMIC_SCOPE_CONTINUE]]:
4646
// AMDGCN-NEXT: ret void
4747
// AMDGCN: [[ACQUIRE]]:
48-
// AMDGCN-NEXT: fence syncscope("workgroup-one-as") acquire
48+
// AMDGCN-NEXT: fence syncscope("workgroup") acquire
4949
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
5050
// AMDGCN: [[RELEASE]]:
51-
// AMDGCN-NEXT: fence syncscope("workgroup-one-as") release
51+
// AMDGCN-NEXT: fence syncscope("workgroup") release
5252
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
5353
// AMDGCN: [[ACQREL]]:
54-
// AMDGCN-NEXT: fence syncscope("workgroup-one-as") acq_rel
54+
// AMDGCN-NEXT: fence syncscope("workgroup") acq_rel
5555
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
5656
// AMDGCN: [[SEQCST]]:
5757
// AMDGCN-NEXT: fence syncscope("workgroup") seq_cst
@@ -134,19 +134,19 @@ void fe1b(int ord) {
134134
// AMDGCN: [[ATOMIC_SCOPE_CONTINUE]]:
135135
// AMDGCN-NEXT: ret void
136136
// AMDGCN: [[DEVICE_SCOPE]]:
137-
// AMDGCN-NEXT: fence syncscope("agent-one-as") release
137+
// AMDGCN-NEXT: fence syncscope("agent") release
138138
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
139139
// AMDGCN: [[SYSTEM_SCOPE]]:
140-
// AMDGCN-NEXT: fence syncscope("one-as") release
140+
// AMDGCN-NEXT: fence release
141141
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
142142
// AMDGCN: [[WORKGROUP_SCOPE]]:
143-
// AMDGCN-NEXT: fence syncscope("workgroup-one-as") release
143+
// AMDGCN-NEXT: fence syncscope("workgroup") release
144144
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
145145
// AMDGCN: [[WAVEFRONT_SCOPE]]:
146-
// AMDGCN-NEXT: fence syncscope("wavefront-one-as") release
146+
// AMDGCN-NEXT: fence syncscope("wavefront") release
147147
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
148148
// AMDGCN: [[SINGLE_SCOPE]]:
149-
// AMDGCN-NEXT: fence syncscope("singlethread-one-as") release
149+
// AMDGCN-NEXT: fence syncscope("singlethread") release
150150
// AMDGCN-NEXT: br label %[[ATOMIC_SCOPE_CONTINUE]]
151151
//
152152
// SPIRV-LABEL: define hidden spir_func void @fe1c(
@@ -237,7 +237,7 @@ void fe2a() {
237237
// AMDGCN-LABEL: define hidden void @fe2b(
238238
// AMDGCN-SAME: ) #[[ATTR0]] {
239239
// AMDGCN-NEXT: [[ENTRY:.*:]]
240-
// AMDGCN-NEXT: fence syncscope("one-as") release
240+
// AMDGCN-NEXT: fence release
241241
// AMDGCN-NEXT: ret void
242242
//
243243
// SPIRV-LABEL: define hidden spir_func void @fe2b(

‎clang/test/CodeGenCUDA/amdgpu-atomic-ops.cu‎

Lines changed: 30 additions & 30 deletions
Original file line numberDiff line numberDiff line change
@@ -26,19 +26,19 @@ __global__ void ffp1(float *p) {
2626
// SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4{{$}}
2727
// SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4{{$}}
2828
// SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4{{$}}
29-
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent-one-as") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE:[0-9]+]]{{$}}
30-
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
31-
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
32-
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
29+
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE:[0-9]+]]{{$}}
30+
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
31+
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
32+
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
3333

3434
// UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.ignore.denormal.mode !{{[0-9]+$}}
3535
// UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
3636
// UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
3737
// UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 4, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
3838
// UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE:[0-9]+]], !amdgpu.no.fine.grained.memory !{{[0-9]+}}, !amdgpu.ignore.denormal.mode !{{[0-9]+$}}
3939
// UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
40-
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
41-
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
40+
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
41+
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
4242

4343
// SAFE: _Z4ffp1Pf
4444
// SAFE: global_atomic_cmpswap
@@ -73,19 +73,19 @@ __global__ void ffp2(double *p) {
7373
// SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8{{$}}
7474
// SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8{{$}}
7575
// SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8{{$}}
76-
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
77-
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
78-
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
79-
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
76+
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
77+
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
78+
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
79+
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
8080

8181
// UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
8282
// UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
8383
// UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
8484
// UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
85-
// UNSAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
86-
// UNSAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
87-
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
88-
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
85+
// UNSAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
86+
// UNSAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
87+
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
88+
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
8989

9090
// SAFE-LABEL: @_Z4ffp2Pd
9191
// SAFE: global_atomic_cmpswap_b64
@@ -119,19 +119,19 @@ __global__ void ffp3(long double *p) {
119119
// SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8{{$}}
120120
// SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8{{$}}
121121
// SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8{{$}}
122-
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
123-
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
124-
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
125-
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
122+
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
123+
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
124+
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
125+
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
126126

127127
// UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
128128
// UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
129129
// UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
130130
// UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 8, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
131-
// UNSAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
132-
// UNSAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
133-
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
134-
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
131+
// UNSAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
132+
// UNSAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
133+
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
134+
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 8, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
135135

136136
// SAFE-LABEL: @_Z4ffp3Pe
137137
// SAFE: global_atomic_cmpswap_b64
@@ -185,19 +185,19 @@ __global__ void ffp6(_Float16 *p) {
185185
// SAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2{{$}}
186186
// SAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2{{$}}
187187
// SAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2{{$}}
188-
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent-one-as") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
189-
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
190-
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
191-
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
188+
// SAFEIR: atomicrmw fadd ptr {{.*}} syncscope("agent") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
189+
// SAFEIR: atomicrmw fsub ptr {{.*}} syncscope("workgroup") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
190+
// SAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
191+
// SAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
192192

193193
// UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
194194
// UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
195195
// UNSAFEIR: atomicrmw fmax ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
196196
// UNSAFEIR: atomicrmw fmin ptr {{.*}} monotonic, align 2, !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
197197
// UNSAFEIR: atomicrmw fadd ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
198198
// UNSAFEIR: atomicrmw fsub ptr {{.*}} monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
199-
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent-one-as") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
200-
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup-one-as") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
199+
// UNSAFEIR: atomicrmw fmax ptr {{.*}} syncscope("agent") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
200+
// UNSAFEIR: atomicrmw fmin ptr {{.*}} syncscope("workgroup") monotonic, align 2, !noalias.addrspace ![[$NO_PRIVATE]], !amdgpu.no.fine.grained.memory !{{[0-9]+$}}
201201

202202
// SAFE: _Z4ffp6PDF16
203203
// SAFE: global_atomic_cmpswap
@@ -228,8 +228,8 @@ __global__ void ffp6(_Float16 *p) {
228228
// CHECK-LABEL: @_Z12test_cmpxchgPiii
229229
// CHECK: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} acquire acquire, align 4{{$}}
230230
// CHECK: cmpxchg weak ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} acquire acquire, align 4{{$}}
231-
// CHECK: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("workgroup-one-as") monotonic monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
232-
// CHECK: cmpxchg weak ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("workgroup-one-as") monotonic monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
231+
// CHECK: cmpxchg ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("workgroup") monotonic monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
232+
// CHECK: cmpxchg weak ptr %{{.+}}, i32 %{{.+}}, i32 %{{.+}} syncscope("workgroup") monotonic monotonic, align 4, !noalias.addrspace ![[$NO_PRIVATE]]{{$}}
233233
__device__ int test_cmpxchg(int *ptr, int cmp, int desired) {
234234
bool flag = __atomic_compare_exchange(ptr, &cmp, &desired, 0, memory_order_acquire, memory_order_acquire);
235235
flag = __atomic_compare_exchange_n(ptr, &cmp, desired, 1, memory_order_acquire, memory_order_acquire);

0 commit comments

Comments
 (0)