Skip to content

Commit 561c17b

Browse files
wenju-hewenju-he
committed
[SYCL] Move Subgroup*INTEL SPIR-V built-in declarations to clang SPIRVBuiltins.td
Motivation is similar to c040f9a : unify with SPV-IR mangling. SubgroupBlockReadINTEL is not handled in this PR. SPV-IR appends return type to mangled function name. It might be simpler for SYCL header to keep current mangling.
1 parent 2e7528e commit 561c17b

File tree

12 files changed

+791
-637
lines changed

12 files changed

+791
-637
lines changed

‎clang/lib/Sema/SPIRVBuiltins.td‎

Lines changed: 36 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -984,6 +984,42 @@ foreach name = ["GroupLogicalAndKHR", "GroupLogicalOrKHR"] in {
984984
def : SPVBuiltin<name, [Bool, Int, Int, Bool], Attr.Convergent>;
985985
}
986986

987+
def SubgroupShuffleINTELVecType
988+
: GenericType<"SubgroupShuffleINTELVecType",
989+
TypeList<[Char, UChar, Short, UShort, Int, UInt, Float]>,
990+
VecNoScalar>;
991+
992+
foreach name = ["SubgroupShuffleINTEL", "SubgroupShuffleXorINTEL"] in {
993+
def : SPVBuiltin<name, [AGenType1, AGenType1, UInt]>;
994+
def : SPVBuiltin<name, [SubgroupShuffleINTELVecType,
995+
SubgroupShuffleINTELVecType, UInt]>;
996+
}
997+
998+
foreach name = ["SubgroupShuffleUpINTEL", "SubgroupShuffleDownINTEL"] in {
999+
def : SPVBuiltin<name, [AGenType1, AGenType1, AGenType1, UInt]>;
1000+
def : SPVBuiltin<name, [SubgroupShuffleINTELVecType,
1001+
SubgroupShuffleINTELVecType,
1002+
SubgroupShuffleINTELVecType, UInt]>;
1003+
}
1004+
1005+
foreach name = ["SubgroupBlockWriteINTEL"] in {
1006+
foreach AS = [GlobalAS, LocalAS] in {
1007+
foreach Ty = TLUnsignedInts.List in {
1008+
def : SPVBuiltin<name, [Void, PointerType<Ty, AS>, Ty]>;
1009+
}
1010+
foreach Ty = [UChar, UShort] in {
1011+
foreach v = [2, 4, 8, 16] in {
1012+
def : SPVBuiltin<name, [Void, PointerType<Ty, AS>, VectorType<Ty, v>]>;
1013+
}
1014+
}
1015+
foreach Ty = [UInt, ULong] in {
1016+
foreach v = [2, 4, 8] in {
1017+
def : SPVBuiltin<name, [Void, PointerType<Ty, AS>, VectorType<Ty, v>]>;
1018+
}
1019+
}
1020+
}
1021+
}
1022+
9871023
// 3.56.24. Non-Uniform Instructions
9881024

