@@ -12,11 +12,11 @@ SYCL_EXTERNAL void intAtomicFunc(int *i) {
1212 atomicInt.fetch_and (1 );
1313 atomicInt.fetch_or (1 );
1414 // CHECK: void{{.*}}intAtomicFunc
15- // CHECK-SAFE: cmpxchg volatile
15+ // CHECK-SAFE: cmpxchg
1616 // CHECK-SAFE-NOT: atomicrmw
17- // CHECK-UNSAFE: atomicrmw volatile xor
18- // CHECK-UNSAFE: atomicrmw volatile and
19- // CHECK-UNSAFE: atomicrmw volatile or
17+ // CHECK-UNSAFE: atomicrmw xor
18+ // CHECK-UNSAFE: atomicrmw and
19+ // CHECK-UNSAFE: atomicrmw or
2020 // CHECK-UNSAFE-NOT: cmpxchg
2121}
2222
@@ -25,15 +25,15 @@ SYCL_EXTERNAL void fpAtomicFunc(float *f, double *d) {
2525 sycl::access::address_space::global_space>(*f)
2626 .fetch_add (1 .0f );
2727 // CHECK: void{{.*}}fpAtomicFunc
28- // CHECK-SAFE: atomicrmw volatile fadd
28+ // CHECK-SAFE: atomicrmw fadd
2929 // CHECK-SAFE-NOT: amdgpu.ignore.denormal.mode
30- // CHECK-UNSAFE-FP: atomicrmw volatile fadd {{.*}}!amdgpu.no.fine.grained.memory{{.*}}!amdgpu.ignore.denormal.mode
30+ // CHECK-UNSAFE-FP: atomicrmw fadd {{.*}}!amdgpu.no.fine.grained.memory{{.*}}!amdgpu.ignore.denormal.mode
3131 sycl::atomic_ref<double , sycl::memory_order_relaxed,
3232 sycl::memory_scope_device,
3333 sycl::access::address_space::global_space>(*d)
3434 .fetch_add (1.0 );
3535 // CHECK-SAFE: cmpxchg
3636 // CHECK-SAFE-NOT: llvm.amdgcn.global.atomic.fadd.f64
37- // CHECK-UNSAFE-FP: atomicrmw volatile fadd {{.*}}!amdgpu.no.fine.grained.memory
37+ // CHECK-UNSAFE-FP: atomicrmw fadd {{.*}}!amdgpu.no.fine.grained.memory
3838 // CHECK: __CLANG_OFFLOAD_BUNDLE____END__ sycl-amdgcn-amd-amdhsa-
3939}
0 commit comments