Skip to content

Commit e2075c7

Browse files
aelovikov-intelaelovikov-intel
authored
[SYCL][NFCI] Refactor traits for fixed-width integer types selection (#15688)
1 parent 9ea0f20 commit e2075c7

File tree

10 files changed

+26
-82
lines changed

10 files changed

+26
-82
lines changed

‎sycl/include/sycl/detail/builtins/relational_functions.inc

Lines changed: 1 addition & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -48,8 +48,7 @@ auto builtin_device_rel_impl(FuncTy F, const Ts &...xs) {
4848
// arguments' element type).
4949
auto ret = F(builtins::convert_arg(xs)...);
5050
vec<signed char, num_elements<T>::value> tmp{ret};
51-
using res_elem_type =
52-
make_type_t<get_elem_type_t<T>, type_list<int16_t, int32_t, int64_t>>;
51+
using res_elem_type = fixed_width_signed<sizeof(get_elem_type_t<T>)>;
5352
static_assert(is_scalar_arithmetic_v<res_elem_type>);
5453
return tmp.template convert<res_elem_type>();
5554
} else if constexpr (std::is_same_v<T, half>) {

‎sycl/include/sycl/detail/generic_type_traits.hpp

Lines changed: 13 additions & 36 deletions
Original file line numberDiff line numberDiff line change
@@ -99,41 +99,19 @@ using is_byte = typename
9999

100100
template <typename T> inline constexpr bool is_byte_v = is_byte<T>::value;
101101

102-
template <typename T>
103-
using make_floating_point_t = make_type_t<T, gtl::scalar_floating_list>;
104-
105-
template <typename T>
106-
using make_singed_integer_t = make_type_t<T, gtl::scalar_signed_integer_list>;
107-
108-
template <typename T>
109-
using make_unsinged_integer_t =
110-
make_type_t<T, gtl::scalar_unsigned_integer_list>;
111-
112102
template <int Size>
113-
using cl_unsigned = std::conditional_t<
114-
Size == 1, opencl::cl_uchar,
103+
using fixed_width_unsigned = std::conditional_t<
104+
Size == 1, uint8_t,
115105
std::conditional_t<
116-
Size == 2, opencl::cl_ushort,
117-
std::conditional_t<Size == 4, opencl::cl_uint, opencl::cl_ulong>>>;
118-
119-
// select_apply_cl_scalar_t selects from T8/T16/T32/T64 basing on
120-
// sizeof(IN). expected to handle scalar types.
121-
template <typename T, typename T8, typename T16, typename T32, typename T64>
122-
using select_apply_cl_scalar_t = std::conditional_t<
123-
sizeof(T) == 1, T8,
124-
std::conditional_t<sizeof(T) == 2, T16,
125-
std::conditional_t<sizeof(T) == 4, T32, T64>>>;
126-
127-
// Shortcuts for selecting scalar int/unsigned int/fp type.
128-
template <typename T>
129-
using select_cl_scalar_integral_signed_t =
130-
select_apply_cl_scalar_t<T, sycl::opencl::cl_char, sycl::opencl::cl_short,
131-
sycl::opencl::cl_int, sycl::opencl::cl_long>;
106+
Size == 2, uint16_t,
107+
std::conditional_t<Size == 4, uint32_t, uint64_t>>>;
132108

133-
template <typename T>
134-
using select_cl_scalar_integral_unsigned_t =
135-
select_apply_cl_scalar_t<T, sycl::opencl::cl_uchar, sycl::opencl::cl_ushort,
136-
sycl::opencl::cl_uint, sycl::opencl::cl_ulong>;
109+
template <int Size>
110+
using fixed_width_signed = std::conditional_t<
111+
Size == 1, int8_t,
112+
std::conditional_t<
113+
Size == 2, int16_t,
114+
std::conditional_t<Size == 4, int32_t, int64_t>>>;
137115

138116
// Use SFINAE so that std::complex specialization could be implemented in
139117
// include/sycl/stl_wrappers/complex that would only be available if STL's
@@ -188,10 +166,9 @@ template <typename T> auto convertToOpenCLType(T &&x) {
188166
return static_cast<uint8_t>(x);
189167
#endif
190168
} else if constexpr (std::is_integral_v<no_ref>) {
191-
using OpenCLType =
192-
std::conditional_t<std::is_signed_v<no_ref>,
193-
select_cl_scalar_integral_signed_t<no_ref>,
194-
select_cl_scalar_integral_unsigned_t<no_ref>>;
169+
using OpenCLType = std::conditional_t<std::is_signed_v<no_ref>,
170+
fixed_width_signed<sizeof(no_ref)>,
171+
fixed_width_unsigned<sizeof(no_ref)>>;
195172
static_assert(sizeof(OpenCLType) == sizeof(T));
196173
return static_cast<OpenCLType>(x);
197174
} else if constexpr (is_half_v<no_ref>) {

‎sycl/include/sycl/detail/spirv.hpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -255,7 +255,7 @@ using EnableIfBitcastBroadcast = std::enable_if_t<
255255
is_bitcast_broadcast<T>::value && std::is_integral<IdT>::value, T>;
256256

257257
template <typename T>
258-
using ConvertToNativeBroadcastType_t = select_cl_scalar_integral_unsigned_t<T>;
258+
using ConvertToNativeBroadcastType_t = fixed_width_unsigned<sizeof(T)>;
259259

260260
// Generic broadcasts may require multiple calls to SPIR-V GroupBroadcast
261261
// 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>
524524
AtomicCompareExchange(multi_ptr<T, AddressSpace, IsDecorated> MPtr,
525525
memory_scope Scope, memory_order Success,
526526
memory_order Failure, T Desired, T Expected) {
527-
using I = detail::make_unsinged_integer_t<T>;
527+
using I = detail::fixed_width_unsigned<sizeof(T)>;
528528
auto SPIRVSuccess = getMemorySemanticsMask(Success);
529529
auto SPIRVFailure = getMemorySemanticsMask(Failure);
530530
auto SPIRVScope = getScope(Scope);
@@ -552,7 +552,7 @@ template <typename T, access::address_space AddressSpace,
552552
inline typename std::enable_if_t<std::is_floating_point<T>::value, T>
553553
AtomicLoad(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
554554
memory_order Order) {
555-
using I = detail::make_unsinged_integer_t<T>;
555+
using I = detail::fixed_width_unsigned<sizeof(T)>;
556556
auto *PtrInt = GetMultiPtrDecoratedAs<I>(MPtr);
557557
auto SPIRVOrder = getMemorySemanticsMask(Order);
558558
auto SPIRVScope = getScope(Scope);
@@ -587,7 +587,7 @@ template <typename T, access::address_space AddressSpace,
587587
inline typename std::enable_if_t<std::is_floating_point<T>::value>
588588
AtomicStore(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
589589
memory_order Order, T Value) {
590-
using I = detail::make_unsinged_integer_t<T>;
590+
using I = detail::fixed_width_unsigned<sizeof(T)>;
591591
auto *PtrInt = GetMultiPtrDecoratedAs<I>(MPtr);
592592
auto SPIRVOrder = getMemorySemanticsMask(Order);
593593
auto SPIRVScope = getScope(Scope);
@@ -622,7 +622,7 @@ template <typename T, access::address_space AddressSpace,
622622
inline typename std::enable_if_t<std::is_floating_point<T>::value, T>
623623
AtomicExchange(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
624624
memory_order Order, T Value) {
625-
using I = detail::make_unsinged_integer_t<T>;
625+
using I = detail::fixed_width_unsigned<sizeof(T)>;
626626
auto *PtrInt = GetMultiPtrDecoratedAs<I>(MPtr);
627627
auto SPIRVOrder = getMemorySemanticsMask(Order);
628628
auto SPIRVScope = getScope(Scope);
@@ -1129,7 +1129,7 @@ EnableIfNonScalarShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
11291129
}
11301130

11311131
template <typename T>
1132-
using ConvertToNativeShuffleType_t = select_cl_scalar_integral_unsigned_t<T>;
1132+
using ConvertToNativeShuffleType_t = fixed_width_unsigned<sizeof(T)>;
11331133

11341134
template <typename GroupT, typename T>
11351135
EnableIfBitcastShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {

‎sycl/include/sycl/detail/type_traits.hpp

Lines changed: 0 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -386,19 +386,6 @@ struct remove_pointer : remove_pointer_impl<std::remove_cv_t<T>> {};
386386

387387
template <typename T> using remove_pointer_t = typename remove_pointer<T>::type;
388388

389-
// make_type_t
390-
template <typename T, typename TL> struct make_type_impl {
391-
using type = find_same_size_type_t<TL, T>;
392-
};
393-
394-
template <typename T, int N, typename TL> struct make_type_impl<vec<T, N>, TL> {
395-
using scalar_type = typename make_type_impl<T, TL>::type;
396-
using type = vec<scalar_type, N>;
397-
};
398-
399-
template <typename T, typename TL>
400-
using make_type_t = typename make_type_impl<T, TL>::type;
401-
402389
#if defined(RESTRICT_WRITE_ACCESS_TO_CONSTANT_PTR)
403390
template <access::address_space AS, class DataT>
404391
using const_if_const_AS =

‎sycl/include/sycl/detail/vector_arith.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -113,7 +113,7 @@ template <typename DataT, int NumElements>
113113
class vec_arith : public vec_arith_common<DataT, NumElements> {
114114
protected:
115115
using vec_t = vec<DataT, NumElements>;
116-
using ocl_t = detail::select_cl_scalar_integral_signed_t<DataT>;
116+
using ocl_t = detail::fixed_width_signed<sizeof(DataT)>;
117117
template <typename T> using vec_data = vec_helper<T>;
118118

119119
// operator!.

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -156,7 +156,7 @@ struct BlockTypeInfo<BlockInfo<IteratorT, ElementsPerWorkItem, Blocked>> {
156156
using BlockInfoTy = BlockInfo<IteratorT, ElementsPerWorkItem, Blocked>;
157157
static_assert(BlockInfoTy::has_builtin);
158158

159-
using block_type = detail::cl_unsigned<BlockInfoTy::block_size>;
159+
using block_type = detail::fixed_width_unsigned<BlockInfoTy::block_size>;
160160

161161
using block_pointer_elem_type = std::conditional_t<
162162
std::is_const_v<std::remove_reference_t<

‎sycl/include/sycl/sub_group.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -34,7 +34,7 @@ namespace sub_group {
3434

3535
// Selects 8, 16, 32, or 64-bit type depending on size of scalar type T.
3636
template <typename T>
37-
using SelectBlockT = select_cl_scalar_integral_unsigned_t<T>;
37+
using SelectBlockT = fixed_width_unsigned<sizeof(T)>;
3838

3939
template <typename MultiPtrTy> auto convertToBlockPtr(MultiPtrTy MultiPtr) {
4040
static_assert(is_multi_ptr_v<MultiPtrTy>);

‎sycl/include/sycl/vector.hpp

Lines changed: 2 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -248,7 +248,7 @@ class __SYCL_EBO vec
248248
}
249249

250250
// Element type for relational operator return value.
251-
using rel_t = detail::select_cl_scalar_integral_signed_t<DataT>;
251+
using rel_t = detail::fixed_width_signed<sizeof(DataT)>;
252252

253253
public:
254254
// Aliases required by SYCL 2020 to make sycl::vec consistent
@@ -492,8 +492,7 @@ template <typename T> class GetScalarOp {
492492
private:
493493
DataT m_Data;
494494
};
495-
template <typename T>
496-
using rel_t = detail::select_cl_scalar_integral_signed_t<T>;
495+
template <typename T> using rel_t = detail::fixed_width_signed<sizeof(T)>;
497496

498497
template <typename T> struct EqualTo {
499498
constexpr rel_t<T> operator()(const T &Lhs, const T &Rhs) const {

‎sycl/source/builtins/relational_functions.cpp

Lines changed: 1 addition & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -112,10 +112,7 @@ HOST_IMPL(bitselect, [](auto x, auto y, auto z) {
112112
static_assert(std::is_same_v<T0, T1> && std::is_same_v<T1, T2> &&
113113
detail::is_scalar_arithmetic_v<T0>);
114114

115-
using utype = detail::make_type_t<
116-
T0, detail::type_list<unsigned char, unsigned short, unsigned int,
117-
unsigned long, unsigned long long>>;
118-
static_assert(sizeof(utype) == sizeof(T0));
115+
using utype = fixed_width_unsigned<sizeof(T0)>;
119116
bitset bx(bit_cast<utype>(x)), by(bit_cast<utype>(y)), bz(bit_cast<utype>(z));
120117
bitset res = (bz & by) | (~bz & bx);
121118
unsigned long long ures = res.to_ullong();

‎sycl/test/type_traits/type_traits.cpp

Lines changed: 0 additions & 15 deletions
Original file line numberDiff line numberDiff line change
@@ -28,12 +28,6 @@ template <typename T, bool Expected = true> void test_is_arithmetic() {
2828
static_assert(d::is_arithmetic<T>::value == Expected, "");
2929
}
3030

31-
template <typename T, typename TL, typename CheckedT, bool Expected = true>
32-
void test_make_type_t() {
33-
static_assert(is_same<d::make_type_t<T, TL>, CheckedT>::value == Expected,
34-
"");
35-
}
36-
3731
template <typename T, typename T2, typename CheckedT, bool Expected = true>
3832
void test_change_base_type_t() {
3933
static_assert(
@@ -102,15 +96,6 @@ int main() {
10296
test_is_arithmetic<s::half>();
10397
test_is_arithmetic<s::half2>();
10498

105-
test_make_type_t<int, d::gtl::scalar_unsigned_int_list, unsigned int>();
106-
test_make_type_t<s::opencl::cl_int, d::gtl::scalar_float_list,
107-
s::opencl::cl_float>();
108-
test_make_type_t<s::vec<s::opencl::cl_int, 3>,
109-
d::gtl::scalar_unsigned_int_list,
110-
s::vec<s::opencl::cl_uint, 3>>();
111-
test_make_type_t<s::vec<s::opencl::cl_int, 3>, d::gtl::scalar_float_list,
112-
s::vec<s::opencl::cl_float, 3>>();
113-
11499
test_change_base_type_t<int, float, float>();
115100
test_change_base_type_t<s::int2, float, s::float2>();
116101
test_change_base_type_t<long, float, float>();

0 commit comments

Comments
 (0)