Skip to content

Commit ef97d89

Browse files
aelovikov-intelaelovikov-intel
authored
[SYCL] Remove underspecified vec::vector_t (#17867)
KhronosGroup/SYCL-Docs#676
1 parent fa3d85f commit ef97d89

File tree

10 files changed

+79
-16
lines changed

10 files changed

+79
-16
lines changed

Diff for: ‎sycl/include/sycl/detail/builtins/helper_macros.hpp

+2-1
Original file line numberDiff line numberDiff line change
@@ -197,7 +197,8 @@
197197
[](NUM_ARGS##_AUTO_ARG) { return (NS::NAME)(NUM_ARGS##_ARG); }, \
198198
NUM_ARGS##_ARG); \
199199
} else { \
200-
return __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG); \
200+
return bit_cast<detail::ENABLER<NUM_ARGS##_TEMPLATE_TYPE>>( \
201+
__VA_ARGS__(NUM_ARGS##_CONVERTED_ARG)); \
201202
} \
202203
}
203204

Diff for: ‎sycl/include/sycl/detail/builtins/math_functions.inc

+3-1
Original file line numberDiff line numberDiff line change
@@ -254,7 +254,9 @@ auto builtin_delegate_ptr_impl(FuncTy F, PtrTy p, Ts... xs) {
254254
detail::NON_SCALAR_ENABLER<SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), \
255255
PtrTy> \
256256
NAME(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG), PtrTy p) { \
257-
return detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \
257+
return bit_cast<detail::NON_SCALAR_ENABLER< \
258+
SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE), PtrTy>>( \
259+
detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p)); \
258260
}
259261

260262
#if __SYCL_DEVICE_ONLY__

Diff for: ‎sycl/include/sycl/detail/builtins/relational_functions.inc

