diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp index a2f934a147c33..c12b9e6781eb9 100644 --- a/sycl/include/sycl/types.hpp +++ b/sycl/include/sycl/types.hpp @@ -987,6 +987,41 @@ template class vec { #endif #ifdef __SYCL_USE_EXT_VECTOR_TYPE__ +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \ + friend vec operator BINOP(const vec &Lhs, const vec &Rhs) { \ + vec Ret; \ + if constexpr (IsUsingArrayOnDevice) { \ + for (size_t I = 0; I < NumElements; ++I) { \ + Ret.setValue(I, (Lhs.getValue(I) BINOP Rhs.getValue(I))); \ + } \ + } else { \ + Ret.m_Data = Lhs.m_Data BINOP Rhs.m_Data; \ + if constexpr (std::is_same_v && CONVERT) { \ + Ret.ConvertToDataT(); \ + } \ + } \ + return Ret; \ + } \ + friend vec operator BINOP(const vec &Lhs, const DataT &Rhs) { \ + return Lhs BINOP vec(Rhs); \ + } \ + friend vec operator BINOP(const DataT &Lhs, const vec &Rhs) { \ + return vec(Lhs) BINOP Rhs; \ + } \ + friend vec &operator OPASSIGN(vec & Lhs, const vec & Rhs) { \ + Lhs = Lhs BINOP Rhs; \ + return Lhs; \ + } \ + template \ + friend typename std::enable_if_t operator OPASSIGN( \ + vec & Lhs, const DataT & Rhs) { \ + Lhs = Lhs BINOP vec(Rhs); \ + return Lhs; \ + } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) + +#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \ template \ vec operator BINOP(const EnableIfNotUsingArrayOnDevice &Rhs) const { \ @@ -1024,38 +1059,37 @@ template class vec { *this = *this BINOP vec(Rhs); \ return *this; \ } +#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES) + #else // __SYCL_USE_EXT_VECTOR_TYPE__ #if defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \ - vec operator BINOP(const vec &Rhs) const { \ + friend vec operator BINOP(const vec &Lhs, const vec &Rhs) { \ vec Ret{}; \ if constexpr (NativeVec) \ - Ret.m_Data = m_Data BINOP Rhs.m_Data; \ + Ret.m_Data = Lhs.m_Data BINOP Rhs.m_Data; \ else \ for (size_t I = 0; I < NumElements; ++I) \ - Ret.setValue(I, (DataT)(vec_data::get(getValue( \ + Ret.setValue(I, (DataT)(vec_data::get(Lhs.getValue( \ I)) BINOP vec_data::get(Rhs.getValue(I)))); \ return Ret; \ } \ - template \ - typename std::enable_if_t< \ - std::is_convertible_v && \ - (std::is_fundamental_v> || \ - detail::is_half_or_bf16_v>), \ - vec> \ - operator BINOP(const T & Rhs) const { \ - return *this BINOP vec(static_cast(Rhs)); \ + friend vec operator BINOP(const vec &Lhs, const DataT &Rhs) { \ + return Lhs BINOP vec(Rhs); \ } \ - vec &operator OPASSIGN(const vec & Rhs) { \ - *this = *this BINOP Rhs; \ - return *this; \ + friend vec operator BINOP(const DataT &Lhs, const vec &Rhs) { \ + return vec(Lhs) BINOP Rhs; \ + } \ + friend vec &operator OPASSIGN(vec & Lhs, const vec & Rhs) { \ + Lhs = Lhs BINOP Rhs; \ + return Lhs; \ } \ template \ - typename std::enable_if_t operator OPASSIGN( \ - const DataT & Rhs) { \ - *this = *this BINOP vec(Rhs); \ - return *this; \ + friend typename std::enable_if_t operator OPASSIGN( \ + vec & Lhs, const DataT & Rhs) { \ + Lhs = Lhs BINOP vec(Rhs); \ + return Lhs; \ } #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) @@ -1120,6 +1154,42 @@ template class vec { // Use __SYCL_DEVICE_ONLY__ macro because cast to OpenCL vector type is defined // by SYCL device compiler only. #ifdef __SYCL_DEVICE_ONLY__ +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#define __SYCL_RELLOGOP(RELLOGOP) \ + friend vec operator RELLOGOP(const vec & Lhs, \ + const vec & Rhs) { \ + vec Ret{}; \ + /* This special case is needed since there are no standard operator|| */ \ + /* or operator&& functions for std::array. */ \ + if constexpr (IsUsingArrayOnDevice && \ + (std::string_view(#RELLOGOP) == "||" || \ + std::string_view(#RELLOGOP) == "&&")) { \ + for (size_t I = 0; I < NumElements; ++I) { \ + /* We cannot use SetValue here as the operator is not a friend of*/ \ + /* Ret on Windows. */ \ + Ret[I] = static_cast(-(vec_data::get( \ + Lhs.getValue(I)) RELLOGOP vec_data::get(Rhs.getValue(I)))); \ + } \ + } else { \ + Ret = vec( \ + (typename vec::vector_t)( \ + Lhs.m_Data RELLOGOP Rhs.m_Data)); \ + if (NumElements == 1) /*Scalar 0/1 logic was applied, invert*/ \ + Ret *= -1; \ + } \ + return Ret; \ + } \ + friend vec operator RELLOGOP(const vec & Lhs, \ + const DataT & Rhs) { \ + return Lhs RELLOGOP vec(Rhs); \ + } \ + friend vec operator RELLOGOP(const DataT & Lhs, \ + const vec & Rhs) { \ + return vec(Lhs) RELLOGOP Rhs; \ + } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) + +#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_RELLOGOP(RELLOGOP) \ vec operator RELLOGOP(const vec & Rhs) const { \ vec Ret{}; \ @@ -1129,9 +1199,10 @@ template class vec { (std::string_view(#RELLOGOP) == "||" || \ std::string_view(#RELLOGOP) == "&&")) { \ for (size_t I = 0; I < NumElements; ++I) { \ - Ret.setValue(I, \ - -(vec_data::get(getValue(I)) \ - RELLOGOP vec_data::get(Rhs.getValue(I)))); \ + /* We cannot use SetValue here as the operator is not a friend of*/ \ + /* Ret on Windows. */ \ + Ret[I] = static_cast(-(vec_data::get( \ + getValue(I)) RELLOGOP vec_data::get(Rhs.getValue(I)))); \ } \ } else { \ Ret = vec( \ @@ -1150,13 +1221,38 @@ template class vec { operator RELLOGOP(const T & Rhs) const { \ return *this RELLOGOP vec(static_cast(Rhs)); \ } +#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES) #else +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#define __SYCL_RELLOGOP(RELLOGOP) \ + friend vec operator RELLOGOP(const vec & Lhs, \ + const vec & Rhs) { \ + vec Ret{}; \ + for (size_t I = 0; I < NumElements; ++I) { \ + /* We cannot use SetValue here as the operator is not a friend of*/ \ + /* Ret on Windows. */ \ + Ret[I] = static_cast(-(vec_data::get( \ + Lhs.getValue(I)) RELLOGOP vec_data::get(Rhs.getValue(I)))); \ + } \ + return Ret; \ + } \ + friend vec operator RELLOGOP(const vec & Lhs, \ + const DataT & Rhs) { \ + return Lhs RELLOGOP vec(Rhs); \ + } \ + friend vec operator RELLOGOP(const DataT & Lhs, \ + const vec & Rhs) { \ + return vec(Lhs) RELLOGOP Rhs; \ + } +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_RELLOGOP(RELLOGOP) \ vec operator RELLOGOP(const vec & Rhs) const { \ vec Ret{}; \ for (size_t I = 0; I < NumElements; ++I) { \ - Ret.setValue(I, -(vec_data::get(getValue(I)) \ - RELLOGOP vec_data::get(Rhs.getValue(I)))); \ + /* We cannot use SetValue here as the operator is not a friend of*/ \ + /* Ret on Windows. */ \ + Ret[I] = static_cast(-(vec_data::get( \ + getValue(I)) RELLOGOP vec_data::get(Rhs.getValue(I)))); \ } \ return Ret; \ } \ @@ -1168,6 +1264,7 @@ template class vec { operator RELLOGOP(const T & Rhs) const { \ return *this RELLOGOP vec(static_cast(Rhs)); \ } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #endif __SYCL_RELLOGOP(==) @@ -1184,6 +1281,18 @@ template class vec { #ifdef __SYCL_UOP #error "Undefine __SYCL_UOP macro" #endif +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#define __SYCL_UOP(UOP, OPASSIGN) \ + friend vec &operator UOP(vec & Rhs) { \ + Rhs OPASSIGN vec_data::get(1); \ + return Rhs; \ + } \ + friend vec operator UOP(vec &Lhs, int) { \ + vec Ret(Lhs); \ + Lhs OPASSIGN vec_data::get(1); \ + return Ret; \ + } +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_UOP(UOP, OPASSIGN) \ vec &operator UOP() { \ *this OPASSIGN vec_data::get(1); \ @@ -1194,6 +1303,7 @@ template class vec { *this OPASSIGN vec_data::get(1); \ return Ret; \ } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) __SYCL_UOP(++, +=) __SYCL_UOP(--, -=) @@ -1203,150 +1313,84 @@ template class vec { // operator~() available only when: dataT != float && dataT != double // && dataT != half - template - typename std::enable_if_t> && - (!IsUsingArrayOnDevice && !IsUsingArrayOnHost), - vec> - operator~() const { - vec Ret{(typename vec::DataType) ~m_Data}; - if constexpr (std::is_same_v) { - Ret.ConvertToDataT(); - } - return Ret; - } - template - typename std::enable_if_t> && - (IsUsingArrayOnDevice || IsUsingArrayOnHost), - vec> - operator~() const { - vec Ret{}; - for (size_t I = 0; I < NumElements; ++I) { - Ret.setValue(I, ~getValue(I)); + friend vec operator~(const vec &Rhs) { + if constexpr (IsUsingArrayOnDevice || IsUsingArrayOnHost) { + vec Ret{}; + for (size_t I = 0; I < NumElements; ++I) { + Ret.setValue(I, ~Rhs.getValue(I)); + } + return Ret; + } else { + vec Ret{(typename vec::DataType) ~Rhs.m_Data}; + if constexpr (std::is_same_v) { + Ret.ConvertToDataT(); + } + return Ret; } - return Ret; } - template -#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) - using OpNotRet = detail::rel_t; -#else - using OpNotRet = T; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES - // operator! - template - EnableIfNotUsingArray, N>> operator!() const { - return vec{(typename vec::DataType) !m_Data} -#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) - .template as, N>>(); -#else - ; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES - } - - // std::byte neither supports ! unary op or casting, so special handling is - // needed. And, worse, Windows has a conflict with 'byte'. + friend vec, NumElements> operator!(const vec &Rhs) { + if constexpr (IsUsingArrayOnDevice || IsUsingArrayOnHost) { + vec Ret{}; + for (size_t I = 0; I < NumElements; ++I) { #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) - template - typename std::enable_if_t && - (IsUsingArrayOnDevice || IsUsingArrayOnHost), - vec, N>> - operator!() const { - vec Ret{}; - for (size_t I = 0; I < NumElements; ++I) { - Ret.setValue(I, std::byte{!vec_data::get(getValue(I))}); + // std::byte neither supports ! unary op or casting, so special handling + // is needed. And, worse, Windows has a conflict with 'byte'. + if constexpr (std::is_same_v) { + Ret.setValue(I, std::byte{!vec_data::get(Rhs.getValue(I))}); + } else +#endif // (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0) + { + Ret.setValue(I, !vec_data::get(Rhs.getValue(I))); + } + } + return Ret.template as, NumElements>>(); + } else { + return vec{(typename vec::DataType) !Rhs.m_Data} + .template as, NumElements>>(); } -#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) - return Ret.template as, N>>(); -#else - return Ret; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES } - template - typename std::enable_if_t && - (IsUsingArrayOnDevice || IsUsingArrayOnHost), - vec, N>> - operator!() const { - vec Ret{}; - for (size_t I = 0; I < NumElements; ++I) - Ret.setValue(I, !vec_data::get(getValue(I))); -#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) - return Ret.template as, N>>(); -#else - return Ret; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES - } -#else - template - EnableIfUsingArray, N>> operator!() const { - vec Ret{}; - for (size_t I = 0; I < NumElements; ++I) - Ret.setValue(I, !vec_data::get(getValue(I))); -#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) - return Ret.template as, N>>(); -#else - return Ret; -#endif // __INTEL_PREVIEW_BREAKING_CHANGES - } -#endif - // operator + - template EnableIfNotUsingArray operator+() const { - return vec{+m_Data}; - } - - template EnableIfUsingArray operator+() const { - vec Ret{}; - for (size_t I = 0; I < NumElements; ++I) - Ret.setValue(I, vec_data::get(+vec_data::get(getValue(I)))); - return Ret; - } - - // operator - - template EnableIfNotUsingArray operator-() const { - namespace oneapi = sycl::ext::oneapi; - if constexpr (IsBfloat16 && NumElements == 1) { - vec Ret{}; - oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data); - oneapi::bfloat16 w = -v; - Ret.m_Data = oneapi::detail::bfloat16ToBits(w); - } else if constexpr (IsBfloat16) { + friend vec operator+(const vec &Lhs) { + if constexpr (IsUsingArrayOnDevice || IsUsingArrayOnHost) { vec Ret{}; - for (size_t I = 0; I < NumElements; ++I) { - oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data[I]); - oneapi::bfloat16 w = -v; - Ret.m_Data[I] = oneapi::detail::bfloat16ToBits(w); - } + for (size_t I = 0; I < NumElements; ++I) + Ret.setValue( + I, vec_data::get(+vec_data::get(Lhs.getValue(I)))); return Ret; } else { - vec Ret{-m_Data}; - if constexpr (std::is_same_v) { - Ret.ConvertToDataT(); - } - return Ret; + return vec{+Lhs.m_Data}; } } - template EnableIfUsingArray operator-() const { + // operator - + friend vec operator-(const vec &Lhs) { namespace oneapi = sycl::ext::oneapi; vec Ret{}; if constexpr (IsBfloat16 && NumElements == 1) { - oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data); + oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(Lhs.m_Data); oneapi::bfloat16 w = -v; Ret.m_Data = oneapi::detail::bfloat16ToBits(w); } else if constexpr (IsBfloat16) { for (size_t I = 0; I < NumElements; I++) { - oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(m_Data[I]); + oneapi::bfloat16 v = oneapi::detail::bitsToBfloat16(Lhs.m_Data[I]); oneapi::bfloat16 w = -v; Ret.m_Data[I] = oneapi::detail::bfloat16ToBits(w); } - } else { + } else if constexpr (IsUsingArrayOnDevice || IsUsingArrayOnHost) { for (size_t I = 0; I < NumElements; ++I) - Ret.setValue(I, - vec_data::get(-vec_data::get(getValue(I)))); + Ret.setValue( + I, vec_data::get(-vec_data::get(Lhs.getValue(I)))); + return Ret; + } else { + Ret = vec{-Lhs.m_Data}; + if constexpr (std::is_same_v) { + Ret.ConvertToDataT(); + } + return Ret; } - return Ret; } #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) @@ -1770,6 +1814,19 @@ class SwizzleOp { #ifdef __SYCL_OPASSIGN #error "Undefine __SYCL_OPASSIGN macro." #endif +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#define __SYCL_OPASSIGN(OPASSIGN, OP) \ + friend SwizzleOp &operator OPASSIGN(SwizzleOp & Lhs, const DataT & Rhs) { \ + Lhs.operatorHelper(vec_t(Rhs)); \ + return Lhs; \ + } \ + template \ + friend SwizzleOp &operator OPASSIGN(SwizzleOp & Lhs, \ + const RhsOperation & Rhs) { \ + Lhs.operatorHelper(Rhs); \ + return Lhs; \ + } +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_OPASSIGN(OPASSIGN, OP) \ SwizzleOp &operator OPASSIGN(const DataT & Rhs) { \ operatorHelper(vec_t(Rhs)); \ @@ -1780,6 +1837,7 @@ class SwizzleOp { operatorHelper(Rhs); \ return *this; \ } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) __SYCL_OPASSIGN(+=, std::plus) __SYCL_OPASSIGN(-=, std::minus) @@ -1796,6 +1854,18 @@ class SwizzleOp { #ifdef __SYCL_UOP #error "Undefine __SYCL_UOP macro" #endif +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#define __SYCL_UOP(UOP, OPASSIGN) \ + friend SwizzleOp &operator UOP(SwizzleOp & Rhs) { \ + Rhs OPASSIGN static_cast(1); \ + return Rhs; \ + } \ + friend vec_t operator UOP(SwizzleOp &Lhs, int) { \ + vec_t Ret = Lhs; \ + Lhs OPASSIGN static_cast(1); \ + return Ret; \ + } +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_UOP(UOP, OPASSIGN) \ SwizzleOp &operator UOP() { \ *this OPASSIGN static_cast(1); \ @@ -1806,11 +1876,36 @@ class SwizzleOp { *this OPASSIGN static_cast(1); \ return Ret; \ } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) __SYCL_UOP(++, +=) __SYCL_UOP(--, -=) #undef __SYCL_UOP +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) + template + friend typename std::enable_if_t< + std::is_same_v && std::is_integral_v>, vec_t> + operator~(const SwizzleOp &Rhs) { + vec_t Tmp = Rhs; + return ~Tmp; + } + + friend vec_rel_t operator!(const SwizzleOp &Rhs) { + vec_t Tmp = Rhs; + return !Tmp; + } + + friend vec_t operator+(const SwizzleOp &Rhs) { + vec_t Tmp = Rhs; + return +Tmp; + } + + friend vec_t operator-(const SwizzleOp &Rhs) { + vec_t Tmp = Rhs; + return -Tmp; + } +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) template typename std::enable_if_t>, vec_t> operator~() { @@ -1832,6 +1927,80 @@ class SwizzleOp { vec_t Tmp = *this; return -Tmp; } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) + +#if defined(__INTEL_PREVIEW_BREAKING_CHANGES) +// scalar BINOP vec<> +// scalar BINOP SwizzleOp +// vec<> BINOP SwizzleOp +#ifdef __SYCL_BINOP +#error "Undefine __SYCL_BINOP macro" +#endif +#define __SYCL_BINOP(BINOP) \ + friend vec_t operator BINOP(const DataT &Lhs, const SwizzleOp &Rhs) { \ + vec_t Tmp = Rhs; \ + return Lhs BINOP Tmp; \ + } \ + friend vec_t operator BINOP(const SwizzleOp &Lhs, const DataT &Rhs) { \ + vec_t Tmp = Lhs; \ + return Tmp BINOP Rhs; \ + } \ + friend vec_t operator BINOP(const vec_t &Lhs, const SwizzleOp &Rhs) { \ + vec_t Tmp = Rhs; \ + return Lhs BINOP Tmp; \ + } \ + friend vec_t operator BINOP(const SwizzleOp &Lhs, const vec_t &Rhs) { \ + vec_t Tmp = Lhs; \ + return Tmp BINOP Rhs; \ + } + + __SYCL_BINOP(+) + __SYCL_BINOP(-) + __SYCL_BINOP(*) + __SYCL_BINOP(/) + __SYCL_BINOP(%) + __SYCL_BINOP(&) + __SYCL_BINOP(|) + __SYCL_BINOP(^) + __SYCL_BINOP(>>) + __SYCL_BINOP(<<) +#undef __SYCL_BINOP + +// scalar RELLOGOP vec<> +// scalar RELLOGOP SwizzleOp +// vec<> RELLOGOP SwizzleOp +#ifdef __SYCL_RELLOGOP +#error "Undefine __SYCL_RELLOGOP macro" +#endif +#define __SYCL_RELLOGOP(RELLOGOP) \ + friend vec_rel_t operator RELLOGOP(const DataT &Lhs, const SwizzleOp &Rhs) { \ + vec_t Tmp = Rhs; \ + return Lhs RELLOGOP Tmp; \ + } \ + friend vec_rel_t operator RELLOGOP(const SwizzleOp &Lhs, const DataT &Rhs) { \ + vec_t Tmp = Lhs; \ + return Tmp RELLOGOP Rhs; \ + } \ + friend vec_rel_t operator RELLOGOP(const vec_t &Lhs, const SwizzleOp &Rhs) { \ + vec_t Tmp = Rhs; \ + return Lhs RELLOGOP Tmp; \ + } \ + friend vec_rel_t operator RELLOGOP(const SwizzleOp &Lhs, const vec_t &Rhs) { \ + vec_t Tmp = Lhs; \ + return Tmp RELLOGOP Rhs; \ + } + + __SYCL_RELLOGOP(==) + __SYCL_RELLOGOP(!=) + __SYCL_RELLOGOP(>) + __SYCL_RELLOGOP(<) + __SYCL_RELLOGOP(>=) + __SYCL_RELLOGOP(<=) + // TODO: limit to integral types. + __SYCL_RELLOGOP(&&) + __SYCL_RELLOGOP(||) +#undef __SYCL_RELLOGOP +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) template > @@ -2265,6 +2434,7 @@ class SwizzleOp { }; } // namespace detail +#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES) // scalar BINOP vec<> // scalar BINOP SwizzleOp // vec<> BINOP SwizzleOp @@ -2374,6 +2544,7 @@ __SYCL_RELLOGOP(<=) __SYCL_RELLOGOP(&&) __SYCL_RELLOGOP(||) #undef __SYCL_RELLOGOP +#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES) namespace detail { diff --git a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp new file mode 100644 index 0000000000000..158b52ab5f27e --- /dev/null +++ b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp @@ -0,0 +1,150 @@ +// REQUIRES: preview-breaking-changes-supported +// RUN: %{build} -fpreview-breaking-changes -o %t.out +// RUN: %{run} %t.out + +// This test currently fails on AMD HIP due to an unresolved memcmp function. +// XFAIL: hip_amd + +// Checks scalar/vec operator ordering. + +#include + +template +using rel_t = std::conditional_t< + sizeof(T) == 1, int8_t, + std::conditional_t< + sizeof(T) == 2, int16_t, + std::conditional_t>>>; + +template +bool CheckResult(sycl::vec V, T2 Ref) { + if constexpr (IsRelOp) { + // Check that all elements have the same boolean representation as the + // scalar. + for (size_t I = 0; I < N; ++I) + if (static_cast(V[I]) != static_cast(Ref)) + return false; + return true; + } else { + // Check that all elements are equal to the scalar. + for (size_t I = 0; I < N; ++I) + if (V[I] != Ref) + return false; + return true; + } +} + +#define CHECK(Q, C, T, N, IS_RELOP, OP) \ + { \ + using VecT = sycl::vec; \ + using ResT = sycl::vec, T>, N>; \ + constexpr T RefVal = 2; \ + VecT InVec{static_cast(RefVal)}; \ + { \ + VecT OutVecsDevice[2]; \ + T OutRefsDevice[2]; \ + { \ + sycl::buffer OutVecsBuff{OutVecsDevice, 2}; \ + sycl::buffer OutRefsBuff{OutRefsDevice, 2}; \ + Q.submit([&](sycl::handler &CGH) { \ + sycl::accessor OutVecsAcc{OutVecsBuff, CGH, sycl::read_write}; \ + sycl::accessor OutRefsAcc{OutRefsBuff, CGH, sycl::read_write}; \ + CGH.single_task([=]() { \ + auto OutVec1 = InVec OP RefVal; \ + auto OutVec2 = RefVal OP InVec; \ + static_assert(std::is_same_v); \ + static_assert(std::is_same_v); \ + OutVecsAcc[0] = OutVec1; \ + OutVecsAcc[1] = OutVec2; \ + OutRefsAcc[0] = RefVal OP RefVal; \ + OutRefsAcc[1] = RefVal OP RefVal; \ + }); \ + }); \ + } \ + if (!CheckResult(OutVecsDevice[0], OutRefsDevice[0])) { \ + std::cout << ("Check of vector " #OP \ + " scalar from device failed for " #T " and " #N) \ + << std::endl; \ + ++C; \ + } \ + if (!CheckResult(OutVecsDevice[1], OutRefsDevice[1])) { \ + std::cout << ("Check of scalar " #OP \ + " vector from device failed for " #T " and " #N) \ + << std::endl; \ + ++C; \ + } \ + } \ + { \ + auto OutVec1 = InVec OP RefVal; \ + auto OutVec2 = RefVal OP InVec; \ + static_assert(std::is_same_v); \ + static_assert(std::is_same_v); \ + if (!CheckResult(OutVec1, RefVal OP RefVal)) { \ + std::cout << ("Check of vector " #OP \ + " scalar from host failed for " #T " and " #N) \ + << std::endl; \ + ++C; \ + } \ + if (!CheckResult(OutVec2, RefVal OP RefVal)) { \ + std::cout << ("Check of scalar " #OP \ + " vector from host failed for " #T " and " #N) \ + << std::endl; \ + ++C; \ + } \ + } \ + } + +#define CHECK_SIZES(Q, C, T, IS_RELOP, OP) \ + CHECK(Q, C, T, 1, IS_RELOP, OP) \ + CHECK(Q, C, T, 2, IS_RELOP, OP) \ + CHECK(Q, C, T, 4, IS_RELOP, OP) \ + CHECK(Q, C, T, 8, IS_RELOP, OP) \ + CHECK(Q, C, T, 16, IS_RELOP, OP) + +// NOTE: For the sake of compile-time we pick only a few operators per category. +#define CHECK_SIZES_AND_COMMON_OPS(Q, C, T) \ + CHECK_SIZES(Q, Failures, T, false, *) \ + CHECK_SIZES(Q, Failures, T, true, &&) \ + CHECK_SIZES(Q, Failures, T, true, ==) \ + CHECK_SIZES(Q, Failures, T, true, <) \ + CHECK_SIZES(Q, Failures, T, true, >=) +#define CHECK_SIZES_AND_INT_ONLY_OPS(Q, C, T) \ + CHECK_SIZES(Q, Failures, T, false, %) \ + CHECK_SIZES(Q, Failures, T, false, >>) \ + CHECK_SIZES(Q, Failures, T, false, ^) + +int main() { + sycl::queue Q; + int Failures = 0; + + // Check operators on types with requirements if they are supported. + if (Q.get_device().has(sycl::aspect::fp16)) { + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, sycl::half); + } + if (Q.get_device().has(sycl::aspect::fp64)) { + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, double); + } + + // Check all operators without requirements. + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, float); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, int8_t); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, int16_t); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, int32_t); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, int64_t); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, uint8_t); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, uint16_t); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, uint32_t); + CHECK_SIZES_AND_COMMON_OPS(Q, Failures, uint64_t); + + // Check integer only operators. + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, int8_t); + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, int16_t); + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, int32_t); + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, int64_t); + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, uint8_t); + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, uint16_t); + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, uint32_t); + CHECK_SIZES_AND_INT_ONLY_OPS(Q, Failures, uint64_t); + return Failures; +} diff --git a/sycl/test/basic_tests/types.cpp b/sycl/test/basic_tests/types.cpp index 6aab1e433c7a7..14a1070567274 100644 --- a/sycl/test/basic_tests/types.cpp +++ b/sycl/test/basic_tests/types.cpp @@ -134,12 +134,13 @@ template inline void checkVecNotReturnType() { using Vector = sycl::vec; #if defined(__INTEL_PREVIEW_BREAKING_CHANGES) using ExpectedVector = sycl::vec; + using OpNotResult = decltype(operator!(std::declval())); #else using ExpectedVector = sycl::vec; -#endif using OpNotResult = decltype(std::declval().operator!()); +#endif static_assert(std::is_same_v, - "Incorrect vec::operator! return type"); + "Incorrect operator! return type"); } // the math built-in testing ensures that the vec binary ops get tested,