diff --git a/sycl/include/sycl/detail/builtins/builtins.hpp b/sycl/include/sycl/detail/builtins/builtins.hpp index 7546c0f0c1dac..10987b7e70708 100644 --- a/sycl/include/sycl/detail/builtins/builtins.hpp +++ b/sycl/include/sycl/detail/builtins/builtins.hpp @@ -93,8 +93,7 @@ template auto convert_arg(T &&x) { using result_type = std::conditional_t; - // TODO: We should have this bit_cast impl inside vec::convert. - return bit_cast(static_cast(x)); + return bit_cast(x); } else if constexpr (is_swizzle_v) { return convert_arg(simplify_if_swizzle_t{x}); } else { @@ -104,14 +103,6 @@ template auto convert_arg(T &&x) { return convertToOpenCLType(std::forward(x)); } } - -template auto convert_result(T &&x) { - if constexpr (is_vec_v) { - return bit_cast(x); - } else { - return std::forward(x); - } -} #endif } // namespace builtins diff --git a/sycl/include/sycl/detail/builtins/integer_functions.inc b/sycl/include/sycl/detail/builtins/integer_functions.inc index c2023d2ee26dd..65b91160ac260 100644 --- a/sycl/include/sycl/detail/builtins/integer_functions.inc +++ b/sycl/include/sycl/detail/builtins/integer_functions.inc @@ -38,8 +38,7 @@ BUILTIN_CREATE_ENABLER(builtin_enable_suint32, default_ret_type, NUM_ARGS, NAME, builtin_enable_integer_t, [](auto... xs) { \ using ret_ty = \ detail::builtin_enable_integer_t; \ - return detail::builtins::convert_result( \ - __spirv_ocl_##NAME(xs...)); \ + return bit_cast(__spirv_ocl_##NAME(xs...)); \ }) #else #define BUILTIN_GENINT(NUM_ARGS, NAME) \ @@ -54,11 +53,10 @@ BUILTIN_CREATE_ENABLER(builtin_enable_suint32, default_ret_type, NUM_ARGS, NAME, builtin_enable_integer_t, [](auto... xs) { \ using ret_ty = \ detail::builtin_enable_integer_t; \ - using detail::builtins::convert_result; \ if constexpr (std::is_signed_v>) \ - return convert_result(__spirv_ocl_s_##NAME(xs...)); \ + return bit_cast(__spirv_ocl_s_##NAME(xs...)); \ else \ - return convert_result(__spirv_ocl_u_##NAME(xs...)); \ + return bit_cast(__spirv_ocl_u_##NAME(xs...)); \ }) #else #define BUILTIN_GENINT_SU(NUM_ARGS, NAME) BUILTIN_GENINT(NUM_ARGS, NAME) @@ -67,15 +65,14 @@ BUILTIN_CREATE_ENABLER(builtin_enable_suint32, default_ret_type, #if __SYCL_DEVICE_ONLY__ DEVICE_IMPL_TEMPLATE(ONE_ARG, abs, builtin_enable_integer_t, [](auto x) { using ret_ty = detail::builtin_enable_integer_t; - using detail::builtins::convert_result; if constexpr (std::is_signed_v>) // SPIR-V builtin returns unsigned type, SYCL's return type is signed // with the following restriction: // > The behavior is undefined if the result cannot be represented by // > the return type - return convert_result(bit_cast(__spirv_ocl_s_abs(x))); + return bit_cast(__spirv_ocl_s_abs(x)); else - return convert_result(__spirv_ocl_u_abs(x)); + return bit_cast(__spirv_ocl_u_abs(x)); }) #else BUILTIN_GENINT_SU(ONE_ARG, abs) @@ -87,25 +84,10 @@ BUILTIN_GENINT_SU(TWO_ARGS, add_sat) DEVICE_IMPL_TEMPLATE( TWO_ARGS, abs_diff, builtin_enable_integer_t, [](auto... xs) { using ret_ty = detail::builtin_enable_integer_t; - using detail::builtins::convert_result; if constexpr (std::is_signed_v>) { - // SPIRV built-in returns [vector of] unsigned type(s). - auto ret = __spirv_ocl_s_abs_diff(xs...); - if constexpr (detail::is_vec_v) { - // SYCL 2020 revision 8's abs_diff returns T0 (or corresponding vec in - // case of a swizzle). The only way to produce signed ext_vector_type - // from unsigned is with C-style case. Also note that element type of - // sycl::vec and ext_vector_type might be different, e.g. - // sycl::vec::vector_t is - // signed char __attribute__((ext_vector_type(N))). - // - // TODO: Shouldn't be different from "abs" above. - return convert_result((typename T0::vector_t)(ret)); - } else { - return convert_result(ret); - } + return bit_cast(__spirv_ocl_s_abs_diff(xs...)); } else { - return convert_result(__spirv_ocl_u_abs_diff(xs...)); + return bit_cast(__spirv_ocl_u_abs_diff(xs...)); } }) #else diff --git a/sycl/include/sycl/detail/builtins/relational_functions.inc b/sycl/include/sycl/detail/builtins/relational_functions.inc index 2fa96c378f970..a765951430bc0 100644 --- a/sycl/include/sycl/detail/builtins/relational_functions.inc +++ b/sycl/include/sycl/detail/builtins/relational_functions.inc @@ -133,8 +133,7 @@ DEVICE_IMPL_TEMPLATE( THREE_ARGS, bitselect, builtin_enable_bitselect_t, [](auto... xs) { using ret_ty = detail::builtin_enable_bitselect_t; - using detail::builtins::convert_result; - return convert_result(__spirv_ocl_bitselect(xs...)); + return bit_cast(__spirv_ocl_bitselect(xs...)); }) #else HOST_IMPL_TEMPLATE(THREE_ARGS, bitselect, builtin_enable_bitselect_t, rel,