Skip to content

Commit 828659c

Browse files
vladimirlazvladimirlaz
committed
Allow load/store only for pointers from local and global AS
1 parent 50e85ce commit 828659c

File tree

5 files changed

+23
-65
lines changed

5 files changed

+23
-65
lines changed

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

Lines changed: 0 additions & 20 deletions
Original file line numberDiff line numberDiff line change
@@ -169,8 +169,6 @@ __SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned long)
169169
__SPIRV_ATOMICS(__SPIRV_ATOMIC_UNSIGNED, unsigned long long)
170170
__SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Min)
171171
__SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
172-
extern SYCL_EXTERNAL __attribute__((opencl_generic)) void *
173-
__spirv_PtrCastToGeneric(const void *Ptr) noexcept;
174172

175173
extern SYCL_EXTERNAL __attribute__((opencl_global)) void *
176174
__spirv_GenericCastToPtrExplicit_ToGlobal(const void *Ptr,
@@ -180,16 +178,6 @@ extern SYCL_EXTERNAL __attribute__((opencl_local)) void *
180178
__spirv_GenericCastToPtrExplicit_ToLocal(const void *Ptr,
181179
__spv::StorageClass::Flag S) noexcept;
182180

183-
extern SYCL_EXTERNAL __attribute__((opencl_private)) void *
184-
__spirv_GenericCastToPtrExplicit_ToPrivate(
185-
const void *Ptr, __spv::StorageClass::Flag S) noexcept;
186-
187-
template <typename dataT>
188-
extern __attribute__((opencl_generic)) dataT *
189-
__spirv_PtrCastToGeneric(const void *Ptr) noexcept {
190-
return (__attribute__((opencl_generic)) dataT *)__spirv_PtrCastToGeneric(Ptr);
191-
}
192-
193181
template <typename dataT>
194182
extern __attribute__((opencl_global)) dataT *
195183
__spirv_GenericCastToPtrExplicit_ToGlobal(
@@ -206,14 +194,6 @@ __spirv_GenericCastToPtrExplicit_ToLocal(const void *Ptr,
206194
dataT *)__spirv_GenericCastToPtrExplicit_ToLocal(Ptr, S);
207195
}
208196

209-
template <typename dataT>
210-
extern __attribute__((opencl_private)) dataT *
211-
__spirv_GenericCastToPtrExplicit_ToPrivate(
212-
const void *Ptr, __spv::StorageClass::Flag S) noexcept {
213-
return (__attribute__((opencl_private))
214-
dataT *)__spirv_GenericCastToPtrExplicit_ToPrivate(Ptr, S);
215-
}
216-
217197
template <typename dataT>
218198
__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT
219199
__spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;

‎sycl/include/CL/sycl/ONEAPI/sub_group.hpp

Lines changed: 5 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -254,13 +254,8 @@ struct sub_group {
254254
if (g)
255255
return load(g);
256256

257-
auto p = __spirv_GenericCastToPtrExplicit_ToPrivate<T>(
258-
src, __spv::StorageClass::Function);
259-
assert((p == nullptr) &&
260-
"Sub-group load() is not supported for private pointers.");
261-
262-
// Fallback for other address spaces to be mapped to global
263-
return load(__spirv_PtrCastToGeneric<T>(src));
257+
assert(!"Sub-group load() is supported for local or global pointers only.");
258+
return 0;
264259
#endif // __NVPTX__
265260
}
266261
#else //__SYCL_DEVICE_ONLY__
@@ -396,13 +391,9 @@ struct sub_group {
396391
return;
397392
}
398393

399-
auto p = __spirv_GenericCastToPtrExplicit_ToPrivate<T>(
400-
dst, __spv::StorageClass::Function);
401-
assert((p == nullptr) &&
402-
"Sub-group store() is not supported for private pointers.");
403-
404-
// Fallback for other address spaces to be mapped to global
405-
store(__spirv_PtrCastToGeneric<T>(dst), x);
394+
assert(
395+
!"Sub-group store() is supported for local or global pointers only.");
396+
return;
406397
#endif // __NVPTX__
407398
}
408399
#else //__SYCL_DEVICE_ONLY__

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

