[SYCL] Update checks in unsafe atomics test for AMDGPU (#16342)

The need for updated check was caused by a move from using AMD builtins
in favour of IR instructions with correct metadata attached, see:
b5e63cc

I travelled back in time to:
* intel/llvm:
358b777
* unified-runtime:
oneapi-src/unified-runtime@6406879

and verified that generated assembly (cpp -> IR -> s) is the same for
before and after the builtin change.

Author: jchlanda

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-
 }