diff --git a/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp b/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp index d8cefd44003f0..b7f7e1e26830e 100644 --- a/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp +++ b/sycl/test/check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp @@ -1,5 +1,4 @@ // REQUIRES: hip -// XFAIL: hip // RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -S -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-SAFE // RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx906 %s -mllvm --amdgpu-oclc-unsafe-int-atomics=true -S -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-UNSAFE // RUN: %clangxx -fsycl -fsycl-targets=amd_gpu_gfx90a %s -mllvm --amdgpu-oclc-unsafe-fp-atomics=true -mllvm --amdgpu-oclc-unsafe-int-atomics=true -S -emit-llvm -o - | FileCheck %s --check-prefixes=CHECK,CHECK-UNSAFE-FP @@ -27,16 +26,14 @@ SYCL_EXTERNAL void fpAtomicFunc(float *f, double *d) { .fetch_add(1.0f); // CHECK: void{{.*}}fpAtomicFunc // CHECK-SAFE: atomicrmw volatile fadd - // CHECK-SAFE-NOT: llvm.amdgcn.global.atomic.fadd.f32 - // CHECK-UNSAFE-FP: llvm.amdgcn.global.atomic.fadd.f32 - // CHECK-UNSAFE-FP-NOT: atomicrmw volatile fadd + // CHECK-SAFE-NOT: amdgpu.ignore.denormal.mode + // CHECK-UNSAFE-FP: atomicrmw volatile fadd {{.*}}!amdgpu.no.fine.grained.memory{{.*}}!amdgpu.ignore.denormal.mode sycl::atomic_ref(*d) .fetch_add(1.0); // CHECK-SAFE: cmpxchg // CHECK-SAFE-NOT: llvm.amdgcn.global.atomic.fadd.f64 - // CHECK-UNSAFE-FP: llvm.amdgcn.global.atomic.fadd.f64 - // CHECK-UNSAFE-FP-NOT: cmpxchg + // CHECK-UNSAFE-FP: atomicrmw volatile fadd {{.*}}!amdgpu.no.fine.grained.memory // CHECK: __CLANG_OFFLOAD_BUNDLE____END__ sycl-amdgcn-amd-amdhsa- }