[SYCL] Change check_device_code HIP tests to use SYCL_EXTERNAL

[ianayl]

Changed the HIP sycl/test/check_device_code lit tests to use SYCL_EXTERNAL functions instead of writing entire programs.

[konradkusiak97]

LGTM.

[hdelan]

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 0b78738a0151..4856f3bcba1a 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
@@ -5,45 +5,37 @@
 
 #include <sycl/sycl.hpp>
 
-class intKernel;
-class fpKernel;
+SYCL_EXTERNAL void intAtomicFunc(int *i) {
+  sycl::atomic_ref<int, sycl::memory_order_relaxed, sycl::memory_scope_device>
+      atomicInt(*i);
+  atomicInt.fetch_xor(1);
+  atomicInt.fetch_and(1);
+  atomicInt.fetch_or(1);
+  // CHECK: void{{.*}}intAtomicFunc
+  // CHECK-SAFE: cmpxchg volatile
+  // CHECK-SAFE-NOT: atomicrmw
+  // CHECK-UNSAFE: atomicrmw volatile xor
+  // CHECK-UNSAFE: atomicrmw volatile and
+  // CHECK-UNSAFE: atomicrmw volatile or
+  // CHECK-UNSAFE-NOT: cmpxchg
+}
 
-int main() {
-  int *i;
-  float *f;
-  double *d;
-  sycl::queue{}.single_task<intKernel>([=] {
-    sycl::atomic_ref<int, sycl::memory_order_relaxed, sycl::memory_scope_device>
-        atomicInt(*i);
-    atomicInt.fetch_xor(1);
-    atomicInt.fetch_and(1);
-    atomicInt.fetch_or(1);
-    // CHECK: amdgpu_kernel void{{.*}}intKernel
-    // CHECK-SAFE: cmpxchg volatile
-    // CHECK-SAFE-NOT: atomicrmw
-    // CHECK-UNSAFE: atomicrmw volatile xor
-    // CHECK-UNSAFE: atomicrmw volatile and
-    // CHECK-UNSAFE: atomicrmw volatile or
-    // CHECK-UNSAFE-NOT: cmpxchg
-  });
-  sycl::queue{}.single_task<fpKernel>([=] {
-    sycl::atomic_ref<float, sycl::memory_order_relaxed,
-                     sycl::memory_scope_device,
-                     sycl::access::address_space::global_space>(*f)
-        .fetch_add(1.0f);
-    // CHECK: amdgpu_kernel void{{.*}}fpKernel
-    // 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
-    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: __CLANG_OFFLOAD_BUNDLE____END__ sycl-amdgcn-amd-amdhsa-
-  });
+SYCL_EXTERNAL void fpAtomicFunc(float *f, double *d) {
+  sycl::atomic_ref<float, sycl::memory_order_relaxed, sycl::memory_scope_device,
+                   sycl::access::address_space::global_space>(*f)
+      .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
+  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: __CLANG_OFFLOAD_BUNDLE____END__ sycl-amdgcn-amd-amdhsa-
 }

This should work for the new changes in amdgpu_unsafe_atomics.cpp

[ianayl]

@intel/llvm-reviewers-cuda The PR is ready, is there anything else I need to do to get it merged?

Thanks!

[konradkusiak97]

Looks good to me. Please merge the sycl tip into this branch again, that will hopefully make the CI green.

[ianayl]

@konradkusiak97 I've merged sycl into the branch again, can you run the CI again? Thanks!

[konradkusiak97]

@intel/llvm-gatekeepers please merge

[sommerlukas]

@ianayl you need to update your Github profile settings to include your mail address in PRs.

Otherwise it will say

This commit will be authored by 54701975+ianayl@users.noreply.github.com

Akin to upstream LLVM policy, we don't allow commits with those mail addresses.

[ianayl]

@sommerlukas My bad! It's public now. I wouldn't need to redo my commits would I?

[sommerlukas]

@sommerlukas My bad! It's public now. I wouldn't need to redo my commits would I?

No, the commit resulting from "Squash and Merge" now has the correct author mail address after you changed the setttings, no worries.

I merged the PR now.

[ianayl]

Thanks!