diff --git a/sycl/include/sycl/detail/vector_arith.hpp b/sycl/include/sycl/detail/vector_arith.hpp index 0b24c71454984..6d92213e825e5 100644 --- a/sycl/include/sycl/detail/vector_arith.hpp +++ b/sycl/include/sycl/detail/vector_arith.hpp @@ -66,6 +66,18 @@ struct IncDec {}; template static constexpr bool not_fp = !is_vgenfloat_v; +#if !__SYCL_USE_LIBSYCL8_VEC_IMPL +// Not using `is_byte_v` to avoid unnecessary dependencies on `half`/`bfloat16` +// headers. +template +static constexpr bool not_byte = +#if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) + !std::is_same_v; +#else + true; +#endif +#endif + // To provide information about operators availability depending on vec/swizzle // element type. template @@ -80,6 +92,7 @@ inline constexpr bool is_op_available_for_type, T> = inline constexpr bool is_op_available_for_type = COND; // clang-format off +#if __SYCL_USE_LIBSYCL8_VEC_IMPL __SYCL_OP_AVAILABILITY(std::plus , true) __SYCL_OP_AVAILABILITY(std::minus , true) __SYCL_OP_AVAILABILITY(std::multiplies , true) @@ -110,6 +123,38 @@ __SYCL_OP_AVAILABILITY(std::bit_not , not_fp) __SYCL_OP_AVAILABILITY(UnaryPlus , true) __SYCL_OP_AVAILABILITY(IncDec , true) +#else +__SYCL_OP_AVAILABILITY(std::plus , not_byte) +__SYCL_OP_AVAILABILITY(std::minus , not_byte) +__SYCL_OP_AVAILABILITY(std::multiplies , not_byte) +__SYCL_OP_AVAILABILITY(std::divides , not_byte) +__SYCL_OP_AVAILABILITY(std::modulus , not_fp) + +__SYCL_OP_AVAILABILITY(std::bit_and , not_fp) +__SYCL_OP_AVAILABILITY(std::bit_or , not_fp) +__SYCL_OP_AVAILABILITY(std::bit_xor , not_fp) + +__SYCL_OP_AVAILABILITY(std::equal_to , true) +__SYCL_OP_AVAILABILITY(std::not_equal_to , true) +__SYCL_OP_AVAILABILITY(std::less , true) +__SYCL_OP_AVAILABILITY(std::greater , true) +__SYCL_OP_AVAILABILITY(std::less_equal , true) +__SYCL_OP_AVAILABILITY(std::greater_equal , true) + +__SYCL_OP_AVAILABILITY(std::logical_and , not_byte && not_fp) +__SYCL_OP_AVAILABILITY(std::logical_or , not_byte && not_fp) + +__SYCL_OP_AVAILABILITY(ShiftLeft , not_byte && not_fp) +__SYCL_OP_AVAILABILITY(ShiftRight , not_byte && not_fp) + +// Unary +__SYCL_OP_AVAILABILITY(std::negate , not_byte) +__SYCL_OP_AVAILABILITY(std::logical_not , not_byte) +__SYCL_OP_AVAILABILITY(std::bit_not , not_fp) +__SYCL_OP_AVAILABILITY(UnaryPlus , not_byte) + +__SYCL_OP_AVAILABILITY(IncDec , not_byte) +#endif // clang-format on #undef __SYCL_OP_AVAILABILITY @@ -188,6 +233,12 @@ template struct VecOperators { using element_type = typename from_incomplete::element_type; static constexpr int N = from_incomplete::size(); +#if !__SYCL_USE_LIBSYCL8_VEC_IMPL + template + static constexpr bool is_compatible_scalar = + std::is_convertible_v::element_type>; +#endif + template using result_t = std::conditional_t< is_logical, vec, N>, Self>; @@ -293,6 +344,7 @@ template struct VecOperators { struct OpMixin>> : public IncDecImpl {}; +#if __SYCL_USE_LIBSYCL8_VEC_IMPL #define __SYCL_VEC_BINOP_MIXIN(OP, OPERATOR) \ template \ struct OpMixin>> { \ @@ -341,6 +393,52 @@ template struct VecOperators { friend auto operator OPERATOR(const Self &v) { return apply(v); } \ }; +#else + +#define __SYCL_VEC_BINOP_MIXIN(OP, OPERATOR) \ + template \ + struct OpMixin>> { \ + friend result_t operator OPERATOR(const Self & lhs, \ + const Self & rhs) { \ + return VecOperators::apply(lhs, rhs); \ + } \ + template \ + friend std::enable_if_t, result_t> \ + operator OPERATOR(const Self & lhs, const T & rhs) { \ + return VecOperators::apply(lhs, Self{static_cast(rhs)}); \ + } \ + template \ + friend std::enable_if_t, result_t> \ + operator OPERATOR(const T & lhs, const Self & rhs) { \ + return VecOperators::apply(Self{static_cast(lhs)}, rhs); \ + } \ + }; + +#define __SYCL_VEC_OPASSIGN_MIXIN(OP, OPERATOR) \ + template \ + struct OpMixin>>> { \ + friend Self &operator OPERATOR(Self & lhs, const Self & rhs) { \ + lhs = OP{}(lhs, rhs); \ + return lhs; \ + } \ + template \ + friend std::enable_if_t, Self &> \ + operator OPERATOR(Self & lhs, const T & rhs) { \ + lhs = OP{}(lhs, rhs); \ + return lhs; \ + } \ + }; + +#define __SYCL_VEC_UOP_MIXIN(OP, OPERATOR) \ + template \ + struct OpMixin>> { \ + friend result_t operator OPERATOR(const Self & v) { \ + return apply(v); \ + } \ + }; + +#endif + __SYCL_INSTANTIATE_OPERATORS(__SYCL_VEC_BINOP_MIXIN, __SYCL_VEC_OPASSIGN_MIXIN, __SYCL_VEC_UOP_MIXIN) @@ -348,6 +446,7 @@ template struct VecOperators { #undef __SYCL_VEC_OPASSIGN_MIXIN #undef __SYCL_VEC_BINOP_MIXIN +#if __SYCL_USE_LIBSYCL8_VEC_IMPL template struct OpMixin>>> { template ::element_type> @@ -356,6 +455,7 @@ template struct VecOperators { return apply>(v); } }; +#endif template struct __SYCL_EBO CombineImpl : public OpMixin... {}; @@ -377,6 +477,7 @@ template struct VecOperators { OpAssign, IncDec> {}; }; +#if __SYCL_USE_LIBSYCL8_VEC_IMPL template class vec_arith : public VecOperators>::Combined {}; @@ -427,6 +528,7 @@ class vec_arith } }; #endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) +#endif #undef __SYCL_INSTANTIATE_OPERATORS diff --git a/sycl/include/sycl/vector.hpp b/sycl/include/sycl/vector.hpp index 7dbc453367fcd..556e91e7912ef 100644 --- a/sycl/include/sycl/vector.hpp +++ b/sycl/include/sycl/vector.hpp @@ -318,14 +318,18 @@ template class vec_base { // Provides a cross-platform vector class template that works efficiently on // SYCL devices as well as in host C++ code. template -class __SYCL_EBO vec - : public detail::vec_arith, - public detail::ApplyIf< - NumElements == 1, - detail::ScalarConversionOperatorsMixIn>>, - public detail::NamedSwizzlesMixinBoth>, - // Keep it last to simplify ABI layout test: - public detail::vec_base { +class __SYCL_EBO vec : +#if __SYCL_USE_LIBSYCL8_VEC_IMPL + public detail::vec_arith, +#else + public detail::VecOperators>::Combined, +#endif + public detail::ApplyIf< + NumElements == 1, + detail::ScalarConversionOperatorsMixIn>>, + public detail::NamedSwizzlesMixinBoth>, + // Keep it last to simplify ABI layout test: + public detail::vec_base { static_assert(std::is_same_v>, "DataT must be cv-unqualified"); @@ -408,6 +412,7 @@ class __SYCL_EBO vec constexpr vec &operator=(const vec &) = default; constexpr vec &operator=(vec &&) = default; +#if __SYCL_USE_LIBSYCL8_VEC_IMPL // Template required to prevent ambiguous overload with the copy assignment // when NumElements == 1. The template prevents implicit conversion from // vec<_, 1> to DataT. @@ -427,6 +432,14 @@ class __SYCL_EBO vec *this = Rhs.template as(); return *this; } +#else + template + typename std::enable_if_t, vec &> + operator=(const T &Rhs) { + *this = vec{static_cast(Rhs)}; + return *this; + } +#endif __SYCL2020_DEPRECATED("get_count() is deprecated, please use size() instead") static constexpr size_t get_count() { return size(); } @@ -536,8 +549,10 @@ class __SYCL_EBO vec int... T5> friend class detail::SwizzleOp; template friend class __SYCL_EBO vec; +#if __SYCL_USE_LIBSYCL8_VEC_IMPL // To allow arithmetic operators access private members of vec. template friend class detail::vec_arith; +#endif }; ///////////////////////// class sycl::vec ///////////////////////// diff --git a/sycl/test-e2e/Basic/vector/byte.cpp b/sycl/test-e2e/Basic/vector/byte.cpp index 6eadf6a177b7b..c2edadf147ed3 100644 --- a/sycl/test-e2e/Basic/vector/byte.cpp +++ b/sycl/test-e2e/Basic/vector/byte.cpp @@ -180,6 +180,7 @@ int main() { assert(SwizByte2Neg[0] == ~SwizByte2B[0]); } +#if __SYCL_USE_LIBSYCL8_VEC_IMPL { // std::byte is not an arithmetic type and it only supports the following // overloads of >> and << operators. @@ -207,6 +208,7 @@ int main() { assert(SwizShiftRight[0] == SwizByte2Shift[0] >> 3 && SwizShiftLeft[1] == SwizByte2Shift[1] << 3); } +#endif } return 0; diff --git a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.hpp b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.hpp index 879a4f6da9463..1ae7733195728 100644 --- a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.hpp +++ b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.hpp @@ -38,10 +38,10 @@ bool CheckResult(sycl::vec V, T2 Ref) { constexpr T RefVal = 2; \ VecT InVec{static_cast(RefVal)}; \ { \ - VecT OutVecsDevice[2]; \ + ResT OutVecsDevice[2]; \ T OutRefsDevice[2]; \ { \ - sycl::buffer OutVecsBuff{OutVecsDevice, 2}; \ + sycl::buffer OutVecsBuff{OutVecsDevice, 2}; \ sycl::buffer OutRefsBuff{OutRefsDevice, 2}; \ Q.submit([&](sycl::handler &CGH) { \ sycl::accessor OutVecsAcc{OutVecsBuff, CGH, sycl::read_write}; \ diff --git a/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp b/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp index 32afba1e95f07..876f800743769 100644 --- a/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp +++ b/sycl/test-e2e/DeviceLib/built-ins/vector_integer.cpp @@ -203,9 +203,9 @@ int main() { // abs { - s::uint2 r{0}; + s::int2 r{0}; { - s::buffer BufR(&r, s::range<1>(1)); + s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { auto AccR = BufR.get_access(cgh); @@ -214,8 +214,8 @@ int main() { }); }); } - unsigned int r1 = r.x(); - unsigned int r2 = r.y(); + int r1 = r.x(); + int r2 = r.y(); assert(r1 == 5); assert(r2 == 2); } @@ -240,9 +240,9 @@ int main() { // abs_diff { - s::uint2 r{0}; + s::int2 r{0}; { - s::buffer BufR(&r, s::range<1>(1)); + s::buffer BufR(&r, s::range<1>(1)); s::queue myQueue; myQueue.submit([&](s::handler &cgh) { auto AccR = BufR.get_access(cgh); @@ -251,8 +251,8 @@ int main() { }); }); } - unsigned int r1 = r.x(); - unsigned int r2 = r.y(); + int r1 = r.x(); + int r2 = r.y(); assert(r1 == 4); assert(r2 == 1); } diff --git a/sycl/test/basic_tests/vectors/assign.cpp b/sycl/test/basic_tests/vectors/assign.cpp index 67ed92971d022..2045eca0303e3 100644 --- a/sycl/test/basic_tests/vectors/assign.cpp +++ b/sycl/test/basic_tests/vectors/assign.cpp @@ -27,8 +27,8 @@ using sw_double_2 = decltype(std::declval>().swizzle<1, 2>()); // EXCEPT_IN_PREVIEW condition<> static_assert( std::is_assignable_v, half>); -static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v, float>); -static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v, double>); +static_assert( std::is_assignable_v, float>); +static_assert( std::is_assignable_v, double>); static_assert( std::is_assignable_v, vec>); static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v, vec>); static_assert(EXCEPT_IN_PREVIEW std::is_assignable_v, vec>);