[SYCL] Implement SYCL 2020 multi_ptr

clang/lib/Sema/SPIRVBuiltins.td

@@ -840,6 +840,12 @@ foreach AS = [GlobalAS, LocalAS, PrivateAS] in {
   def : SPVBuiltin<"GenericCastToPtrExplicit", [PointerType<Char, AS>, PointerType<Char, GenericAS>], Attr.Const>;
 }
 
+foreach Ty = [Void, ConstType<Void>, VolatileType<Void>, VolatileType<ConstType<Void>>] in {
+  def : SPVBuiltin<"GenericCastToPtrExplicit_ToGlobal", [PointerType<Ty, GlobalAS>, PointerType<Ty, DefaultAS>, Int], Attr.Const>;
+  def : SPVBuiltin<"GenericCastToPtrExplicit_ToLocal", [PointerType<Ty, LocalAS>, PointerType<Ty, DefaultAS>, Int], Attr.Const>;
+  def : SPVBuiltin<"GenericCastToPtrExplicit_ToPrivate", [PointerType<Ty, PrivateAS>, PointerType<Ty, DefaultAS>, Int], Attr.Const>;
+}
+
 foreach Type = TLFloat.List in {
   foreach v = [2, 3, 4, 8, 16] in {
     def : SPVBuiltin<"VectorTimesScalar", [VectorType<Type, v>, VectorType<Type, v>, Type], Attr.Const>;

sycl/doc/extensions/experimental/sycl_ext_oneapi_group_load_store.asciidoc

@@ -78,8 +78,8 @@ location at `src` + `get_local_id()`.
 
 [source,c++]
 ----
-template <typename T, access::address_space Space>
-T load(const multi_ptr<T, Space>* src)
+template <typename T, access::address_space Space, access::decorated IsDecorated>
+T load(const multi_ptr<T, Space, IsDecorated>* src)
 ----
 _Constraints_: `T` must be a _NumericType_. `Space` must be
 `access::address_space::global_space` or `access::address_space::local_space`.
@@ -92,8 +92,8 @@ location at `src` + `get_local_id()`.
 
 [source,c++]
 ----
-template <int N, typename T, access::address_space Space>
-vec<T, N> load(const multi_ptr<T, Space> src)
+template <int N, typename T, access::address_space Space, access::decorated IsDecorated>
+vec<T, N> load(const multi_ptr<T, Space, IsDecorated> src)
 ----
 _Constraints_: `T` must be a _NumericType_. `Space` must be
 `access::address_space::global_space` or `access::address_space::local_space`.
@@ -122,8 +122,8 @@ _Effects_: Writes the value of `x` from each work-item to the memory location at
 
 [source,c++]
 ----
-template <typename T, access::address_space Space>
-void store(multi_ptr<T, Space> dst, const T& x)
+template <typename T, access::address_space Space, access::decorated IsDecorated>
+void store(multi_ptr<T, Space, IsDecorated> dst, const T& x)
 ----
 _Constraints_: `T` must be a _NumericType_. `Space` must be
 `access::address_space::global_space` or `access::address_space::local_space`.
@@ -136,8 +136,8 @@ _Effects_: Writes the value of `x` from each work-item to the memory location at
 
 [source,c++]
 ----
-template <typename T, access::address_space Space>
-void store(multi_ptr<T, Space> dst, const vec<T, N>& x)
+template <typename T, access::address_space Space, access::decorated IsDecorated>
+void store(multi_ptr<T, Space, IsDecorated> dst, const vec<T, N>& x)
 ----
 _Constraints_: `T` must be a _NumericType_. `Space` must be
 `access::address_space::global_space` or `access::address_space::local_space`.

sycl/doc/extensions/experimental/sycl_ext_oneapi_matrix.asciidoc

@@ -130,9 +130,10 @@ IMPORTANT: In the current implementation, only the subgroup scope is supported.
 namespace sycl::ext::oneapi::experimental::matrix {
   template <typename Group, typename T, size_t NumRows, size_t NumCols,
           matrix_layout Layout,
-          access::address_space Space>
+          access::address_space Space,
+          access::decorated IsDecorated>
   void joint_matrix_load(Group sg, joint_matrix<T, NumRows, NumCols, Layout, Group> &res,
-		    multi_ptr<T, Space> src, size_t stride, matrix_layout MemLayout);
+		    multi_ptr<T, Space, IsDecorated> src, size_t stride, matrix_layout MemLayout);
 }
 ```
 This function loads data from memory to the 2d tiles/registers of Intel AMX/DPAS.
@@ -143,9 +144,10 @@ This function loads data from memory to the 2d tiles/registers of Intel AMX/DPAS
 namespace sycl::ext::oneapi::experimental::matrix {
   template <typename Group, typename T, size_t NumRows, size_t NumCols,
           matrix_layout L,
-          access::address_space Space>	  
+          access::address_space Space,
+          access::decorated IsDecorated>
   void joint_matrix_store(Group sg, joint_matrix<T, NumRows, NumCols, L, Group> &res,
-		     multi_ptr<T, Space> src, size_t stride, matrix_layout memL);
+		     multi_ptr<T, Space, IsDecorated> src, size_t stride, matrix_layout memL);
 }
 ```
 This function stores the data from the 2d tiles back to memory.

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -315,30 +315,78 @@ __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
 #undef __SPIRV_ATOMIC_UNSIGNED
 #undef __SPIRV_ATOMIC_XOR
 
-extern SYCL_EXTERNAL __attribute__((opencl_global)) void *
-__spirv_GenericCastToPtrExplicit_ToGlobal(const void *Ptr,
-                                          __spv::StorageClass::Flag S) noexcept;
-
-extern SYCL_EXTERNAL __attribute__((opencl_local)) void *
-__spirv_GenericCastToPtrExplicit_ToLocal(const void *Ptr,
-                                         __spv::StorageClass::Flag S) noexcept;
-
 template <typename dataT>
 extern __attribute__((opencl_global)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
+__SYCL_GenericCastToPtrExplicit_ToGlobal(void *Ptr) noexcept {
   return (__attribute__((opencl_global)) dataT *)
       __spirv_GenericCastToPtrExplicit_ToGlobal(
           Ptr, __spv::StorageClass::CrossWorkgroup);
 }
 
+template <typename dataT>
+extern const __attribute__((opencl_global)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
+  return (const __attribute__((opencl_global)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToGlobal(
+          Ptr, __spv::StorageClass::CrossWorkgroup);
+}
+
+template <typename dataT>
+extern const volatile __attribute__((opencl_global)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToGlobal(const volatile void *Ptr) noexcept {
+  return (const volatile __attribute__((opencl_global)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToGlobal(
+          Ptr, __spv::StorageClass::CrossWorkgroup);
+}
+
 template <typename dataT>
 extern __attribute__((opencl_local)) dataT *
-__SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
+__SYCL_GenericCastToPtrExplicit_ToLocal(void *Ptr) noexcept {
   return (__attribute__((opencl_local)) dataT *)
       __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
                                                __spv::StorageClass::Workgroup);
 }
 
+template <typename dataT>
+extern const __attribute__((opencl_local)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
+  return (const __attribute__((opencl_local)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
+                                               __spv::StorageClass::Workgroup);
+}
+
+template <typename dataT>
+extern const volatile __attribute__((opencl_local)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToLocal(const volatile void *Ptr) noexcept {
+  return (const volatile __attribute__((opencl_local)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
+                                               __spv::StorageClass::Workgroup);
+}
+
+template <typename dataT>
+extern __attribute__((opencl_private)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToPrivate(void *Ptr) noexcept {
+  return (__attribute__((opencl_private)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
+                                                 __spv::StorageClass::Function);
+}
+
+template <typename dataT>
+extern const __attribute__((opencl_private)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToPrivate(const void *Ptr) noexcept {
+  return (const __attribute__((opencl_private)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
+                                                 __spv::StorageClass::Function);
+}
+
+template <typename dataT>
+extern const volatile __attribute__((opencl_private)) dataT *
+__SYCL_GenericCastToPtrExplicit_ToPrivate(const volatile void *Ptr) noexcept {
+  return (const volatile __attribute__((opencl_private)) dataT *)
+      __spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
+                                                 __spv::StorageClass::Function);
+}
+
 template <typename dataT>
 __SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT
 __spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;