From e2075c714ee6125f34bc81d81ff8fcf180155c15 Mon Sep 17 00:00:00 2001 From: aelovikov-intel Date: Mon, 14 Oct 2024 11:24:02 -0700 Subject: [PATCH] [SYCL][NFCI] Refactor traits for fixed-width integer types selection (#15688) --- .../detail/builtins/relational_functions.inc | 3 +- .../sycl/detail/generic_type_traits.hpp | 49 +++++-------------- sycl/include/sycl/detail/spirv.hpp | 12 ++--- sycl/include/sycl/detail/type_traits.hpp | 13 ----- sycl/include/sycl/detail/vector_arith.hpp | 2 +- .../oneapi/experimental/group_load_store.hpp | 2 +- sycl/include/sycl/sub_group.hpp | 2 +- sycl/include/sycl/vector.hpp | 5 +- sycl/source/builtins/relational_functions.cpp | 5 +- sycl/test/type_traits/type_traits.cpp | 15 ------ 10 files changed, 26 insertions(+), 82 deletions(-) diff --git a/sycl/include/sycl/detail/builtins/relational_functions.inc b/sycl/include/sycl/detail/builtins/relational_functions.inc index db1f43a0b4998..2fa96c378f970 100644 --- a/sycl/include/sycl/detail/builtins/relational_functions.inc +++ b/sycl/include/sycl/detail/builtins/relational_functions.inc @@ -48,8 +48,7 @@ auto builtin_device_rel_impl(FuncTy F, const Ts &...xs) { // arguments' element type). auto ret = F(builtins::convert_arg(xs)...); vec::value> tmp{ret}; - using res_elem_type = - make_type_t, type_list>; + using res_elem_type = fixed_width_signed)>; static_assert(is_scalar_arithmetic_v); return tmp.template convert(); } else if constexpr (std::is_same_v) { diff --git a/sycl/include/sycl/detail/generic_type_traits.hpp b/sycl/include/sycl/detail/generic_type_traits.hpp index 89d1dd6f3a0ad..aafef77a12067 100644 --- a/sycl/include/sycl/detail/generic_type_traits.hpp +++ b/sycl/include/sycl/detail/generic_type_traits.hpp @@ -99,41 +99,19 @@ using is_byte = typename template inline constexpr bool is_byte_v = is_byte::value; -template -using make_floating_point_t = make_type_t; - -template -using make_singed_integer_t = make_type_t; - -template -using make_unsinged_integer_t = - make_type_t; - template -using cl_unsigned = std::conditional_t< - Size == 1, opencl::cl_uchar, +using fixed_width_unsigned = std::conditional_t< + Size == 1, uint8_t, std::conditional_t< - Size == 2, opencl::cl_ushort, - std::conditional_t>>; - -// select_apply_cl_scalar_t selects from T8/T16/T32/T64 basing on -// sizeof(IN). expected to handle scalar types. -template -using select_apply_cl_scalar_t = std::conditional_t< - sizeof(T) == 1, T8, - std::conditional_t>>; - -// Shortcuts for selecting scalar int/unsigned int/fp type. -template -using select_cl_scalar_integral_signed_t = - select_apply_cl_scalar_t; + Size == 2, uint16_t, + std::conditional_t>>; -template -using select_cl_scalar_integral_unsigned_t = - select_apply_cl_scalar_t; +template +using fixed_width_signed = std::conditional_t< + Size == 1, int8_t, + std::conditional_t< + Size == 2, int16_t, + std::conditional_t>>; // Use SFINAE so that std::complex specialization could be implemented in // include/sycl/stl_wrappers/complex that would only be available if STL's @@ -188,10 +166,9 @@ template auto convertToOpenCLType(T &&x) { return static_cast(x); #endif } else if constexpr (std::is_integral_v) { - using OpenCLType = - std::conditional_t, - select_cl_scalar_integral_signed_t, - select_cl_scalar_integral_unsigned_t>; + using OpenCLType = std::conditional_t, + fixed_width_signed, + fixed_width_unsigned>; static_assert(sizeof(OpenCLType) == sizeof(T)); return static_cast(x); } else if constexpr (is_half_v) { diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp index 98ca92b6c988d..32d87c7562b77 100644 --- a/sycl/include/sycl/detail/spirv.hpp +++ b/sycl/include/sycl/detail/spirv.hpp @@ -255,7 +255,7 @@ using EnableIfBitcastBroadcast = std::enable_if_t< is_bitcast_broadcast::value && std::is_integral::value, T>; template -using ConvertToNativeBroadcastType_t = select_cl_scalar_integral_unsigned_t; +using ConvertToNativeBroadcastType_t = fixed_width_unsigned; // Generic broadcasts may require multiple calls to SPIR-V GroupBroadcast // intrinsics, and should use the fewest broadcasts possible @@ -524,7 +524,7 @@ inline typename std::enable_if_t::value, T> AtomicCompareExchange(multi_ptr MPtr, memory_scope Scope, memory_order Success, memory_order Failure, T Desired, T Expected) { - using I = detail::make_unsinged_integer_t; + using I = detail::fixed_width_unsigned; auto SPIRVSuccess = getMemorySemanticsMask(Success); auto SPIRVFailure = getMemorySemanticsMask(Failure); auto SPIRVScope = getScope(Scope); @@ -552,7 +552,7 @@ template ::value, T> AtomicLoad(multi_ptr MPtr, memory_scope Scope, memory_order Order) { - using I = detail::make_unsinged_integer_t; + using I = detail::fixed_width_unsigned; auto *PtrInt = GetMultiPtrDecoratedAs(MPtr); auto SPIRVOrder = getMemorySemanticsMask(Order); auto SPIRVScope = getScope(Scope); @@ -587,7 +587,7 @@ template ::value> AtomicStore(multi_ptr MPtr, memory_scope Scope, memory_order Order, T Value) { - using I = detail::make_unsinged_integer_t; + using I = detail::fixed_width_unsigned; auto *PtrInt = GetMultiPtrDecoratedAs(MPtr); auto SPIRVOrder = getMemorySemanticsMask(Order); auto SPIRVScope = getScope(Scope); @@ -622,7 +622,7 @@ template ::value, T> AtomicExchange(multi_ptr MPtr, memory_scope Scope, memory_order Order, T Value) { - using I = detail::make_unsinged_integer_t; + using I = detail::fixed_width_unsigned; auto *PtrInt = GetMultiPtrDecoratedAs(MPtr); auto SPIRVOrder = getMemorySemanticsMask(Order); auto SPIRVScope = getScope(Scope); @@ -1129,7 +1129,7 @@ EnableIfNonScalarShuffle ShuffleUp(GroupT g, T x, uint32_t delta) { } template -using ConvertToNativeShuffleType_t = select_cl_scalar_integral_unsigned_t; +using ConvertToNativeShuffleType_t = fixed_width_unsigned; template EnableIfBitcastShuffle Shuffle(GroupT g, T x, id<1> local_id) { diff --git a/sycl/include/sycl/detail/type_traits.hpp b/sycl/include/sycl/detail/type_traits.hpp index 56c8571467483..90790633faa9f 100644 --- a/sycl/include/sycl/detail/type_traits.hpp +++ b/sycl/include/sycl/detail/type_traits.hpp @@ -386,19 +386,6 @@ struct remove_pointer : remove_pointer_impl> {}; template using remove_pointer_t = typename remove_pointer::type; -// make_type_t -template struct make_type_impl { - using type = find_same_size_type_t; -}; - -template struct make_type_impl, TL> { - using scalar_type = typename make_type_impl::type; - using type = vec; -}; - -template -using make_type_t = typename make_type_impl::type; - #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR) template using const_if_const_AS = diff --git a/sycl/include/sycl/detail/vector_arith.hpp b/sycl/include/sycl/detail/vector_arith.hpp index bde6ce270afb5..4267febad8e30 100644 --- a/sycl/include/sycl/detail/vector_arith.hpp +++ b/sycl/include/sycl/detail/vector_arith.hpp @@ -113,7 +113,7 @@ template class vec_arith : public vec_arith_common { protected: using vec_t = vec; - using ocl_t = detail::select_cl_scalar_integral_signed_t; + using ocl_t = detail::fixed_width_signed; template using vec_data = vec_helper; // operator!. diff --git a/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp b/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp index e98e8848f0e59..a4af373753c10 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp @@ -156,7 +156,7 @@ struct BlockTypeInfo> { using BlockInfoTy = BlockInfo; static_assert(BlockInfoTy::has_builtin); - using block_type = detail::cl_unsigned; + using block_type = detail::fixed_width_unsigned; using block_pointer_elem_type = std::conditional_t< std::is_const_v -using SelectBlockT = select_cl_scalar_integral_unsigned_t; +using SelectBlockT = fixed_width_unsigned; template auto convertToBlockPtr(MultiPtrTy MultiPtr) { static_assert(is_multi_ptr_v); diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp index 371108a6eae4f..c8d2273094c60 100644 --- a/sycl/include/sycl/vector.hpp +++ b/sycl/include/sycl/vector.hpp @@ -248,7 +248,7 @@ class __SYCL_EBO vec } // Element type for relational operator return value. - using rel_t = detail::select_cl_scalar_integral_signed_t; + using rel_t = detail::fixed_width_signed; public: // Aliases required by SYCL 2020 to make sycl::vec consistent @@ -492,8 +492,7 @@ template class GetScalarOp { private: DataT m_Data; }; -template -using rel_t = detail::select_cl_scalar_integral_signed_t; +template using rel_t = detail::fixed_width_signed; template struct EqualTo { constexpr rel_t operator()(const T &Lhs, const T &Rhs) const { diff --git a/sycl/source/builtins/relational_functions.cpp b/sycl/source/builtins/relational_functions.cpp index 0119d8e15c5ad..0c7f2f6229153 100644 --- a/sycl/source/builtins/relational_functions.cpp +++ b/sycl/source/builtins/relational_functions.cpp @@ -112,10 +112,7 @@ HOST_IMPL(bitselect, [](auto x, auto y, auto z) { static_assert(std::is_same_v && std::is_same_v && detail::is_scalar_arithmetic_v); - using utype = detail::make_type_t< - T0, detail::type_list>; - static_assert(sizeof(utype) == sizeof(T0)); + using utype = fixed_width_unsigned; bitset bx(bit_cast(x)), by(bit_cast(y)), bz(bit_cast(z)); bitset res = (bz & by) | (~bz & bx); unsigned long long ures = res.to_ullong(); diff --git a/sycl/test/type_traits/type_traits.cpp b/sycl/test/type_traits/type_traits.cpp index be4cb8c9ba670..c01bd07b32239 100644 --- a/sycl/test/type_traits/type_traits.cpp +++ b/sycl/test/type_traits/type_traits.cpp @@ -28,12 +28,6 @@ template void test_is_arithmetic() { static_assert(d::is_arithmetic::value == Expected, ""); } -template -void test_make_type_t() { - static_assert(is_same, CheckedT>::value == Expected, - ""); -} - template void test_change_base_type_t() { static_assert( @@ -102,15 +96,6 @@ int main() { test_is_arithmetic(); test_is_arithmetic(); - test_make_type_t(); - test_make_type_t(); - test_make_type_t, - d::gtl::scalar_unsigned_int_list, - s::vec>(); - test_make_type_t, d::gtl::scalar_float_list, - s::vec>(); - test_change_base_type_t(); test_change_base_type_t(); test_change_base_type_t();