[SYCL] Add atomic64 aspect decoration to atomic_ref<T *> (#14052)

atomic_ref<T *> uses 64-bit atomics and it should be decorated with the
corresponding aspect.

fixes: #12743

Author: maksimsab

sycl/include/sycl/atomic_ref.hpp

@@ -568,9 +568,14 @@ class [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]] atomic_ref_impl<
 // Partial specialization for pointer types
 // Arithmetic is emulated because target's representation of T* is unknown
 // TODO: Find a way to use intptr_t or uintptr_t atomics instead
-template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder, memory_scope DefaultScope,
-          access::address_space AddressSpace>
-class atomic_ref_impl<T *, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope, AddressSpace>
+template <typename T, bool IsAspectAtomic64AttrUsed, memory_order DefaultOrder,
+          memory_scope DefaultScope, access::address_space AddressSpace>
+#ifndef __SYCL_DEVICE_ONLY__
+class atomic_ref_impl<
+#else
+class [[__sycl_detail__::__uses_aspects__(aspect::atomic64)]] atomic_ref_impl<
+#endif
+    T *, IsAspectAtomic64AttrUsed, DefaultOrder, DefaultScope, AddressSpace>
     : public atomic_ref_base<uintptr_t, DefaultOrder, DefaultScope,
                              AddressSpace> {
 

sycl/test-e2e/syclcompat/atomic/atomic_arith.cpp

@@ -32,9 +32,7 @@
 
 // UNSUPPORTED: hip
 
-// FIXME: Remove "-fsycl-device-code-split=per_kernel" option after fixing
-// https://github.com/intel/llvm/issues/12743.
-// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} -fsycl-device-code-split=per_kernel %s -o %t.out
+// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
 // RUN: %{run} %t.out
 
 #include <cstddef>

sycl/test-e2e/syclcompat/atomic/atomic_comp_exchange.cpp

@@ -32,9 +32,7 @@
 
 // UNSUPPORTED: hip
 
-// FIXME: Remove "-fsycl-device-code-split=per_kernel" option after fixing
-// https://github.com/intel/llvm/issues/12743.
-// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} -fsycl-device-code-split=per_kernel %s -o %t.out
+// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
 // RUN: %{run} %t.out
 
 #include <type_traits>

sycl/test-e2e/syclcompat/math/math_ops.cpp

@@ -20,9 +20,7 @@
  *    tests for non-vectorized math helper functions
  **************************************************************************/
 
-// FIXME: Remove "-fsycl-device-code-split=per_kernel" option after fixing
-// https://github.com/intel/llvm/issues/12743.
-// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} -fsycl-device-code-split=per_kernel %s -o %t.out
+// RUN: %clangxx -std=c++20 -fsycl -fsycl-targets=%{sycl_triple} %s -o %t.out
 // RUN: %{run} %t.out
 
 #include <syclcompat/dims.hpp>

sycl/test/optional_kernel_features/atomic_ref-atomic64-aspect.cpp

@@ -1,10 +1,11 @@
 // RUN: %clangxx %s -S -o %t.ll -fsycl-device-only -Xclang -disable-llvm-passes
 // RUN: FileCheck %s --input-file %t.ll
 
-// CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM1:]], ![[#MDNUM2:]], ![[#MDNUM3:]]}
-// CHECK: ![[#MDNUM1]] = !{!"class.sycl::_V1::detail::atomic_ref_impl", i32 [[#ASPECT_NUM:]]}
-// CHECK: ![[#MDNUM2]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.2", i32 [[#ASPECT_NUM:]]}
-// CHECK: ![[#MDNUM3]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.7", i32 [[#ASPECT_NUM:]]}
+// CHECK: !sycl_types_that_use_aspects = !{![[#MDNUM1:]], ![[#MDNUM2:]], ![[#MDNUM3:]], ![[#MDNUM4:]]}
+// CHECK: ![[#MDNUM1]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.20", i32 [[#ASPECT_NUM:]]}
+// CHECK-NEXT: ![[#MDNUM2]] = !{!"class.sycl::_V1::detail::atomic_ref_impl", i32 [[#ASPECT_NUM:]]}
+// CHECK-NEXT: ![[#MDNUM3]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.2", i32 [[#ASPECT_NUM:]]}
+// CHECK-NEXT: ![[#MDNUM4]] = !{!"class.sycl::_V1::detail::atomic_ref_impl.7", i32 [[#ASPECT_NUM:]]}
 // CHECK: !{{.*}} = !{!"atomic64", i32 [[#ASPECT_NUM]]}
 
 #include <sycl/sycl.hpp>
@@ -46,6 +47,12 @@ int main() {
           sycl::atomic_ref<int, sycl::memory_order_acq_rel,
                            sycl::memory_scope_device,
                            sycl::access::address_space::local_space>(val_int);
+
+      double *ptr = nullptr;
+      auto ref_double_ptr =
+          sycl::atomic_ref<double *, sycl::memory_order_acq_rel,
+                           sycl::memory_scope_device,
+                           sycl::access::address_space::local_space>(ptr);
     });
   });
   return 0;