Skip to content

[SYCL] Move Subgroup*INTEL SPIR-V built-in declarations to clang SPIRVBuiltins.td #17732

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 7 commits into from
Apr 11, 2025
Merged

[SYCL] Move Subgroup*INTEL SPIR-V built-in declarations to clang SPIRVBuiltins.td #17732

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
561c17b
[SYCL] Move Subgroup*INTEL SPIR-V built-in declarations to clang SPIR…
wenju-he Mar 31, 2025
eed108a
refine format
wenju-he Apr 1, 2025
c750ebf
remove SubgroupShuffle*INTEL declarations from libspirv
wenju-he Apr 1, 2025
9aa422a
update test: spir64 -> spirv64
wenju-he Apr 2, 2025
2f5b830
add back old mangling implementation for backward compatibility
wenju-he Apr 2, 2025
c152a66
Merge branch 'sycl' into move-intel-subgroup-builtin-to-SPIRVBuiltins.td
wenju-he Apr 10, 2025
3518a82
clang-format
wenju-he Apr 10, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
36 changes: 36 additions & 0 deletions clang/lib/Sema/SPIRVBuiltins.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -984,6 +984,42 @@ foreach name = ["GroupLogicalAndKHR", "GroupLogicalOrKHR"] in {
def : SPVBuiltin<name, [Bool, Int, Int, Bool], Attr.Convergent>;
}

def SubgroupShuffleINTELVecType
: GenericType<"SubgroupShuffleINTELVecType",
TypeList<[Char, UChar, Short, UShort, Int, UInt, Float]>,
VecNoScalar>;

foreach name = ["SubgroupShuffleINTEL", "SubgroupShuffleXorINTEL"] in {
def : SPVBuiltin<name, [AGenType1, AGenType1, UInt]>;
def : SPVBuiltin<name, [SubgroupShuffleINTELVecType,
SubgroupShuffleINTELVecType, UInt]>;
}

foreach name = ["SubgroupShuffleUpINTEL", "SubgroupShuffleDownINTEL"] in {
def : SPVBuiltin<name, [AGenType1, AGenType1, AGenType1, UInt]>;
def : SPVBuiltin<name, [SubgroupShuffleINTELVecType,
SubgroupShuffleINTELVecType,
SubgroupShuffleINTELVecType, UInt]>;
}

foreach name = ["SubgroupBlockWriteINTEL"] in {
foreach AS = [GlobalAS, LocalAS] in {
foreach Ty = TLUnsignedInts.List in {
def : SPVBuiltin<name, [Void, PointerType<Ty, AS>, Ty]>;
}
foreach Ty = [UChar, UShort] in {
foreach v = [2, 4, 8, 16] in {
def : SPVBuiltin<name, [Void, PointerType<Ty, AS>, VectorType<Ty, v>]>;
}
}
foreach Ty = [UInt, ULong] in {
foreach v = [2, 4, 8] in {
def : SPVBuiltin<name, [Void, PointerType<Ty, AS>, VectorType<Ty, v>]>;
}
}
}
}

// 3.56.24. Non-Uniform Instructions