9891025
foreach name = ["GroupNonUniformElect"] in {

‎clang/test/CodeGenSPIRV/spirv-builtin-lookup-intel-subgroup.cpp‎

Lines changed: 252 additions & 0 deletions
Large diffs are not rendered by default.

‎libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl‎

Lines changed: 130 additions & 134 deletions
Large diffs are not rendered by default.

‎libclc/libspirv/lib/amdgcn-amdhsa/misc/sub_group_shuffle.cl‎

Lines changed: 329 additions & 398 deletions
Large diffs are not rendered by default.

‎libdevice/nativecpu_utils.cpp‎

Lines changed: 7 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -73,10 +73,8 @@ DefGenericCastToPtrExpl(ToGlobal, OCL_GLOBAL);
7373
__spirv_SubgroupBlockReadINTEL<Type>(const OCL_GLOBAL PType *Ptr) noexcept { \
7474
return Ptr[__spirv_SubgroupLocalInvocationId()]; \
7575
} \
76-
template <> \
77-
__SYCL_CONVERGENT__ DEVICE_EXTERNAL void \
78-
__spirv_SubgroupBlockWriteINTEL<Type>(PType OCL_GLOBAL * ptr, \
79-
Type v) noexcept { \
76+
__SYCL_CONVERGENT__ DEVICE_EXTERNAL void __spirv_SubgroupBlockWriteINTEL( \
77+
PType OCL_GLOBAL *ptr, Type v) noexcept { \
8078
((Type *)ptr)[__spirv_SubgroupLocalInvocationId()] = v; \
8179
} \
8280
static_assert(true)
@@ -252,18 +250,16 @@ DefineBroadCast(int64_t, i64, int64_t);
252250
#define DefShuffleINTEL(Type, Sfx, MuxType) \
253251
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_##Sfx(MuxType val, \
254252
int32_t lid) noexcept; \
255-
template <> \
256-
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleINTEL<Type>( \
257-
Type val, unsigned id) noexcept { \
253+
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleINTEL(Type val, \
254+
unsigned id) noexcept { \
258255
return (Type)__mux_sub_group_shuffle_##Sfx((MuxType)val, id); \
259256
} \
260257
static_assert(true)
261258

262259
#define DefShuffleUpINTEL(Type, Sfx, MuxType) \
263260
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_up_##Sfx( \
264261
MuxType prev, MuxType curr, int32_t delta) noexcept; \
265-
template <> \
266-
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleUpINTEL<Type>( \
262+
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleUpINTEL( \
267263
Type prev, Type curr, unsigned delta) noexcept { \
268264
return (Type)__mux_sub_group_shuffle_up_##Sfx((MuxType)prev, \
269265
(MuxType)curr, delta); \
@@ -273,8 +269,7 @@ DefineBroadCast(int64_t, i64, int64_t);
273269
#define DefShuffleDownINTEL(Type, Sfx, MuxType) \
274270
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_down_##Sfx( \
275271
MuxType curr, MuxType next, int32_t delta) noexcept; \
276-
template <> \
277-
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleDownINTEL<Type>( \
272+
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleDownINTEL( \
278273
Type curr, Type next, unsigned delta) noexcept { \
279274
return (Type)__mux_sub_group_shuffle_down_##Sfx((MuxType)curr, \
280275
(MuxType)next, delta); \
@@ -284,8 +279,7 @@ DefineBroadCast(int64_t, i64, int64_t);
284279
#define DefShuffleXorINTEL(Type, Sfx, MuxType) \
285280
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_xor_##Sfx(MuxType val, \
286281
int32_t xor_val); \
287-
template <> \
288-
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleXorINTEL<Type>( \
282+
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleXorINTEL( \
289283
Type data, unsigned value) noexcept { \
290284
return (Type)__mux_sub_group_shuffle_xor_##Sfx((MuxType)data, value); \
291285
} \

‎llvm/lib/SYCLNativeCPUUtils/FixABIMuxBuiltinsSYCLNativeCPU.cpp‎

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -37,25 +37,25 @@ PreservedAnalyses FixABIMuxBuiltinsPass::run(Module &M,
3737
return false;
3838
}
3939
Updates.clear();
40-
auto LIDvPos = F.getName().find("ELIDv");
40+
auto LIDvPos = F.getName().find("ELDv");
4141
llvm::StringRef NameToMatch;
4242
if (LIDvPos != llvm::StringRef::npos) {
4343
// Add sizeof ELIDv to get num characters to match against
44-
NameToMatch = F.getName().take_front(LIDvPos + 5);
44+
NameToMatch = F.getName().take_front(LIDvPos + 4);
4545
} else {
4646
return false;
4747
}
4848

4949
unsigned int StartIdx = 0;
5050
unsigned int EndIdx = 1;
51-
if (NameToMatch == "_Z32__spirv_SubgroupShuffleDownINTELIDv") {
51+
if (NameToMatch == "_Z32__spirv_SubgroupShuffleDownINTELDv") {
5252
MuxFuncNameToCall = "__mux_sub_group_shuffle_down_";
53-
} else if (NameToMatch == "_Z30__spirv_SubgroupShuffleUpINTELIDv") {
53+
} else if (NameToMatch == "_Z30__spirv_SubgroupShuffleUpINTELDv") {
5454
MuxFuncNameToCall = "__mux_sub_group_shuffle_up_";
55-
} else if (NameToMatch == "_Z28__spirv_SubgroupShuffleINTELIDv") {
55+
} else if (NameToMatch == "_Z28__spirv_SubgroupShuffleINTELDv") {
5656
MuxFuncNameToCall = "__mux_sub_group_shuffle_";
5757
EndIdx = 0;
58-
} else if (NameToMatch == "_Z31__spirv_SubgroupShuffleXorINTELIDv") {
58+
} else if (NameToMatch == "_Z31__spirv_SubgroupShuffleXorINTELDv") {
5959
MuxFuncNameToCall = "__mux_sub_group_shuffle_xor_";
6060
EndIdx = 0;
6161
} else {

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

Lines changed: 0 additions & 55 deletions
Original file line numberDiff line numberDiff line change
@@ -395,101 +395,46 @@ __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
395395
#undef __SPIRV_ATOMIC_UNSIGNED
396396
#undef __SPIRV_ATOMIC_XOR
397397

398-
template <typename dataT>
399-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
400-
__spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;
401-
template <typename dataT>
402-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
403-
__spirv_SubgroupShuffleDownINTEL(dataT Current, dataT Next,
404-
uint32_t Delta) noexcept;
405-
template <typename dataT>
406-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
407-
__spirv_SubgroupShuffleUpINTEL(dataT Previous, dataT Current,
408-
uint32_t Delta) noexcept;
409-
template <typename dataT>
410-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
411-
__spirv_SubgroupShuffleXorINTEL(dataT Data, uint32_t Value) noexcept;
412-
413398
template <typename dataT>
414399
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
415400
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
416401
uint8_t *Ptr) noexcept;
417402

418-
template <typename dataT>
419-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
420-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint8_t *Ptr,
421-
dataT Data) noexcept;
422-
423403
template <typename dataT>
424404
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
425405
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
426406
uint16_t *Ptr) noexcept;
427407

428-
template <typename dataT>
429-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
430-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint16_t *Ptr,
431-
dataT Data) noexcept;
432-
433408
template <typename dataT>
434409
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
435410
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
436411
uint32_t *Ptr) noexcept;
437412

