diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp index 5a1215dbe2f8a..54bb7d229c372 100644 --- a/sycl/include/sycl/detail/spirv.hpp +++ b/sycl/include/sycl/detail/spirv.hpp @@ -803,8 +803,7 @@ EnableIfGenericShuffle SubgroupShuffleUp(T x, uint32_t delta); template EnableIfNativeShuffle SubgroupShuffle(T x, id<1> local_id) { #ifndef __NVPTX__ - using OCLT = detail::ConvertToOpenCLType_t; - return __spirv_SubgroupShuffleINTEL(OCLT(x), + return __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), static_cast(local_id.get(0))); #else return __nvvm_shfl_sync_idx_i32(membermask(), x, local_id.get(0), 0x1f); @@ -814,9 +813,8 @@ EnableIfNativeShuffle SubgroupShuffle(T x, id<1> local_id) { template EnableIfNativeShuffle SubgroupShuffleXor(T x, id<1> local_id) { #ifndef __NVPTX__ - using OCLT = detail::ConvertToOpenCLType_t; return __spirv_SubgroupShuffleXorINTEL( - OCLT(x), static_cast(local_id.get(0))); + convertToOpenCLType(x), static_cast(local_id.get(0))); #else return __nvvm_shfl_sync_bfly_i32(membermask(), x, local_id.get(0), 0x1f); #endif @@ -825,8 +823,8 @@ EnableIfNativeShuffle SubgroupShuffleXor(T x, id<1> local_id) { template EnableIfNativeShuffle SubgroupShuffleDown(T x, uint32_t delta) { #ifndef __NVPTX__ - using OCLT = detail::ConvertToOpenCLType_t; - return __spirv_SubgroupShuffleDownINTEL(OCLT(x), OCLT(x), delta); + return __spirv_SubgroupShuffleDownINTEL(convertToOpenCLType(x), + convertToOpenCLType(x), delta); #else return __nvvm_shfl_sync_down_i32(membermask(), x, delta, 0x1f); #endif @@ -835,8 +833,8 @@ EnableIfNativeShuffle SubgroupShuffleDown(T x, uint32_t delta) { template EnableIfNativeShuffle SubgroupShuffleUp(T x, uint32_t delta) { #ifndef __NVPTX__ - using OCLT = detail::ConvertToOpenCLType_t; - return __spirv_SubgroupShuffleUpINTEL(OCLT(x), OCLT(x), delta); + return __spirv_SubgroupShuffleUpINTEL(convertToOpenCLType(x), + convertToOpenCLType(x), delta); #else return __nvvm_shfl_sync_up_i32(membermask(), x, delta, 0); #endif diff --git a/sycl/include/sycl/ext/oneapi/experimental/non_uniform_groups.hpp b/sycl/include/sycl/ext/oneapi/experimental/non_uniform_groups.hpp index a2f2bfb321842..af68ce0e10e0f 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/non_uniform_groups.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/non_uniform_groups.hpp @@ -34,11 +34,9 @@ inline sycl::vec ExtractMask(ext::oneapi::sub_group_mask Mask) { // TODO: This may need to be generalized beyond uint32_t for big masks inline uint32_t CallerPositionInMask(ext::oneapi::sub_group_mask Mask) { sycl::vec MemberMask = ExtractMask(Mask); - auto OCLMask = - sycl::detail::ConvertToOpenCLType_t>(MemberMask); return __spirv_GroupNonUniformBallotBitCount( __spv::Scope::Subgroup, (int)__spv::GroupOperation::ExclusiveScan, - OCLMask); + sycl::detail::convertToOpenCLType(MemberMask)); } #endif diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp index bf28a1792ce86..a72c517fe27b9 100644 --- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp +++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp @@ -138,10 +138,9 @@ struct sub_group_mask { for (int i = 0; i < 4; ++i) { MemberMask[i] = TmpMArray[i]; } - auto OCLMask = - sycl::detail::ConvertToOpenCLType_t>(MemberMask); return __spirv_GroupNonUniformBallotBitCount( - __spv::Scope::Subgroup, (int)__spv::GroupOperation::Reduce, OCLMask); + __spv::Scope::Subgroup, (int)__spv::GroupOperation::Reduce, + sycl::detail::convertToOpenCLType(MemberMask)); #else unsigned int count = 0; auto word = (Bits & valuable_bits(bits_num)); diff --git a/sycl/include/sycl/group.hpp b/sycl/include/sycl/group.hpp index 39b1f0a0293c8..ceb0c58dcf99c 100644 --- a/sycl/include/sycl/group.hpp +++ b/sycl/include/sycl/group.hpp @@ -14,7 +14,7 @@ #include // for NDLoop, __SYCL_ASSERT #include // for __SYCL_TYPE #include // for __SYCL2020_DEPRECATED -#include // for ConvertToOpenCLType_t +#include // for convertToOpenCLType #include // for Builder, getSPIRVMemo... #include // for id, range #include // for is_bool, change_base_...