+1-1
Original file line numberDiff line numberDiff line change
@@ -63,7 +63,7 @@ auto builtin_device_rel_impl(FuncTy F, const Ts &...xs) {
6363
// the relation builtin (vector of int16_t/int32_t/int64_t depending on the
6464
// arguments' element type).
6565
auto ret = F(builtins::convert_arg(xs)...);
66-
vec<signed char, num_elements<T>::value> tmp{ret};
66+
auto tmp = bit_cast<vec<signed char, num_elements<T>::value>>(ret);
6767
using res_elem_type = fixed_width_signed<sizeof(get_elem_type_t<T>)>;
6868
static_assert(is_scalar_arithmetic_v<res_elem_type>);
6969
return tmp.template convert<res_elem_type>();

Diff for: ‎sycl/include/sycl/detail/spirv.hpp

+10-9
Original file line numberDiff line numberDiff line change
@@ -946,11 +946,12 @@ EnableIfNativeShuffle<T> Shuffle(GroupT g, T x, id<1> local_id) {
946946
return result;
947947
} else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
948948
GroupT>) {
949-
return __spirv_GroupNonUniformShuffle(group_scope<GroupT>::value,
950-
convertToOpenCLType(x), LocalId);
949+
return convertFromOpenCLTypeFor<T>(__spirv_GroupNonUniformShuffle(
950+
group_scope<GroupT>::value, convertToOpenCLType(x), LocalId));
951951
} else {
952952
// Subgroup.
953-
return __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId);
953+
return convertFromOpenCLTypeFor<T>(
954+
__spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId));
954955
}
955956
#else
956957
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -987,8 +988,8 @@ EnableIfNativeShuffle<T> ShuffleXor(GroupT g, T x, id<1> mask) {
987988
convertToOpenCLType(x), TargetId);
988989
} else {
989990
// Subgroup.
990-
return __spirv_SubgroupShuffleXorINTEL(convertToOpenCLType(x),
991-
static_cast<uint32_t>(mask.get(0)));
991+
return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleXorINTEL(
992+
convertToOpenCLType(x), static_cast<uint32_t>(mask.get(0))));
992993
}
993994
#else
994995
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -1035,8 +1036,8 @@ EnableIfNativeShuffle<T> ShuffleDown(GroupT g, T x, uint32_t delta) {
10351036
convertToOpenCLType(x), TargetId);
10361037
} else {
10371038
// Subgroup.
1038-
return __spirv_SubgroupShuffleDownINTEL(convertToOpenCLType(x),
1039-
convertToOpenCLType(x), delta);
1039+
return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleDownINTEL(
1040+
convertToOpenCLType(x), convertToOpenCLType(x), delta));
10401041
}
10411042
#else
10421043
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<
@@ -1079,8 +1080,8 @@ EnableIfNativeShuffle<T> ShuffleUp(GroupT g, T x, uint32_t delta) {
10791080
convertToOpenCLType(x), TargetId);
10801081
} else {
10811082
// Subgroup.
1082-
return __spirv_SubgroupShuffleUpINTEL(convertToOpenCLType(x),
1083-
convertToOpenCLType(x), delta);
1083+
return convertFromOpenCLTypeFor<T>(__spirv_SubgroupShuffleUpINTEL(
1084+
convertToOpenCLType(x), convertToOpenCLType(x), delta));
10841085
}
10851086
#else
10861087
if constexpr (ext::oneapi::experimental::is_user_constructed_group_v<

Diff for: ‎sycl/include/sycl/detail/vector_convert.hpp

+30
Original file line numberDiff line numberDiff line change
@@ -882,6 +882,36 @@ using ConvertBoolAndByteT =
882882
template <typename DataT, int NumElements>
883883
template <typename convertT, rounding_mode roundingMode>
884884
vec<convertT, NumElements> vec<DataT, NumElements>::convert() const {
885+
#if !__SYCL_USE_LIBSYCL8_VEC_IMPL
886+
auto getValue = [this](int Index) {
887+
using RetType = typename std::conditional_t<
888+
detail::is_byte_v<DataT>, int8_t,
889+
#ifdef __SYCL_DEVICE_ONLY__
890+
typename detail::map_type<
891+
DataT,
892+
#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
893+
std::byte, /*->*/ std::uint8_t, //
894+
#endif
895+
bool, /*->*/ std::uint8_t, //
896+
sycl::half, /*->*/ sycl::detail::half_impl::StorageT, //
897+
sycl::ext::oneapi::bfloat16,
898+
/*->*/ sycl::ext::oneapi::bfloat16::Bfloat16StorageT, //
899+
char, /*->*/ detail::ConvertToOpenCLType_t<char>, //
900+
DataT, /*->*/ DataT //
901+
>::type
902+
#else
903+
DataT
904+
#endif
905+
>;
906+
907+
#ifdef __SYCL_DEVICE_ONLY__
908+
if constexpr (std::is_same_v<DataT, sycl::ext::oneapi::bfloat16>)
909+
return sycl::bit_cast<RetType>(this->m_Data[Index]);
910+
else
911+
#endif
912+
return static_cast<RetType>(this->m_Data[Index]);
913+
};
914+
#endif
885915
using T = detail::ConvertBoolAndByteT<DataT>;
886916
using R = detail::ConvertBoolAndByteT<convertT>;
887917
using bfloat16 = sycl::ext::oneapi::bfloat16;

Diff for: ‎sycl/include/sycl/vector.hpp

+7
Original file line numberDiff line numberDiff line change
@@ -423,9 +423,11 @@ class __SYCL_EBO Swizzle
423423
using element_type = DataT;
424424
using value_type = DataT;
425425

426+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
426427
#ifdef __SYCL_DEVICE_ONLY__
427428
using vector_t = typename vec<DataT, NumElements>::vector_t;
428429
#endif // __SYCL_DEVICE_ONLY__
430+
#endif
429431

430432
Swizzle() = delete;
431433
Swizzle(const Swizzle &) = delete;
@@ -497,6 +499,7 @@ class __SYCL_EBO vec :
497499

498500
using Base = detail::vec_base<DataT, NumElements>;
499501

502+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
500503
#ifdef __SYCL_DEVICE_ONLY__
501504
using element_type_for_vector_t = typename detail::map_type<
502505
DataT,
@@ -541,6 +544,7 @@ class __SYCL_EBO vec :
541544

542545
private:
543546
#endif // __SYCL_DEVICE_ONLY__
547+
#endif
544548

545549
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
546550
template <int... Indexes>
@@ -618,6 +622,7 @@ class __SYCL_EBO vec :
618622
static constexpr size_t get_size() { return byte_size(); }
619623
static constexpr size_t byte_size() noexcept { return sizeof(Base); }
620624

625+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
621626
private:
622627
// getValue should be able to operate on different underlying
623628
// types: enum cl_float#N , builtin vector float#N, builtin type float.
@@ -640,6 +645,8 @@ class __SYCL_EBO vec :
640645
}
641646

642647
public:
648+
#endif
649+
643650
// Out-of-class definition is in `sycl/detail/vector_convert.hpp`
644651
template <typename convertT,
645652
rounding_mode roundingMode = rounding_mode::automatic>

Diff for: ‎sycl/test-e2e/DeviceLib/built-ins/printf.cpp

+3-2
Original file line numberDiff line numberDiff line change
@@ -59,10 +59,11 @@ int main() {
5959
sycl::vec<int, 4> v4{5, 6, 7, 8};
6060
#if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__))
6161
// On SPIRV devices, vectors can be printed via native OpenCL types:
62-
using ocl_int4 = sycl::vec<int, 4>::vector_t;
62+
using ocl_int4 = int __attribute__((ext_vector_type(4)));
6363
{
6464
static const CONSTANT char format[] = "%v4hld\n";
65-
ext::oneapi::experimental::printf(format, (ocl_int4)v4);
65+
ext::oneapi::experimental::printf(format,
66+
sycl::bit_cast<ocl_int4>(v4));
6667
}
6768

6869
// However, you are still able to print them by-element:

Diff for: ‎sycl/test/basic_tests/generic_type_traits.cpp

+19
Original file line numberDiff line numberDiff line change
@@ -134,18 +134,36 @@ int main() {
134134
#endif
135135

136136
#ifdef __SYCL_DEVICE_ONLY__
137+
static_assert(
138+
std::is_same_v<d::ConvertToOpenCLType_t<s::vec<s::opencl::cl_int, 2>>,
139+
s::opencl::cl_int __attribute__((ext_vector_type(2)))>);
140+
static_assert(
141+
std::is_same_v<d::ConvertToOpenCLType_t<s::vec<long long, 2>>,
142+
s::opencl::cl_long __attribute__((ext_vector_type(2)))>);
143+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
137144
static_assert(
138145
std::is_same_v<d::ConvertToOpenCLType_t<s::vec<s::opencl::cl_int, 2>>,
139146
s::vec<s::opencl::cl_int, 2>::vector_t>);
140147
static_assert(std::is_same_v<d::ConvertToOpenCLType_t<s::vec<long long, 2>>,
141148
s::vec<s::opencl::cl_long, 2>::vector_t>);
149+
#endif
142150
static_assert(std::is_same_v<
143151
d::ConvertToOpenCLType_t<s::multi_ptr<
144152
s::opencl::cl_int, s::access::address_space::global_space,
145153
s::access::decorated::yes>>,
146154
s::multi_ptr<s::opencl::cl_int,
147155
s::access::address_space::global_space,
148156
s::access::decorated::yes>::pointer>);
157+
static_assert(
158+
std::is_same_v<
159+
d::ConvertToOpenCLType_t<
160+
s::multi_ptr<s::vec<s::opencl::cl_int, 4>,
161+
s::access::address_space::global_space,
162+
s::access::decorated::yes>>,
163+
s::multi_ptr<s::opencl::cl_int __attribute__((ext_vector_type(4))),
164+
s::access::address_space::global_space,
165+
s::access::decorated::yes>::pointer>);
166+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
149167
static_assert(
150168
std::is_same_v<d::ConvertToOpenCLType_t<
151169
s::multi_ptr<s::vec<s::opencl::cl_int, 4>,
@@ -154,6 +172,7 @@ int main() {
154172
s::multi_ptr<s::vec<s::opencl::cl_int, 4>::vector_t,
155173
s::access::address_space::global_space,
156174
s::access::decorated::yes>::pointer>);
175+
#endif
157176
#endif
158177
static_assert(std::is_same_v<d::ConvertToOpenCLType_t<s::half>,
159178
d::half_impl::BIsRepresentationT>);

Diff for: ‎sycl/test/basic_tests/vectors/assign.cpp

+2-2
Original file line numberDiff line numberDiff line change
@@ -56,7 +56,7 @@ static_assert( std::is_assignable_v<vec<float, 1>, half>);
5656
static_assert( std::is_assignable_v<vec<float, 1>, float>);
5757
static_assert( std::is_assignable_v<vec<float, 1>, double>);
5858
#if __SYCL_DEVICE_ONLY__
59-
static_assert( std::is_assignable_v<vec<float, 1>, vec<half, 1>>);
59+
static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<float, 1>, vec<half, 1>>);
6060
#else
6161
static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<float, 1>, vec<half, 1>>);
6262
#endif
@@ -73,7 +73,7 @@ static_assert( std::is_assignable_v<vec<float, 2>, half>);
7373
static_assert( std::is_assignable_v<vec<float, 2>, float>);
7474
static_assert( std::is_assignable_v<vec<float, 2>, double>);
7575
#if __SYCL_DEVICE_ONLY__
76-
static_assert( std::is_assignable_v<vec<float, 2>, vec<half, 1>>);
76+
static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v<vec<float, 2>, vec<half, 1>>);
7777
#else
7878
static_assert( !std::is_assignable_v<vec<float, 2>, vec<half, 1>>);
7979
#endif

Diff for: ‎sycl/test/basic_tests/vectors/swizzle_aliases.cpp

+2
Original file line numberDiff line numberDiff line change
@@ -8,10 +8,12 @@ int main() {
88
sycl::vec<int, 4> X{1};
99
static_assert(std::is_same_v<decltype(X.swizzle<0>())::element_type, int>);
1010
static_assert(std::is_same_v<decltype(X.swizzle<0>())::value_type, int>);
11+
#if __SYCL_USE_LIBSYCL8_VEC_IMPL
1112
#ifdef __SYCL_DEVICE_ONLY__
1213
static_assert(std::is_same_v<decltype(X.swizzle<0>())::vector_t,
1314
sycl::vec<int, 1>::vector_t>);
1415
#endif // __SYCL_DEVICE_ONLY__
16+
#endif
1517
});
1618
return 0;
1719
}

0 commit comments

Comments
 (0)