foreach name = ["GroupNonUniformElect"] in {
Expand Down
252 changes: 252 additions & 0 deletions clang/test/CodeGenSPIRV/spirv-builtin-lookup-intel-subgroup.cpp
View file
Open in desktop

Large diffs are not rendered by default.

264 changes: 130 additions & 134 deletions libclc/libspirv/lib/amdgcn-amdhsa/group/collectives.cl
View file
Open in desktop

Large diffs are not rendered by default.

727 changes: 329 additions & 398 deletions libclc/libspirv/lib/amdgcn-amdhsa/misc/sub_group_shuffle.cl
View file
Open in desktop

Large diffs are not rendered by default.

20 changes: 7 additions & 13 deletions libdevice/nativecpu_utils.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -73,10 +73,8 @@ DefGenericCastToPtrExpl(ToGlobal, OCL_GLOBAL);
__spirv_SubgroupBlockReadINTEL<Type>(const OCL_GLOBAL PType *Ptr) noexcept { \
return Ptr[__spirv_SubgroupLocalInvocationId()]; \
} \
template <> \
__SYCL_CONVERGENT__ DEVICE_EXTERNAL void \
__spirv_SubgroupBlockWriteINTEL<Type>(PType OCL_GLOBAL * ptr, \
Type v) noexcept { \
__SYCL_CONVERGENT__ DEVICE_EXTERNAL void __spirv_SubgroupBlockWriteINTEL( \
PType OCL_GLOBAL *ptr, Type v) noexcept { \
((Type *)ptr)[__spirv_SubgroupLocalInvocationId()] = v; \
} \
static_assert(true)
Expand Down Expand Up @@ -252,18 +250,16 @@ DefineBroadCast(int64_t, i64, int64_t);
#define DefShuffleINTEL(Type, Sfx, MuxType) \
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_##Sfx(MuxType val, \
int32_t lid) noexcept; \
template <> \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleINTEL<Type>( \
Type val, unsigned id) noexcept { \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleINTEL(Type val, \
unsigned id) noexcept { \
return (Type)__mux_sub_group_shuffle_##Sfx((MuxType)val, id); \
} \
static_assert(true)

#define DefShuffleUpINTEL(Type, Sfx, MuxType) \
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_up_##Sfx( \
MuxType prev, MuxType curr, int32_t delta) noexcept; \
template <> \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleUpINTEL<Type>( \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleUpINTEL( \
Type prev, Type curr, unsigned delta) noexcept { \
return (Type)__mux_sub_group_shuffle_up_##Sfx((MuxType)prev, \
(MuxType)curr, delta); \
Expand All @@ -273,8 +269,7 @@ DefineBroadCast(int64_t, i64, int64_t);
#define DefShuffleDownINTEL(Type, Sfx, MuxType) \
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_down_##Sfx( \
MuxType curr, MuxType next, int32_t delta) noexcept; \
template <> \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleDownINTEL<Type>( \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleDownINTEL( \
Type curr, Type next, unsigned delta) noexcept { \
return (Type)__mux_sub_group_shuffle_down_##Sfx((MuxType)curr, \
(MuxType)next, delta); \
Expand All @@ -284,8 +279,7 @@ DefineBroadCast(int64_t, i64, int64_t);
#define DefShuffleXorINTEL(Type, Sfx, MuxType) \
DEVICE_EXTERN_C MuxType __mux_sub_group_shuffle_xor_##Sfx(MuxType val, \
int32_t xor_val); \
template <> \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleXorINTEL<Type>( \
DEVICE_EXTERNAL Type __spirv_SubgroupShuffleXorINTEL( \
Type data, unsigned value) noexcept { \
return (Type)__mux_sub_group_shuffle_xor_##Sfx((MuxType)data, value); \
} \
Expand Down
12 changes: 6 additions & 6 deletions llvm/lib/SYCLNativeCPUUtils/FixABIMuxBuiltinsSYCLNativeCPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -37,25 +37,25 @@ PreservedAnalyses FixABIMuxBuiltinsPass::run(Module &M,
return false;
}
Updates.clear();
auto LIDvPos = F.getName().find("ELIDv");
auto LIDvPos = F.getName().find("ELDv");
llvm::StringRef NameToMatch;
if (LIDvPos != llvm::StringRef::npos) {
// Add sizeof ELIDv to get num characters to match against
NameToMatch = F.getName().take_front(LIDvPos + 5);
NameToMatch = F.getName().take_front(LIDvPos + 4);
} else {
return false;
}

unsigned int StartIdx = 0;
unsigned int EndIdx = 1;
if (NameToMatch == "_Z32__spirv_SubgroupShuffleDownINTELIDv") {
if (NameToMatch == "_Z32__spirv_SubgroupShuffleDownINTELDv") {
MuxFuncNameToCall = "__mux_sub_group_shuffle_down_";
} else if (NameToMatch == "_Z30__spirv_SubgroupShuffleUpINTELIDv") {
} else if (NameToMatch == "_Z30__spirv_SubgroupShuffleUpINTELDv") {
MuxFuncNameToCall = "__mux_sub_group_shuffle_up_";
} else if (NameToMatch == "_Z28__spirv_SubgroupShuffleINTELIDv") {
} else if (NameToMatch == "_Z28__spirv_SubgroupShuffleINTELDv") {
MuxFuncNameToCall = "__mux_sub_group_shuffle_";
EndIdx = 0;
} else if (NameToMatch == "_Z31__spirv_SubgroupShuffleXorINTELIDv") {
} else if (NameToMatch == "_Z31__spirv_SubgroupShuffleXorINTELDv") {
MuxFuncNameToCall = "__mux_sub_group_shuffle_xor_";
EndIdx = 0;
} else {
Expand Down
55 changes: 0 additions & 55 deletions sycl/include/sycl/__spirv/spirv_ops.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -395,101 +395,46 @@ __SPIRV_ATOMICS(__SPIRV_ATOMIC_MINMAX, Max)
#undef __SPIRV_ATOMIC_UNSIGNED
#undef __SPIRV_ATOMIC_XOR

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupShuffleINTEL(dataT Data, uint32_t InvocationId) noexcept;
template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupShuffleDownINTEL(dataT Current, dataT Next,
uint32_t Delta) noexcept;
template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupShuffleUpINTEL(dataT Previous, dataT Current,
uint32_t Delta) noexcept;
template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupShuffleXorINTEL(dataT Data, uint32_t Value) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
uint8_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint8_t *Ptr,
dataT Data) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
uint16_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint16_t *Ptr,
dataT Data) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
uint32_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint32_t *Ptr,
dataT Data) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_global))
uint64_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint64_t *Ptr,
dataT Data) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
uint8_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint8_t *Ptr,
dataT Data) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
uint16_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint16_t *Ptr,
dataT Data) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
uint32_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint32_t *Ptr,
dataT Data) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL dataT
__spirv_SubgroupBlockReadINTEL(const __attribute__((opencl_local))
uint64_t *Ptr) noexcept;

