From 8cc1d47a9fd9c4eb5cb59c3e732e16510eb90060 Mon Sep 17 00:00:00 2001
From: Jakub Chlanda <jakub@codeplay.com>
Date: Thu, 12 Dec 2024 10:40:25 +0000
Subject: [PATCH 1/2] [SYCL] Update checks in unsafe atomics test for AMDGPU

---
 .../hip/atomic/amdgpu_unsafe_atomics.cpp                 | 9 +++------
 1 file changed, 3 insertions(+), 6 deletions(-)

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..b51227f9aed8d 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<double, sycl::memory_order_relaxed,
                    sycl::memory_scope_device,
                    sycl::access::address_space::global_space>(*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-
 }

From 7771d32fcecbd5457b9ba7e940144ae941643bbd Mon Sep 17 00:00:00 2001
From: Jakub Chlanda <jakub@codeplay.com>
Date: Thu, 12 Dec 2024 11:12:53 +0000
Subject: [PATCH 2/2] !

---
 .../check_device_code/hip/atomic/amdgpu_unsafe_atomics.cpp    | 4 ++--
 1 file changed, 2 insertions(+), 2 deletions(-)

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 b51227f9aed8d..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
@@ -27,13 +27,13 @@ SYCL_EXTERNAL void fpAtomicFunc(float *f, double *d) {
   // CHECK: void{{.*}}fpAtomicFunc
   // CHECK-SAFE: 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
+  // CHECK-UNSAFE-FP: atomicrmw volatile fadd {{.*}}!amdgpu.no.fine.grained.memory{{.*}}!amdgpu.ignore.denormal.mode
   sycl::atomic_ref<double, sycl::memory_order_relaxed,
                    sycl::memory_scope_device,
                    sycl::access::address_space::global_space>(*d)
       .fetch_add(1.0);
   // CHECK-SAFE: cmpxchg
   // CHECK-SAFE-NOT: llvm.amdgcn.global.atomic.fadd.f64
-  // CHECK-UNSAFE-FP: atomicrmw volatile fadd {{.*}}amdgpu.no.fine.grained.memory
+  // CHECK-UNSAFE-FP: atomicrmw volatile fadd {{.*}}!amdgpu.no.fine.grained.memory
   // CHECK: __CLANG_OFFLOAD_BUNDLE____END__ sycl-amdgcn-amd-amdhsa-
 }