Lines changed: 0 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -115,15 +115,13 @@ constexpr bool modeWritesNewData(access::mode m) {
115115
#define __OPENCL_LOCAL_AS__ __attribute__((opencl_local))
116116
#define __OPENCL_CONSTANT_AS__ __attribute__((opencl_constant))
117117
#define __OPENCL_PRIVATE_AS__ __attribute__((opencl_private))
118-
#define __OPENCL_GENERIC_AS__ __attribute__((opencl_generic))
119118
#else
120119
#define __OPENCL_GLOBAL_AS__
121120
#define __OPENCL_GLOBAL_DEVICE_AS__
122121
#define __OPENCL_GLOBAL_HOST_AS__
123122
#define __OPENCL_LOCAL_AS__
124123
#define __OPENCL_CONSTANT_AS__
125124
#define __OPENCL_PRIVATE_AS__
126-
#define __OPENCL_GENERIC_AS__
127125
#endif
128126

129127
template <access::target accessTarget> struct TargetToAS {
@@ -229,10 +227,6 @@ template <class T> struct remove_AS<__OPENCL_CONSTANT_AS__ T> {
229227
typedef T type;
230228
};
231229

232-
template <class T> struct remove_AS<__OPENCL_GENERIC_AS__ T> {
233-
typedef T type;
234-
};
235-
236230
template <class T> struct deduce_AS<__OPENCL_PRIVATE_AS__ T> {
237231
static const access::address_space value =
238232
access::address_space::private_space;
@@ -246,11 +240,6 @@ template <class T> struct deduce_AS<__OPENCL_CONSTANT_AS__ T> {
246240
static const access::address_space value =
247241
access::address_space::constant_space;
248242
};
249-
250-
template <class T> struct deduce_AS<__OPENCL_GENERIC_AS__ T> {
251-
static const access::address_space value =
252-
access::address_space::global_space;
253-
};
254243
#endif
255244

256245
#undef __OPENCL_GLOBAL_AS__

‎sycl/test/extensions/sub_group_as.cpp

Lines changed: 17 additions & 19 deletions
Original file line numberDiff line numberDiff line change
@@ -29,13 +29,6 @@ int main(int argc, char *argv[]) {
2929
sycl::access::target::local>
3030
local(N, cgh);
3131

32-
// Check that load/store functions for raw pointer was called
33-
//
34-
// CHECK: spir_kernel void @"_ZTSZZ4mainENK3$_0clERN2cl4sycl7handlerEE4test"
35-
// CHECK-COUNT-3: call spir_func i32
36-
// {{.*}}loadIiEENSt9enable_ifIXsr3std7is_sameINS0_6detail9remove_ASIT_E4typeES7_EE5valueES7_E4typeEPS7_{{.*}}i32
37-
// addrspace(4)*
38-
// CHECK: call spir_func void {{.*}}storeIiEENSt9enable_ifIXsr3std7is_sameINS0_6detail9remove_ASIT_E4typeES7_EE5valueEvE4typeEPS7_RKS9_({{.*}} addrspace(4)*{{.*}},
3932
cgh.parallel_for<class test>(
4033
cl::sycl::nd_range<1>(N, 32), [=](cl::sycl::nd_item<1> it) {
4134
int v[N] = {0, 1, 2, 3, 4, 5, 6, 7, 8, 9, 10, 11, 12,
@@ -50,37 +43,42 @@ int main(int argc, char *argv[]) {
5043
local[i] = i;
5144
}
5245
}
46+
// CHECK: call void @_Z22__spirv_ControlBarrierjjj
5347
it.barrier();
5448

5549
int i = (it.get_global_id(0) / sg.get_max_local_range()[0]) *
5650
sg.get_max_local_range()[0];
5751

58-
// CHECK: spir_func i32{{.*}}loadIiEENSt9enable_ifIXsr3std7is_sameINS0_6detail9remove_ASIT_E4typeES7_EE5valueES7_E4typeEPS7_
52+
// load for global address space
5953
// CHECK: call spir_func i8 addrspace(3)* @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
60-
// CHECK: call spir_func i32 @_Z33__spirv_SubgroupLocalInvocationIdv
54+
// CHECK: call spir_func i32 @_Z33__spirv_SubgroupLocalInvocationIdv()
6155
// CHECK: call spir_func i8 addrspace(1)* @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
6256
// CHECK: call spir_func i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(i32 addrspace(1)*
63-
// CHECK: call spir_func i8* @_Z42__spirv_GenericCastToPtrExplicit_ToPrivatePKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
6457
// CHECK: call spir_func void {{.*}}assert
65-
// CHECK: call spir_func i8 addrspace(4)* @_Z24__spirv_PtrCastToGenericPKv(i8 addrspace(4)*
66-
// CHECK: call spir_func i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(i32 addrspace(1)*
67-
// Global address space
6858
auto x = sg.load(&global[i]);
6959

70-
// Local address space
60+
// load() for local address space
61+
// CHECK: call spir_func i8 addrspace(3)* @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
62+
// CHECK: call spir_func i32 @_Z33__spirv_SubgroupLocalInvocationIdv()
63+
// CHECK: call spir_func i8 addrspace(1)* @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
64+
// CHECK: call spir_func i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(i32 addrspace(1)*
65+
// CHECK: call spir_func void {{.*}}assert
7166
auto y = sg.load(&local[i]);
7267

68+
// load() for private address space
69+
// CHECK: call spir_func i8 addrspace(3)* @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
70+
// CHECK: call spir_func i32 @_Z33__spirv_SubgroupLocalInvocationIdv()
71+
// CHECK: call spir_func i8 addrspace(1)* @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
72+
// CHECK: call spir_func i32 @_Z30__spirv_SubgroupBlockReadINTELIjET_PU3AS1Kj(i32 addrspace(1)*
73+
// CHECK: call spir_func void {{.*}}assert
7374
auto z = sg.load(v + i);
7475

75-
// CHECK: spir_func void {{.*}}storeIiEENSt9enable_ifIXsr3std7is_sameINS0_6detail9remove_ASIT_E4typeES7_EE5valueEvE4typeEPS7_RKS9_
76+
// store() for global address space
7677
// CHECK: call spir_func i8 addrspace(3)* @_Z40__spirv_GenericCastToPtrExplicit_ToLocalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
77-
// CHECK: call spir_func i32 @_Z33__spirv_SubgroupLocalInvocationIdv
78+
// CHECK: call spir_func i32 @_Z33__spirv_SubgroupLocalInvocationIdv() #7, !noalias !29
7879
// CHECK: call spir_func i8 addrspace(1)* @_Z41__spirv_GenericCastToPtrExplicit_ToGlobalPKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
7980
// CHECK: call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(i32 addrspace(1)*
80-
// CHECK: call spir_func i8* @_Z42__spirv_GenericCastToPtrExplicit_ToPrivatePKvN5__spv12StorageClass4FlagE(i8 addrspace(4)*
8181
// CHECK: call spir_func void {{.*}}assert
82-
// CHECK: call spir_func i8 addrspace(4)* @_Z24__spirv_PtrCastToGenericPKv(i8 addrspace(4)*
83-
// CHECK: call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(i32 addrspace(1)*
8482
sg.store(&global[i], x + y + z);
8583
});
8684
});

‎sycl/test/on-device/extensions/sub_group_as_private.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -62,7 +62,7 @@ int main(int argc, char *argv[]) {
6262
// Local address space
6363
auto y = sg.load(&local[i]);
6464

65-
// CHECK: Sub-group load() is not supported for private pointers.
65+
// CHECK: Sub-group load() is supported for local or global pointers only
6666
auto z = sg.load(v + i);
6767

6868
sg.store(&global[i], x + y);

0 commit comments

Comments
 (0)