[SYCL] Add basic support for the generic_space address space

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -239,7 +239,7 @@ extern SYCL_EXTERNAL TempRetT __spirv_ImageSampleExplicitLod(SampledType,
 
 #define __SPIRV_ATOMICS(macro, Arg)                                            \
   macro(__attribute__((opencl_global)), Arg)                                   \
-      macro(__attribute__((opencl_local)), Arg)
+      macro(__attribute__((opencl_local)), Arg) macro(, Arg)
 
 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, float)
 __SPIRV_ATOMICS(__SPIRV_ATOMIC_FLOAT, double)

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

@@ -162,6 +162,11 @@ struct DecoratedType<ElementType, access::address_space::private_space> {
   using type = __OPENCL_PRIVATE_AS__ ElementType;
 };
 
+template <typename ElementType>
+struct DecoratedType<ElementType, access::address_space::generic_space> {
+  using type = ElementType;
+};
+
 template <typename ElementType>
 struct DecoratedType<ElementType, access::address_space::global_space> {
   using type = __OPENCL_GLOBAL_AS__ ElementType;

sycl/include/CL/sycl/atomic_ref.hpp

@@ -114,11 +114,6 @@ template <> struct bit_equal<double> {
 template <typename T, memory_order DefaultOrder, memory_scope DefaultScope,
           access::address_space AddressSpace>
 class atomic_ref_base {
-  static_assert(
-      AddressSpace != access::address_space::generic_space,
-      "access::address_space::generic_space is a valid address space but the "
-      "address space is not supported yet.");
-
   static_assert(
       detail::IsValidAtomicRefType<T>::value,
       "Invalid atomic type.  Valid types are int, unsigned int, long, "

sycl/include/CL/sycl/multi_ptr.hpp

@@ -51,6 +51,13 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   multi_ptr(const multi_ptr &rhs) = default;
   multi_ptr(multi_ptr &&) = default;
 #ifdef __SYCL_DEVICE_ONLY__
+  // The generic address space have no corresponding 'opencl_...' attribute and
+  // this constructor is considered as a duplicate for the
+  // multi_ptr(ElementType *pointer) one, so the check is required.
+  template <
+      access::address_space _Space = Space,
+      typename = typename detail::enable_if_t<
+          _Space == Space && Space != access::address_space::generic_space>>
   multi_ptr(pointer_t pointer) : m_Pointer(pointer) {}
 #endif
 
@@ -71,6 +78,13 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   multi_ptr &operator=(multi_ptr &&) = default;
 
 #ifdef __SYCL_DEVICE_ONLY__
+  // The generic address space have no corresponding 'opencl_...' attribute and
+  // this operator is considered as a duplicate for the
+  // multi_ptr &operator=(ElementType *pointer) one, so the check is required.
+  template <
+      access::address_space _Space = Space,
+      typename = typename detail::enable_if_t<
+          _Space == Space && Space != access::address_space::generic_space>>
   multi_ptr &operator=(pointer_t pointer) {
     m_Pointer = pointer;
     return *this;
@@ -109,26 +123,28 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
     return reinterpret_cast<ReturnPtr>(m_Pointer)[index];
   }
 
-  // Only if Space == global_space || global_device_space
+  // Only if Space == global_space || global_device_space || generic_space
   template <int dimensions, access::mode Mode,
             access::placeholder isPlaceholder, typename PropertyListT,
             access::address_space _Space = Space,
             typename = typename detail::enable_if_t<
                 _Space == Space &&
-                (Space == access::address_space::global_space ||
+                (Space == access::address_space::generic_space ||
+                 Space == access::address_space::global_space ||
                  Space == access::address_space::global_device_space)>>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
                      isPlaceholder, PropertyListT>
                 Accessor) {
     m_Pointer = (pointer_t)(Accessor.get_pointer().get());
   }
 
-  // Only if Space == local_space
-  template <int dimensions, access::mode Mode,
-            access::placeholder isPlaceholder, typename PropertyListT,
-            access::address_space _Space = Space,
-            typename = typename detail::enable_if_t<
-                _Space == Space && Space == access::address_space::local_space>>
+  // Only if Space == local_space || generic_space
+  template <
+      int dimensions, access::mode Mode, access::placeholder isPlaceholder,
+      typename PropertyListT, access::address_space _Space = Space,
+      typename = typename detail::enable_if_t<
+          _Space == Space && (Space == access::address_space::generic_space ||
+                              Space == access::address_space::local_space)>>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
                      isPlaceholder, PropertyListT>
                 Accessor)
@@ -154,29 +170,32 @@ template <typename ElementType, access::address_space Space> class multi_ptr {
   //    2. from multi_ptr<ElementType, Space> to multi_ptr<const ElementType,
   //    Space>
 
-  // Only if Space == global_space || global_device_space and element type is
-  // const
+  // Only if Space == global_space || global_device_space || generic_space and
+  // element type is const
   template <
       int dimensions, access::mode Mode, access::placeholder isPlaceholder,
       typename PropertyListT, access::address_space _Space = Space,
       typename ET = ElementType,
       typename = typename detail::enable_if_t<
           _Space == Space &&
-          (Space == access::address_space::global_space ||
+          (Space == access::address_space::generic_space ||
+           Space == access::address_space::global_space ||
            Space == access::address_space::global_device_space) &&
           std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
   multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
                      access::target::device, isPlaceholder, PropertyListT>
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
-  // Only if Space == local_space and element type is const
+  // Only if Space == local_space || generic_space and element type is const
   template <
       int dimensions, access::mode Mode, access::placeholder isPlaceholder,
       typename PropertyListT, access::address_space _Space = Space,
       typename ET = ElementType,
       typename = typename detail::enable_if_t<
-          _Space == Space && Space == access::address_space::local_space &&
+          _Space == Space &&
+          (Space == access::address_space::generic_space ||
+           Space == access::address_space::local_space) &&
           std::is_const<ET>::value && std::is_same<ET, ElementType>::value>>
   multi_ptr(accessor<typename detail::remove_const_t<ET>, dimensions, Mode,
                      access::target::local, isPlaceholder, PropertyListT>
@@ -373,23 +392,26 @@ template <access::address_space Space> class multi_ptr<void, Space> {
     return *this;
   }
 
-  // Only if Space == global_space || global_device_space
+  // Only if Space == global_space || global_device_space || generic_space
   template <typename ElementType, int dimensions, access::mode Mode,
             typename PropertyListT, access::address_space _Space = Space,
             typename = typename detail::enable_if_t<
                 _Space == Space &&
-                (Space == access::address_space::global_space ||
+                (Space == access::address_space::generic_space ||
+                 Space == access::address_space::global_space ||
                  Space == access::address_space::global_device_space)>>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
                      access::placeholder::false_t, PropertyListT>
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
-  // Only if Space == local_space
-  template <typename ElementType, int dimensions, access::mode Mode,
-            typename PropertyListT, access::address_space _Space = Space,
-            typename = typename detail::enable_if_t<
-                _Space == Space && Space == access::address_space::local_space>>
+  // Only if Space == local_space || generic_space
+  template <
+      typename ElementType, int dimensions, access::mode Mode,
+      typename PropertyListT, access::address_space _Space = Space,
+      typename = typename detail::enable_if_t<
+          _Space == Space && (Space == access::address_space::generic_space ||
+                              Space == access::address_space::local_space)>>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
                      access::placeholder::false_t, PropertyListT>
                 Accessor)
@@ -493,23 +515,26 @@ class multi_ptr<const void, Space> {
     return *this;
   }
 
-  // Only if Space == global_space || global_device_space
+  // Only if Space == global_space || global_device_space || generic_space
   template <typename ElementType, int dimensions, access::mode Mode,
             typename PropertyListT, access::address_space _Space = Space,
             typename = typename detail::enable_if_t<
                 _Space == Space &&
-                (Space == access::address_space::global_space ||
+                (Space == access::address_space::generic_space ||
+                 Space == access::address_space::global_space ||
                  Space == access::address_space::global_device_space)>>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::device,
                      access::placeholder::false_t, PropertyListT>
                 Accessor)
       : multi_ptr(Accessor.get_pointer()) {}
 
-  // Only if Space == local_space
-  template <typename ElementType, int dimensions, access::mode Mode,
-            typename PropertyListT, access::address_space _Space = Space,
-            typename = typename detail::enable_if_t<
-                _Space == Space && Space == access::address_space::local_space>>
+  // Only if Space == local_space || generic_space
+  template <
+      typename ElementType, int dimensions, access::mode Mode,
+      typename PropertyListT, access::address_space _Space = Space,
+      typename = typename detail::enable_if_t<
+          _Space == Space && (Space == access::address_space::generic_space ||
+                              Space == access::address_space::local_space)>>
   multi_ptr(accessor<ElementType, dimensions, Mode, access::target::local,
                      access::placeholder::false_t, PropertyListT>
                 Accessor)

sycl/include/CL/sycl/pointers.hpp

@@ -16,6 +16,10 @@ namespace sycl {
 template <typename ElementType, access::address_space Space> class multi_ptr;
 // Template specialization aliases for different pointer address spaces
 
+template <typename ElementType>
+using generic_ptr =
+    multi_ptr<ElementType, access::address_space::generic_space>;
+
 template <typename ElementType>
 using global_ptr = multi_ptr<ElementType, access::address_space::global_space>;
 

sycl/test/basic_tests/atomic-ref-instantiation.cpp

@@ -1,5 +1,5 @@
 // RUN: %clangxx -fsycl -fsyntax-only -Xclang -verify %s -o %t.out -Xclang -verify-ignore-unexpected=note
-
+// expected-no-diagnostics
 #include <CL/sycl/atomic_ref.hpp>
 
 struct A {};
@@ -20,6 +20,5 @@ int main() {
   A* p = &a;
   auto ref_p = sycl::atomic_ref<A *, sycl::memory_order_acq_rel,
                                 sycl::memory_scope_device>(p);
-  // expected-error@CL/sycl/atomic_ref.hpp:* {{"access::address_space::generic_space is a valid address space but the address space is not supported yet."}}
   return 0;
 }