438-
template <typename dataT>
439-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
440-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint32_t *Ptr,
441-
dataT Data) noexcept;
442-
443413
template <typename dataT>
444414
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
445415
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
446416
uint64_t *Ptr) noexcept;
447417

448-
template <typename dataT>
449-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
450-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint64_t *Ptr,
451-
dataT Data) noexcept;
452-
453418
template <typename dataT>
454419
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
455420
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
456421
uint8_t *Ptr) noexcept;
457422

458-
template <typename dataT>
459-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
460-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint8_t *Ptr,
461-
dataT Data) noexcept;
462-
463423
template <typename dataT>
464424
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
465425
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
466426
uint16_t *Ptr) noexcept;
467427

468-
template <typename dataT>
469-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
470-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint16_t *Ptr,
471-
dataT Data) noexcept;
472-
473428
template <typename dataT>
474429
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
475430
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
476431
uint32_t *Ptr) noexcept;
477432

478-
template <typename dataT>
479-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
480-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint32_t *Ptr,
481-
dataT Data) noexcept;
482-
483433
template <typename dataT>
484434
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
485435
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
486436
uint64_t *Ptr) noexcept;
487437

488-
template <typename dataT>
489-
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
490-
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint64_t *Ptr,
491-
dataT Data) noexcept;
492-
493438
template <int W, int rW>
494439
extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
495440
__spirv_FixedSqrtINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,

‎sycl/test/check_device_code/group_load_store_alignment.cpp‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -82,7 +82,7 @@ SYCL_EXTERNAL void test_load_with_alignment_hint(sycl::sub_group &sg,
8282
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR3]]
8383
// CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_RKSN_SO_SQ__EXIT:%.*]]
8484
// CHECK-GLOBAL: if.end.i.i:
85-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR3]]
85+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR3]]
8686
// CHECK-GLOBAL-NEXT: br label [[_ZN4SYCL3_V13EXT6ONEAPI12EXPERIMENTAL11GROUP_STOREINS0_9SUB_GROUPEIPU3AS1INS3_10PROPERTIESINS3_6DETAIL20PROPERTIES_TYPE_LISTIJNS3_14PROPERTY_VALUEINS3_18DATA_PLACEMENT_KEYEJST17INTEGRAL_CONSTANTIILI0EEEEENSB_INS3_21CONTIGUOUS_MEMORY_KEYEJEEENSB_INS3_14FULL_GROUP_KEYEJEEEEEEEEEENST9ENABLE_IFIXAAAASR6DETAILE18VERIFY_STORE_TYPESIT0_T1_ESR6DETAILE18IS_GENERIC_GROUP_VIT_E18IS_PROPERTY_LIST_VIT2_EEVE4TYPEESP_RKSN_SO_SQ__EXIT]]
8787
// CHECK-GLOBAL: _ZN4sycl3_V13ext6oneapi12experimental11group_storeINS0_9sub_groupEiPU3AS1iNS3_10propertiesINS3_6detail20properties_type_listIJNS3_14property_valueINS3_18data_placement_keyEJSt17integral_constantIiLi0EEEEENSB_INS3_21contiguous_memory_keyEJEEENSB_INS3_14full_group_keyEJEEEEEEEEEENSt9enable_ifIXaaaasr6detailE18verify_store_typesIT0_T1_Esr6detailE18is_generic_group_vIT_E18is_property_list_vIT2_EEvE4typeESP_RKSN_SO_SQ_.exit:
8888
// CHECK-GLOBAL-NEXT: ret void
@@ -96,7 +96,7 @@ SYCL_EXTERNAL void test_store_without_alignment_hint(sycl::sub_group &sg, int v,
9696
// CHECK-GLOBAL-NEXT: entry:
9797
// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null
9898
// CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]])
99-
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V:%.*]]) #[[ATTR3]]
99+
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V:%.*]]) #[[ATTR3]]
100100
// CHECK-GLOBAL-NEXT: ret void
101101
//
102102
SYCL_EXTERNAL void test_store_with_alignment_hint(sycl::sub_group &sg, int v,

0 commit comments

Comments
 (0)