[SYCL][NFCI] Refactor traits for fixed-width integer types selection (i…

…ntel#15688)

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<signed char, num_elements<T>::value> tmp{ret};
-    using res_elem_type =
-        make_type_t<get_elem_type_t<T>, type_list<int16_t, int32_t, int64_t>>;
+    using res_elem_type = fixed_width_signed<sizeof(get_elem_type_t<T>)>;
     static_assert(is_scalar_arithmetic_v<res_elem_type>);
     return tmp.template convert<res_elem_type>();
   } else if constexpr (std::is_same_v<T, half>) {

sycl/include/sycl/detail/generic_type_traits.hpp

@@ -99,41 +99,19 @@ using is_byte = typename
 
 template <typename T> inline constexpr bool is_byte_v = is_byte<T>::value;
 
-template <typename T>
-using make_floating_point_t = make_type_t<T, gtl::scalar_floating_list>;
-
-template <typename T>
-using make_singed_integer_t = make_type_t<T, gtl::scalar_signed_integer_list>;
-
-template <typename T>
-using make_unsinged_integer_t =
-    make_type_t<T, gtl::scalar_unsigned_integer_list>;
-
 template <int Size>
-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<Size == 4, opencl::cl_uint, opencl::cl_ulong>>>;
-
-// select_apply_cl_scalar_t selects from T8/T16/T32/T64 basing on
-// sizeof(IN).  expected to handle scalar types.
-template <typename T, typename T8, typename T16, typename T32, typename T64>
-using select_apply_cl_scalar_t = std::conditional_t<
-    sizeof(T) == 1, T8,
-    std::conditional_t<sizeof(T) == 2, T16,
-                       std::conditional_t<sizeof(T) == 4, T32, T64>>>;
-
-// Shortcuts for selecting scalar int/unsigned int/fp type.
-template <typename T>
-using select_cl_scalar_integral_signed_t =
-    select_apply_cl_scalar_t<T, sycl::opencl::cl_char, sycl::opencl::cl_short,
-                             sycl::opencl::cl_int, sycl::opencl::cl_long>;
+        Size == 2, uint16_t,
+        std::conditional_t<Size == 4, uint32_t, uint64_t>>>;
 
-template <typename T>
-using select_cl_scalar_integral_unsigned_t =
-    select_apply_cl_scalar_t<T, sycl::opencl::cl_uchar, sycl::opencl::cl_ushort,
-                             sycl::opencl::cl_uint, sycl::opencl::cl_ulong>;
+template <int Size>
+using fixed_width_signed = std::conditional_t<
+    Size == 1, int8_t,
+    std::conditional_t<
+        Size == 2, int16_t,
+        std::conditional_t<Size == 4, int32_t, int64_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 <typename T> auto convertToOpenCLType(T &&x) {
     return static_cast<uint8_t>(x);
 #endif
   } else if constexpr (std::is_integral_v<no_ref>) {
-    using OpenCLType =
-        std::conditional_t<std::is_signed_v<no_ref>,
-                           select_cl_scalar_integral_signed_t<no_ref>,
-                           select_cl_scalar_integral_unsigned_t<no_ref>>;
+    using OpenCLType = std::conditional_t<std::is_signed_v<no_ref>,
+                                          fixed_width_signed<sizeof(no_ref)>,
+                                          fixed_width_unsigned<sizeof(no_ref)>>;
     static_assert(sizeof(OpenCLType) == sizeof(T));
     return static_cast<OpenCLType>(x);
   } else if constexpr (is_half_v<no_ref>) {

sycl/include/sycl/detail/spirv.hpp

@@ -255,7 +255,7 @@ using EnableIfBitcastBroadcast = std::enable_if_t<
     is_bitcast_broadcast<T>::value && std::is_integral<IdT>::value, T>;
 
 template <typename T>
-using ConvertToNativeBroadcastType_t = select_cl_scalar_integral_unsigned_t<T>;
+using ConvertToNativeBroadcastType_t = fixed_width_unsigned<sizeof(T)>;
 
 // 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<std::is_floating_point<T>::value, T>
 AtomicCompareExchange(multi_ptr<T, AddressSpace, IsDecorated> MPtr,
                       memory_scope Scope, memory_order Success,
                       memory_order Failure, T Desired, T Expected) {
-  using I = detail::make_unsinged_integer_t<T>;
+  using I = detail::fixed_width_unsigned<sizeof(T)>;
   auto SPIRVSuccess = getMemorySemanticsMask(Success);
   auto SPIRVFailure = getMemorySemanticsMask(Failure);
   auto SPIRVScope = getScope(Scope);
@@ -552,7 +552,7 @@ template <typename T, access::address_space AddressSpace,
 inline typename std::enable_if_t<std::is_floating_point<T>::value, T>
 AtomicLoad(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
            memory_order Order) {
-  using I = detail::make_unsinged_integer_t<T>;
+  using I = detail::fixed_width_unsigned<sizeof(T)>;
   auto *PtrInt = GetMultiPtrDecoratedAs<I>(MPtr);
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -587,7 +587,7 @@ template <typename T, access::address_space AddressSpace,
 inline typename std::enable_if_t<std::is_floating_point<T>::value>
 AtomicStore(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
             memory_order Order, T Value) {
-  using I = detail::make_unsinged_integer_t<T>;
+  using I = detail::fixed_width_unsigned<sizeof(T)>;
   auto *PtrInt = GetMultiPtrDecoratedAs<I>(MPtr);
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -622,7 +622,7 @@ template <typename T, access::address_space AddressSpace,
 inline typename std::enable_if_t<std::is_floating_point<T>::value, T>
 AtomicExchange(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
                memory_order Order, T Value) {
-  using I = detail::make_unsinged_integer_t<T>;
+  using I = detail::fixed_width_unsigned<sizeof(T)>;
   auto *PtrInt = GetMultiPtrDecoratedAs<I>(MPtr);
   auto SPIRVOrder = getMemorySemanticsMask(Order);
   auto SPIRVScope = getScope(Scope);
@@ -1129,7 +1129,7 @@ EnableIfNonScalarShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
 }
 
 template <typename T>
-using ConvertToNativeShuffleType_t = select_cl_scalar_integral_unsigned_t<T>;
+using ConvertToNativeShuffleType_t = fixed_width_unsigned<sizeof(T)>;
 
 template <typename GroupT, typename T>
 EnableIfBitcastShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {

sycl/include/sycl/detail/type_traits.hpp

@@ -386,19 +386,6 @@ struct remove_pointer : remove_pointer_impl<std::remove_cv_t<T>> {};
 
 template <typename T> using remove_pointer_t = typename remove_pointer<T>::type;
 
-// make_type_t
-template <typename T, typename TL> struct make_type_impl {
-  using type = find_same_size_type_t<TL, T>;
-};
-
-template <typename T, int N, typename TL> struct make_type_impl<vec<T, N>, TL> {
-  using scalar_type = typename make_type_impl<T, TL>::type;
-  using type = vec<scalar_type, N>;
-};
-
-template <typename T, typename TL>
-using make_type_t = typename make_type_impl<T, TL>::type;
-
 #if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
 template <access::address_space AS, class DataT>
 using const_if_const_AS =

sycl/include/sycl/detail/vector_arith.hpp

@@ -113,7 +113,7 @@ template <typename DataT, int NumElements>
 class vec_arith : public vec_arith_common<DataT, NumElements> {
 protected:
   using vec_t = vec<DataT, NumElements>;
-  using ocl_t = detail::select_cl_scalar_integral_signed_t<DataT>;
+  using ocl_t = detail::fixed_width_signed<sizeof(DataT)>;
   template <typename T> using vec_data = vec_helper<T>;
 
   // operator!.

sycl/include/sycl/ext/oneapi/experimental/group_load_store.hpp

@@ -156,7 +156,7 @@ struct BlockTypeInfo<BlockInfo<IteratorT, ElementsPerWorkItem, Blocked>> {
   using BlockInfoTy = BlockInfo<IteratorT, ElementsPerWorkItem, Blocked>;
   static_assert(BlockInfoTy::has_builtin);
 
-  using block_type = detail::cl_unsigned<BlockInfoTy::block_size>;
+  using block_type = detail::fixed_width_unsigned<BlockInfoTy::block_size>;
 
   using block_pointer_elem_type = std::conditional_t<
       std::is_const_v<std::remove_reference_t<

sycl/include/sycl/sub_group.hpp

@@ -34,7 +34,7 @@ namespace sub_group {
 
 // Selects 8, 16, 32, or 64-bit type depending on size of scalar type T.
 template <typename T>
-using SelectBlockT = select_cl_scalar_integral_unsigned_t<T>;
+using SelectBlockT = fixed_width_unsigned<sizeof(T)>;
 
 template <typename MultiPtrTy> auto convertToBlockPtr(MultiPtrTy MultiPtr) {
   static_assert(is_multi_ptr_v<MultiPtrTy>);

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<DataT>;
+  using rel_t = detail::fixed_width_signed<sizeof(DataT)>;
 
 public:
   // Aliases required by SYCL 2020 to make sycl::vec consistent
@@ -492,8 +492,7 @@ template <typename T> class GetScalarOp {
 private:
   DataT m_Data;
 };
-template <typename T>
-using rel_t = detail::select_cl_scalar_integral_signed_t<T>;
+template <typename T> using rel_t = detail::fixed_width_signed<sizeof(T)>;
 
 template <typename T> struct EqualTo {
   constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {

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<T0, T1> && std::is_same_v<T1, T2> &&
                 detail::is_scalar_arithmetic_v<T0>);
 
-  using utype = detail::make_type_t<
-      T0, detail::type_list<unsigned char, unsigned short, unsigned int,
-                            unsigned long, unsigned long long>>;
-  static_assert(sizeof(utype) == sizeof(T0));
+  using utype = fixed_width_unsigned<sizeof(T0)>;
   bitset bx(bit_cast<utype>(x)), by(bit_cast<utype>(y)), bz(bit_cast<utype>(z));
   bitset res = (bz & by) | (~bz & bx);
   unsigned long long ures = res.to_ullong();

sycl/test/type_traits/type_traits.cpp

@@ -28,12 +28,6 @@ template <typename T, bool Expected = true> void test_is_arithmetic() {
   static_assert(d::is_arithmetic<T>::value == Expected, "");
 }
 
-template <typename T, typename TL, typename CheckedT, bool Expected = true>
-void test_make_type_t() {
-  static_assert(is_same<d::make_type_t<T, TL>, CheckedT>::value == Expected,
-                "");
-}
-
 template <typename T, typename T2, typename CheckedT, bool Expected = true>
 void test_change_base_type_t() {
   static_assert(
@@ -102,15 +96,6 @@ int main() {
   test_is_arithmetic<s::half>();
   test_is_arithmetic<s::half2>();
 
-  test_make_type_t<int, d::gtl::scalar_unsigned_int_list, unsigned int>();
-  test_make_type_t<s::opencl::cl_int, d::gtl::scalar_float_list,
-                   s::opencl::cl_float>();
-  test_make_type_t<s::vec<s::opencl::cl_int, 3>,
-                   d::gtl::scalar_unsigned_int_list,
-                   s::vec<s::opencl::cl_uint, 3>>();
-  test_make_type_t<s::vec<s::opencl::cl_int, 3>, d::gtl::scalar_float_list,
-                   s::vec<s::opencl::cl_float, 3>>();
-
   test_change_base_type_t<int, float, float>();
   test_change_base_type_t<s::int2, float, s::float2>();
   test_change_base_type_t<long, float, float>();