template <typename dataT>
__SYCL_CONVERGENT__ extern __DPCPP_SYCL_EXTERNAL void
__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_local)) uint64_t *Ptr,
dataT Data) noexcept;

template <int W, int rW>
extern __DPCPP_SYCL_EXTERNAL sycl::detail::ap_int<rW>
__spirv_FixedSqrtINTEL(sycl::detail::ap_int<W> a, bool S, int32_t I, int32_t rI,
Expand Down
4 changes: 2 additions & 2 deletions sycl/test/check_device_code/group_load_store_alignment.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -82,7 +82,7 @@ SYCL_EXTERNAL void test_load_with_alignment_hint(sycl::sub_group &sg,
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z22__spirv_ControlBarrieriii(i32 noundef 3, i32 noundef 3, i32 noundef 912) #[[ATTR3]]
// 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:%.*]]
// CHECK-GLOBAL: if.end.i.i:
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR3]]
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V]]) #[[ATTR3]]
// 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]]
// 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:
// CHECK-GLOBAL-NEXT: ret void
Expand All @@ -96,7 +96,7 @@ SYCL_EXTERNAL void test_store_without_alignment_hint(sycl::sub_group &sg, int v,
// CHECK-GLOBAL-NEXT: entry:
// CHECK-GLOBAL-NEXT: [[CMP_I_I_I:%.*]] = icmp ne ptr addrspace(1) [[P:%.*]], null
// CHECK-GLOBAL-NEXT: tail call void @llvm.assume(i1 [[CMP_I_I_I]])
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELIjEvPU3AS1jT_(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V:%.*]]) #[[ATTR3]]
// CHECK-GLOBAL-NEXT: tail call spir_func void @_Z31__spirv_SubgroupBlockWriteINTELPU3AS1jj(ptr addrspace(1) noundef nonnull [[P]], i32 noundef [[V:%.*]]) #[[ATTR3]]
// CHECK-GLOBAL-NEXT: ret void
//
SYCL_EXTERNAL void test_store_with_alignment_hint(sycl::sub_group &sg, int v,
Expand Down
Loading
Loading