[SYCL][ESIMD] Add support for different types for lsc functions #6952
Conversation
|
Complementary test PR intel/llvm-test-suite#1305 |
| using type = | ||
| typename std::conditional<sizeof(T) <= 4, int32_t, int64_t>::type; |
There was a problem hiding this comment.
Minor comment: this can be shortened to 1 line:
| using type = | |
| typename std::conditional<sizeof(T) <= 4, int32_t, int64_t>::type; | |
| using type = std::conditional_t<sizeof(T) <= 4, uint32_t, uint64_t>; |
There was a problem hiding this comment.
without typename it gives a compilation error for some reason
There was a problem hiding this comment.
it fails if only remove 'typename', but it should pass if also add "_t" to std::conditional
| public: | ||
| using type = | ||
| typename std::conditional<sizeof(_type2) == 1, _type2, _type1>::type; | ||
| using type = typename std::conditional_t< |
There was a problem hiding this comment.
L194: 'typename' is not needed with std::conditional_t
| using type = typename std::conditional_t< | ||
| sizeof(T) == 1, int8_t, | ||
| std::conditional_t< | ||
| sizeof(T) == 2, int16_t, | ||
| std::conditional_t<sizeof(T) == 4, int32_t, | ||
| std::conditional_t<sizeof(T) == 8, int64_t, T>>>>; |
There was a problem hiding this comment.
Was the switch from unsigned types to signed intentional? I see in other places unsigned type is still used.
If either (signed or unsigned) works fine, then you can just re-use the existing __ESIMD_DNS::uint_type_t
| using type = typename std::conditional_t< | |
| sizeof(T) == 1, int8_t, | |
| std::conditional_t< | |
| sizeof(T) == 2, int16_t, | |
| std::conditional_t<sizeof(T) == 4, int32_t, | |
| std::conditional_t<sizeof(T) == 8, int64_t, T>>>>; | |
| using type = __ESIMD_DNS::uint_type_t<sizeof(T)>; |
| @@ -449,7 +449,8 @@ __ESIMD_API std::enable_if_t<!std::is_pointer<AccessorTy>::value, | |||
| lsc_gather(AccessorTy acc, __ESIMD_NS::simd<uint32_t, N> offsets, | |||
| __ESIMD_NS::simd_mask<N> pred = 1) { | |||
| #ifdef __ESIMD_FORCE_STATELESS_MEM | |||
| return lsc_gather<T, N, DS, L1H>(acc.get_pointer().get(), offsets, pred); | |||
| return lsc_gather<T, NElts, DS, L1H, L3H>(acc.get_pointer().get(), offsets, | |||
There was a problem hiding this comment.
Good catch, thank you for fixing this!
| __esimd_lsc_load_bti<uint32_t, L1H, L3H, _AddressScale, _ImmOffset, | ||
| lsc_data_size::u32, _VS, _Transposed, N>( | ||
| pred.data(), offsets.data(), si); | ||
| return result.template bit_cast_view<T>(); |
There was a problem hiding this comment.
@kbobrovs - accordingly to your experience, is this bit_cast_view() always optimized away if the type is already integral?
If not, then it probably makes sense to write something like:
if constexpr (std::is_integral_v<T>) {
return __esimd_lsc_load_bit<....>(....);
} else {
result = __esimd_lsc_load_bit<uint..., >(....);
return result.template bit_cast_view<T>();
}
What is your recommendation here?
There was a problem hiding this comment.
It is optimized, but I'd prefer sycl::bit_cast instead in such cases.
| __esimd_lsc_load_stateless<uint32_t, L1H, L3H, _AddressScale, | ||
| _ImmOffset, lsc_data_size::u32, _VS, | ||
| _Transposed, N>(pred.data(), addrs.data()); | ||
| return result.template bit_cast_view<T>(); |
There was a problem hiding this comment.
Nit: I suggest to use sycl::bit_cast when view is not really needed (like in this case)
There was a problem hiding this comment.
Approved taking into account that the functionality is needed and also the accidental/wrong merge of intel/llvm-test-suite#1305
Please though make additional effort trying using sycl::bit_cast instead of bit_cast_view() in follow-up PR.
No description provided.