Skip to content

Commit 8700b76

Browse files
steffenlarsensteffenlarsen
authored
[SYCL] Implement SYCL 2020 multi_ptr (#6893)
This commit adds implementations of [SYCL 2020 `multi_ptr`](https://registry.khronos.org/SYCL/specs/sycl-2020/html/sycl-2020.html#sec:multiptr) and `address_space_cast`. Likewise it adds the `legacy` decoration for SYCL 1.2.1 style `multi_ptr` as deprecated. To prevent breaking user code the legacy decoration is made the default. Test-suite changes: intel/llvm-test-suite#1293 Signed-off-by: Larsen, Steffen <[email protected]>
1 parent c657d06 commit 8700b76

35 files changed

+1685
-507
lines changed

‎clang/lib/Sema/SPIRVBuiltins.td

+6
Original file line numberDiff line numberDiff line change
@@ -840,6 +840,12 @@ foreach AS = [GlobalAS, LocalAS, PrivateAS] in {
840840
def : SPVBuiltin<"GenericCastToPtrExplicit", [PointerType<Char, AS>, PointerType<Char, GenericAS>], Attr.Const>;
841841
}
842842

843+
foreach Ty = [Void, ConstType<Void>, VolatileType<Void>, VolatileType<ConstType<Void>>] in {
844+
def : SPVBuiltin<"GenericCastToPtrExplicit_ToGlobal", [PointerType<Ty, GlobalAS>, PointerType<Ty, DefaultAS>, Int], Attr.Const>;
845+
def : SPVBuiltin<"GenericCastToPtrExplicit_ToLocal", [PointerType<Ty, LocalAS>, PointerType<Ty, DefaultAS>, Int], Attr.Const>;
846+
def : SPVBuiltin<"GenericCastToPtrExplicit_ToPrivate", [PointerType<Ty, PrivateAS>, PointerType<Ty, DefaultAS>, Int], Attr.Const>;
847+
}
848+
843849
foreach Type = TLFloat.List in {
844850
foreach v = [2, 3, 4, 8, 16] in {
845851
def : SPVBuiltin<"VectorTimesScalar", [VectorType<Type, v>, VectorType<Type, v>, Type], Attr.Const>;

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_group_load_store.asciidoc

+8-8
Original file line numberDiff line numberDiff line change
@@ -78,8 +78,8 @@ location at `src` + `get_local_id()`.
7878

7979
[source,c++]
8080
----
81-
template <typename T, access::address_space Space>
82-
T load(const multi_ptr<T, Space>* src)
81+
template <typename T, access::address_space Space, access::decorated IsDecorated>
82+
T load(const multi_ptr<T, Space, IsDecorated>* src)
8383
----
8484
_Constraints_: `T` must be a _NumericType_. `Space` must be
8585
`access::address_space::global_space` or `access::address_space::local_space`.
@@ -92,8 +92,8 @@ location at `src` + `get_local_id()`.
9292

9393
[source,c++]
9494
----
95-
template <int N, typename T, access::address_space Space>
96-
vec<T, N> load(const multi_ptr<T, Space> src)
95+
template <int N, typename T, access::address_space Space, access::decorated IsDecorated>
96+
vec<T, N> load(const multi_ptr<T, Space, IsDecorated> src)
9797
----
9898
_Constraints_: `T` must be a _NumericType_. `Space` must be
9999
`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
122122

123123
[source,c++]
124124
----
125-
template <typename T, access::address_space Space>
126-
void store(multi_ptr<T, Space> dst, const T& x)
125+
template <typename T, access::address_space Space, access::decorated IsDecorated>
126+
void store(multi_ptr<T, Space, IsDecorated> dst, const T& x)
127127
----
128128
_Constraints_: `T` must be a _NumericType_. `Space` must be
129129
`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
136136

137137
[source,c++]
138138
----
139-
template <typename T, access::address_space Space>
140-
void store(multi_ptr<T, Space> dst, const vec<T, N>& x)
139+
template <typename T, access::address_space Space, access::decorated IsDecorated>
140+
void store(multi_ptr<T, Space, IsDecorated> dst, const vec<T, N>& x)
141141
----
142142
_Constraints_: `T` must be a _NumericType_. `Space` must be
143143
`access::address_space::global_space` or `access::address_space::local_space`.

‎sycl/doc/extensions/experimental/sycl_ext_oneapi_matrix.asciidoc

+6-4
Original file line numberDiff line numberDiff line change
@@ -130,9 +130,10 @@ IMPORTANT: In the current implementation, only the subgroup scope is supported.
130130
namespace sycl::ext::oneapi::experimental::matrix {
131131
template <typename Group, typename T, size_t NumRows, size_t NumCols,
132132
matrix_layout Layout,
133-
access::address_space Space>
133+
access::address_space Space,
134+
access::decorated IsDecorated>
134135
void joint_matrix_load(Group sg, joint_matrix<T, NumRows, NumCols, Layout, Group> &res,
135-
multi_ptr<T, Space> src, size_t stride, matrix_layout MemLayout);
136+
multi_ptr<T, Space, IsDecorated> src, size_t stride, matrix_layout MemLayout);
136137
}
137138
```
138139
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
143144
namespace sycl::ext::oneapi::experimental::matrix {
144145
template <typename Group, typename T, size_t NumRows, size_t NumCols,
145146
matrix_layout L,
146-
access::address_space Space>
147+
access::address_space Space,
148+
access::decorated IsDecorated>
147149
void joint_matrix_store(Group sg, joint_matrix<T, NumRows, NumCols, L, Group> &res,
148-
multi_ptr<T, Space> src, size_t stride, matrix_layout memL);
150+
multi_ptr<T, Space, IsDecorated> src, size_t stride, matrix_layout memL);
149151
}
150152
```
151153
This function stores the data from the 2d tiles back to memory.

‎sycl/include/CL/__spirv/spirv_ops.hpp

+58-10
Original file line numberDiff line numberDiff line change
@@ -328,30 +328,78 @@ __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
328328
#undef __SPIRV_ATOMIC_UNSIGNED
329329
#undef __SPIRV_ATOMIC_XOR
330330

331-
extern SYCL_EXTERNAL __attribute__((opencl_global)) void *
332-
__spirv_GenericCastToPtrExplicit_ToGlobal(const void *Ptr,
333-
__spv::StorageClass::Flag S) noexcept;
334-
335-
extern SYCL_EXTERNAL __attribute__((opencl_local)) void *
336-
__spirv_GenericCastToPtrExplicit_ToLocal(const void *Ptr,
337-
__spv::StorageClass::Flag S) noexcept;
338-
339331
template <typename dataT>
340332
extern __attribute__((opencl_global)) dataT *
341-
__SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
333+
__SYCL_GenericCastToPtrExplicit_ToGlobal(void *Ptr) noexcept {
342334
return (__attribute__((opencl_global)) dataT *)
343335
__spirv_GenericCastToPtrExplicit_ToGlobal(
344336
Ptr, __spv::StorageClass::CrossWorkgroup);
345337
}
346338

339+
template <typename dataT>
340+
extern const __attribute__((opencl_global)) dataT *
341+
__SYCL_GenericCastToPtrExplicit_ToGlobal(const void *Ptr) noexcept {
342+
return (const __attribute__((opencl_global)) dataT *)
343+
__spirv_GenericCastToPtrExplicit_ToGlobal(
344+
Ptr, __spv::StorageClass::CrossWorkgroup);
345+
}
346+
347+
template <typename dataT>
348+
extern const volatile __attribute__((opencl_global)) dataT *
349+
__SYCL_GenericCastToPtrExplicit_ToGlobal(const volatile void *Ptr) noexcept {
350+
return (const volatile __attribute__((opencl_global)) dataT *)
351+
__spirv_GenericCastToPtrExplicit_ToGlobal(
352+
Ptr, __spv::StorageClass::CrossWorkgroup);
353+
}
354+
347355
template <typename dataT>
348356
extern __attribute__((opencl_local)) dataT *
349-
__SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
357+
__SYCL_GenericCastToPtrExplicit_ToLocal(void *Ptr) noexcept {
350358
return (__attribute__((opencl_local)) dataT *)
351359
__spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
352360
__spv::StorageClass::Workgroup);
353361
}
354362

363+
template <typename dataT>
364+
extern const __attribute__((opencl_local)) dataT *
365+
__SYCL_GenericCastToPtrExplicit_ToLocal(const void *Ptr) noexcept {
366+
return (const __attribute__((opencl_local)) dataT *)
367+
__spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
368+
__spv::StorageClass::Workgroup);
369+
}
370+
371+
template <typename dataT>
372+
extern const volatile __attribute__((opencl_local)) dataT *
373+
__SYCL_GenericCastToPtrExplicit_ToLocal(const volatile void *Ptr) noexcept {
374+
return (const volatile __attribute__((opencl_local)) dataT *)
375+
__spirv_GenericCastToPtrExplicit_ToLocal(Ptr,
376+
__spv::StorageClass::Workgroup);
377+
}
378+
379+
template <typename dataT>
380+
extern __attribute__((opencl_private)) dataT *
381+
__SYCL_GenericCastToPtrExplicit_ToPrivate(void *Ptr) noexcept {
382+
return (__attribute__((opencl_private)) dataT *)
383+
__spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
384+
__spv::StorageClass::Function);
385+
}
386+
387+
template <typename dataT>
388+
extern const __attribute__((opencl_private)) dataT *
389+
__SYCL_GenericCastToPtrExplicit_ToPrivate(const void *Ptr) noexcept {
390+
return (const __attribute__((opencl_private)) dataT *)
391+
__spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
392+
__spv::StorageClass::Function);
393+
}
394+
395+
template <typename dataT>
396+
extern const volatile __attribute__((opencl_private)) dataT *
397+
__SYCL_GenericCastToPtrExplicit_ToPrivate(const volatile void *Ptr) noexcept {
398+
return (const volatile __attribute__((opencl_private)) dataT *)
399+
__spirv_GenericCastToPtrExplicit_ToPrivate(Ptr,
400+
__spv::StorageClass::Function);
401+
}
402+
355403
template <typename dataT>
356404
__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT
357405
__spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;

0 commit comments

Comments
 (0)