diff --git a/sycl/include/sycl/detail/builtins/helper_macros.hpp b/sycl/include/sycl/detail/builtins/helper_macros.hpp index 6469f35da3e6c..1f83870c5b70c 100644 --- a/sycl/include/sycl/detail/builtins/helper_macros.hpp +++ b/sycl/include/sycl/detail/builtins/helper_macros.hpp @@ -197,7 +197,8 @@ [](NUM_ARGS##_AUTO_ARG) { return (NS::NAME)(NUM_ARGS##_ARG); }, \ NUM_ARGS##_ARG); \ } else { \ - return __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG); \ + return bit_cast>( \ + __VA_ARGS__(NUM_ARGS##_CONVERTED_ARG)); \ } \ } diff --git a/sycl/include/sycl/detail/builtins/math_functions.inc b/sycl/include/sycl/detail/builtins/math_functions.inc index d0e9931dc11d6..72a965c47e274 100644 --- a/sycl/include/sycl/detail/builtins/math_functions.inc +++ b/sycl/include/sycl/detail/builtins/math_functions.inc @@ -254,7 +254,9 @@ auto builtin_delegate_ptr_impl(FuncTy F, PtrTy p, Ts... xs) { detail::NON_SCALAR_ENABLER \ NAME(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _TEMPLATE_TYPE_ARG), PtrTy p) { \ - return detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p); \ + return bit_cast>( \ + detail::NAME##_impl(SYCL_CONCAT(LESS_ONE(NUM_ARGS), _ARG), p)); \ } #if __SYCL_DEVICE_ONLY__ diff --git a/sycl/include/sycl/detail/builtins/relational_functions.inc b/sycl/include/sycl/detail/builtins/relational_functions.inc index e958e5c510425..fea8f59e2993f 100644 --- a/sycl/include/sycl/detail/builtins/relational_functions.inc +++ b/sycl/include/sycl/detail/builtins/relational_functions.inc @@ -63,7 +63,7 @@ auto builtin_device_rel_impl(FuncTy F, const Ts &...xs) { // the relation builtin (vector of int16_t/int32_t/int64_t depending on the // arguments' element type). auto ret = F(builtins::convert_arg(xs)...); - vec::value> tmp{ret}; + auto tmp = bit_cast::value>>(ret); using res_elem_type = fixed_width_signed)>; static_assert(is_scalar_arithmetic_v); return tmp.template convert(); diff --git a/sycl/include/sycl/detail/spirv.hpp b/sycl/include/sycl/detail/spirv.hpp index 4b25ec22393b5..541668ed374d9 100644 --- a/sycl/include/sycl/detail/spirv.hpp +++ b/sycl/include/sycl/detail/spirv.hpp @@ -946,11 +946,12 @@ EnableIfNativeShuffle Shuffle(GroupT g, T x, id<1> local_id) { return result; } else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v< GroupT>) { - return __spirv_GroupNonUniformShuffle(group_scope::value, - convertToOpenCLType(x), LocalId); + return convertFromOpenCLTypeFor(__spirv_GroupNonUniformShuffle( + group_scope::value, convertToOpenCLType(x), LocalId)); } else { // Subgroup. - return __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId); + return convertFromOpenCLTypeFor( + __spirv_SubgroupShuffleINTEL(convertToOpenCLType(x), LocalId)); } #else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v< @@ -987,8 +988,8 @@ EnableIfNativeShuffle ShuffleXor(GroupT g, T x, id<1> mask) { convertToOpenCLType(x), TargetId); } else { // Subgroup. - return __spirv_SubgroupShuffleXorINTEL(convertToOpenCLType(x), - static_cast(mask.get(0))); + return convertFromOpenCLTypeFor(__spirv_SubgroupShuffleXorINTEL( + convertToOpenCLType(x), static_cast(mask.get(0)))); } #else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v< @@ -1035,8 +1036,8 @@ EnableIfNativeShuffle ShuffleDown(GroupT g, T x, uint32_t delta) { convertToOpenCLType(x), TargetId); } else { // Subgroup. - return __spirv_SubgroupShuffleDownINTEL(convertToOpenCLType(x), - convertToOpenCLType(x), delta); + return convertFromOpenCLTypeFor(__spirv_SubgroupShuffleDownINTEL( + convertToOpenCLType(x), convertToOpenCLType(x), delta)); } #else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v< @@ -1079,8 +1080,8 @@ EnableIfNativeShuffle ShuffleUp(GroupT g, T x, uint32_t delta) { convertToOpenCLType(x), TargetId); } else { // Subgroup. - return __spirv_SubgroupShuffleUpINTEL(convertToOpenCLType(x), - convertToOpenCLType(x), delta); + return convertFromOpenCLTypeFor(__spirv_SubgroupShuffleUpINTEL( + convertToOpenCLType(x), convertToOpenCLType(x), delta)); } #else if constexpr (ext::oneapi::experimental::is_user_constructed_group_v< diff --git a/sycl/include/sycl/detail/vector_convert.hpp b/sycl/include/sycl/detail/vector_convert.hpp index bd414b01510a8..d66f869885efc 100644 --- a/sycl/include/sycl/detail/vector_convert.hpp +++ b/sycl/include/sycl/detail/vector_convert.hpp @@ -882,6 +882,36 @@ using ConvertBoolAndByteT = template template vec vec::convert() const { +#if !__SYCL_USE_LIBSYCL8_VEC_IMPL + auto getValue = [this](int Index) { + using RetType = typename std::conditional_t< + detail::is_byte_v, int8_t, +#ifdef __SYCL_DEVICE_ONLY__ + typename detail::map_type< + DataT, +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) + std::byte, /*->*/ std::uint8_t, // +#endif + bool, /*->*/ std::uint8_t, // + sycl::half, /*->*/ sycl::detail::half_impl::StorageT, // + sycl::ext::oneapi::bfloat16, + /*->*/ sycl::ext::oneapi::bfloat16::Bfloat16StorageT, // + char, /*->*/ detail::ConvertToOpenCLType_t, // + DataT, /*->*/ DataT // + >::type +#else + DataT +#endif + >; + +#ifdef __SYCL_DEVICE_ONLY__ + if constexpr (std::is_same_v) + return sycl::bit_cast(this->m_Data[Index]); + else +#endif + return static_cast(this->m_Data[Index]); + }; +#endif using T = detail::ConvertBoolAndByteT; using R = detail::ConvertBoolAndByteT; using bfloat16 = sycl::ext::oneapi::bfloat16; diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp index 0f274566ebd91..5671dfb514b1d 100644 --- a/sycl/include/sycl/vector.hpp +++ b/sycl/include/sycl/vector.hpp @@ -423,9 +423,11 @@ class __SYCL_EBO Swizzle using element_type = DataT; using value_type = DataT; +#if __SYCL_USE_LIBSYCL8_VEC_IMPL #ifdef __SYCL_DEVICE_ONLY__ using vector_t = typename vec::vector_t; #endif // __SYCL_DEVICE_ONLY__ +#endif Swizzle() = delete; Swizzle(const Swizzle &) = delete; @@ -497,6 +499,7 @@ class __SYCL_EBO vec : using Base = detail::vec_base; +#if __SYCL_USE_LIBSYCL8_VEC_IMPL #ifdef __SYCL_DEVICE_ONLY__ using element_type_for_vector_t = typename detail::map_type< DataT, @@ -541,6 +544,7 @@ class __SYCL_EBO vec : private: #endif // __SYCL_DEVICE_ONLY__ +#endif #if __SYCL_USE_LIBSYCL8_VEC_IMPL template @@ -618,6 +622,7 @@ class __SYCL_EBO vec : static constexpr size_t get_size() { return byte_size(); } static constexpr size_t byte_size() noexcept { return sizeof(Base); } +#if __SYCL_USE_LIBSYCL8_VEC_IMPL private: // getValue should be able to operate on different underlying // types: enum cl_float#N , builtin vector float#N, builtin type float. @@ -640,6 +645,8 @@ class __SYCL_EBO vec : } public: +#endif + // Out-of-class definition is in `sycl/detail/vector_convert.hpp` template diff --git a/sycl/test-e2e/DeviceLib/built-ins/printf.cpp b/sycl/test-e2e/DeviceLib/built-ins/printf.cpp index d7015109d1704..a0e7bff0d939e 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/printf.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/printf.cpp @@ -59,10 +59,11 @@ int main() { sycl::vec v4{5, 6, 7, 8}; #if defined(__SYCL_DEVICE_ONLY__) && (defined(__SPIR__) || defined(__SPIRV__)) // On SPIRV devices, vectors can be printed via native OpenCL types: - using ocl_int4 = sycl::vec::vector_t; + using ocl_int4 = int __attribute__((ext_vector_type(4))); { static const CONSTANT char format[] = "%v4hld\n"; - ext::oneapi::experimental::printf(format, (ocl_int4)v4); + ext::oneapi::experimental::printf(format, + sycl::bit_cast(v4)); } // However, you are still able to print them by-element: diff --git a/sycl/test/basic_tests/generic_type_traits.cpp b/sycl/test/basic_tests/generic_type_traits.cpp index 3586d25cab588..0d29d635caba1 100644 --- a/sycl/test/basic_tests/generic_type_traits.cpp +++ b/sycl/test/basic_tests/generic_type_traits.cpp @@ -134,11 +134,19 @@ int main() { #endif #ifdef __SYCL_DEVICE_ONLY__ + static_assert( + std::is_same_v>, + s::opencl::cl_int __attribute__((ext_vector_type(2)))>); + static_assert( + std::is_same_v>, + s::opencl::cl_long __attribute__((ext_vector_type(2)))>); +#if __SYCL_USE_LIBSYCL8_VEC_IMPL static_assert( std::is_same_v>, s::vec::vector_t>); static_assert(std::is_same_v>, s::vec::vector_t>); +#endif static_assert(std::is_same_v< d::ConvertToOpenCLType_t::pointer>); + static_assert( + std::is_same_v< + d::ConvertToOpenCLType_t< + s::multi_ptr, + s::access::address_space::global_space, + s::access::decorated::yes>>, + s::multi_ptr::pointer>); +#if __SYCL_USE_LIBSYCL8_VEC_IMPL static_assert( std::is_same_v, @@ -154,6 +172,7 @@ int main() { s::multi_ptr::vector_t, s::access::address_space::global_space, s::access::decorated::yes>::pointer>); +#endif #endif static_assert(std::is_same_v, d::half_impl::BIsRepresentationT>); diff --git a/sycl/test/basic_tests/vectors/assign.cpp b/sycl/test/basic_tests/vectors/assign.cpp index 115bb9af9f2dd..df02b9420d7a4 100644 --- a/sycl/test/basic_tests/vectors/assign.cpp +++ b/sycl/test/basic_tests/vectors/assign.cpp @@ -56,7 +56,7 @@ static_assert( std::is_assignable_v, half>); static_assert( std::is_assignable_v, float>); static_assert( std::is_assignable_v, double>); #if __SYCL_DEVICE_ONLY__ -static_assert( std::is_assignable_v, vec>); +static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v, vec>); #else static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v, vec>); #endif @@ -73,7 +73,7 @@ static_assert( std::is_assignable_v, half>); static_assert( std::is_assignable_v, float>); static_assert( std::is_assignable_v, double>); #if __SYCL_DEVICE_ONLY__ -static_assert( std::is_assignable_v, vec>); +static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v, vec>); #else static_assert( !std::is_assignable_v, vec>); #endif diff --git a/sycl/test/basic_tests/vectors/swizzle_aliases.cpp b/sycl/test/basic_tests/vectors/swizzle_aliases.cpp index 6b331c9d4a9bf..0078370f7ed00 100644 --- a/sycl/test/basic_tests/vectors/swizzle_aliases.cpp +++ b/sycl/test/basic_tests/vectors/swizzle_aliases.cpp @@ -8,10 +8,12 @@ int main() { sycl::vec X{1}; static_assert(std::is_same_v())::element_type, int>); static_assert(std::is_same_v())::value_type, int>); +#if __SYCL_USE_LIBSYCL8_VEC_IMPL #ifdef __SYCL_DEVICE_ONLY__ static_assert(std::is_same_v())::vector_t, sycl::vec::vector_t>); #endif // __SYCL_DEVICE_ONLY__ +#endif }); return 0; }