[SYCL] Move an accessor pointer to global_device space

With this patch an accessor pointer to global buffer is moved from
global space to global_device space. That is done to distinguish this
pointer from those USM pointers, that are allocated global space or
global_host space, in compile time.

In addition to this change there are added explicit conversion operator
from global_device to global space for multi_ptr class and implicit
convertion for atomic class from global_device for global space.
The last change isn't covered by specification published here:
#1840 , but is required to pass
atomic_api CTS.

Signed-off-by: Dmitry Sidorov <[email protected]>

Author: MrSidims

sycl/include/CL/sycl/access/access.hpp

@@ -124,6 +124,11 @@ template <access::target accessTarget> struct TargetToAS {
       access::address_space::global_space;
 };
 
+template <> struct TargetToAS<access::target::global_buffer> {
+  constexpr static access::address_space AS =
+      access::address_space::global_device_space;
+};
+
 template <> struct TargetToAS<access::target::local> {
   constexpr static access::address_space AS =
       access::address_space::local_space;

sycl/include/CL/sycl/atomic.hpp

@@ -184,6 +184,7 @@ class atomic {
       detail::GetSpirvMemoryScope<addressSpace>::scope;
 
 public:
+  using AtomicPtrType = typename detail::PtrValueType<T, addressSpace>::type;
   template <typename pointerT>
 #ifdef __SYCL_DEVICE_ONLY__
   atomic(multi_ptr<pointerT, addressSpace> ptr)
@@ -197,6 +198,37 @@ class atomic {
                   "T and pointerT must be same size");
   }
 
+#ifdef __SYCL_DEVICE_ONLY__
+  AtomicPtrType *getPtr() const { return Ptr; }
+#else
+  std::atomic<T> *getPtr() const { return Ptr; }
+#endif
+
+  // Create atomic in global_space with one from global_device_space
+  template <access::address_space _Space = addressSpace,
+            typename = typename std::enable_if<
+                _Space == addressSpace &&
+                addressSpace == access::address_space::global_space>::type>
+  atomic(const atomic<T, access::address_space::global_device_space> &RHS) {
+#ifdef __SYCL_DEVICE_ONLY__
+    Ptr = RHS.getPtr();
+#else
+    Ptr = reinterpret_cast<std::atomic<T> *>(RHS.getPtr());
+#endif
+  }
+
+  template <access::address_space _Space = addressSpace,
+            typename = typename std::enable_if<
+                _Space == addressSpace &&
+                addressSpace == access::address_space::global_space>::type>
+  atomic(const atomic<T, access::address_space::global_device_space> &&RHS) {
+#ifdef __SYCL_DEVICE_ONLY__
+    Ptr = RHS.getPtr();
+#else
+    Ptr = reinterpret_cast<std::atomic<T> *>(RHS.getPtr());
+#endif
+  }
+
   void store(T Operand, memory_order Order = memory_order::relaxed) {
     __spirv_AtomicStore(
         Ptr, SpirvScope, detail::getSPIRVMemorySemanticsMask(Order), Operand);
@@ -299,7 +331,7 @@ class atomic {
 
 private:
 #ifdef __SYCL_DEVICE_ONLY__
-  typename detail::PtrValueType<T, addressSpace>::type *Ptr;
+  AtomicPtrType *Ptr;
 #else
   std::atomic<T> *Ptr;
 #endif

sycl/include/CL/sycl/handler.hpp

@@ -506,7 +506,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, T>
   readFromFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Src) const {
-    atomic<T, access::address_space::global_space> AtomicSrc = Src;
+    atomic<T, access::address_space::global_device_space> AtomicSrc = Src;
     return AtomicSrc.load();
   }
 
@@ -529,7 +529,7 @@ class __SYCL_EXPORT handler {
             access::placeholder IsPH>
   detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, void>
   writeToFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Dst, T V) const {
-    atomic<T, access::address_space::global_space> AtomicDst = Dst;
+    atomic<T, access::address_space::global_device_space> AtomicDst = Dst;
     AtomicDst.store(V);
   }
 

sycl/include/CL/sycl/intel/atomic_ref.hpp

@@ -43,7 +43,8 @@ using IsValidAtomicType =
 template <cl::sycl::access::address_space AS>
 using IsValidAtomicAddressSpace =
     bool_constant<AS == access::address_space::global_space ||
-                  AS == access::address_space::local_space>;
+                  AS == access::address_space::local_space ||
+                  AS == access::address_space::global_device_space>;
 
 // DefaultOrder parameter is limited to read-modify-write orders
 template <memory_order Order>
@@ -138,7 +139,7 @@ class atomic_ref_base {
                 "intel::atomic_ref does not yet support pointer types");
   static_assert(detail::IsValidAtomicAddressSpace<AddressSpace>::value,
                 "Invalid atomic address_space.  Valid address spaces are: "
-                "global_space, local_space");
+                "global_space, local_space, global_device_space");
   static_assert(
       detail::IsValidDefaultOrder<DefaultOrder>::value,
       "Invalid default memory_order for atomics.  Valid defaults are: "

sycl/include/CL/sycl/multi_ptr.hpp

@@ -275,6 +275,22 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return multi_ptr(m_Pointer - r);
   }
 
+  // Explicit conversion to global_space
+  // Only available if Space == address_space::global_device_space ||
+  // Space == address_space::global_host_space
+  template <access::address_space _Space = Space,
+            typename = typename std::enable_if<
+                _Space == Space &&
+                (Space == access::address_space::global_device_space ||
+                 Space == access::address_space::global_host_space)>::type>
+  explicit
+  operator multi_ptr<ElementType, access::address_space::global_space>() const {
+    using global_pointer_t = typename detail::PtrValueType<
+        ElementType, access::address_space::global_space>::type *;
+    return multi_ptr<ElementType, access::address_space::global_space>(
+        (global_pointer_t)m_Pointer);
+  }
+
   // Only if Space == global_space
   template <access::address_space _Space = Space,
             typename = typename std::enable_if<

sycl/test/check_device_code/kernel_arguments_as.cpp

@@ -5,7 +5,7 @@
 //
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
 // CHECK: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}AccessorImplDevice", %[[UNION:.*]] }
-// CHECK: %[[UNION]] = type { i32 addrspace(1)* }
+// CHECK: %[[UNION]] = type { i32 addrspace(5)* }
 // CHECK: %struct{{.*}}AccWrapper = type { %"class{{.*}}cl::sycl::accessor" }
 // CHECK-NEXT: %"class{{.*}}cl::sycl::accessor" = type { %"class{{.*}}LocalAccessorBaseDevice", i32 addrspace(3)* }
 //

sycl/test/multi_ptr/multi_ptr.cpp

@@ -98,6 +98,8 @@ template <typename T> void testMultPtr() {
         global_ptr<void> ptr_8 = global_ptr<void>(ptr_7);
         host_ptr<void> ptr_9((void *)RawPtr);
         global_ptr<void> ptr_10 = global_ptr<void>(ptr_9);
+        device_ptr<T> ptr_11(accessorData_1);
+        global_ptr<T> ptr_12 = global_ptr<T>(ptr_11);
 
         innerFunc<T>(wiID.get(0), ptr_1, ptr_2, local_ptr);
       });

sycl/test/regression/implicit_atomic_conversion.cpp

@@ -0,0 +1,29 @@
+// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+void test_conversion(queue q) {
+  int init = 0;
+  {
+    buffer<int> in_buf(&init, 1);
+
+    q.submit([&](handler &cgh) {
+      auto in = in_buf.template get_access<access::mode::atomic>(cgh);
+      cgh.single_task<class conversion>([=]() {
+        cl::sycl::atomic<int, access::address_space::global_space> atm = in[0];
+        atm.store(42);
+      });
+    });
+  }
+  assert(init == 42 && "verification failed");
+}
+
+int main() {
+  queue q;
+  test_conversion(q);
+}