From 166b625a7dc19d14d498eccc891621d6a176d369 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Thu, 11 Jan 2024 10:32:49 -0800 Subject: [PATCH 1/6] [SYCL] Change vec operators to be friends This commit changes operators for sycl::vec to be defined like they are in the SYCL specification, i.e. friend functions instead of members. Signed-off-by: Larsen, Steffen --- sycl/include/sycl/types.hpp | 424 +++++++++++++----- .../Basic/vector/vec_binary_scalar_order.cpp | 148 ++++++ sycl/test/basic_tests/types.cpp | 3 +- 3 files changed, 451 insertions(+), 124 deletions(-) create mode 100644 sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp index b9ec19748e05f..74761c9ddb893 100644 --- a/sycl/include/sycl/types.hpp +++ b/sycl/include/sycl/types.hpp @@ -940,6 +940,39 @@ 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; \ + } +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \ template \ vec operator BINOP(const EnableIfNotUsingArrayOnDevice &Rhs) const { \ @@ -977,42 +1010,39 @@ 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> || \ - std::is_same_v, half>), \ - 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) - -#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \ vec operator BINOP(const vec &Rhs) const { \ vec Ret; \ @@ -1040,7 +1070,7 @@ template class vec { *this = *this BINOP vec(Rhs); \ return *this; \ } -#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #endif // __SYCL_USE_EXT_VECTOR_TYPE__ @@ -1073,6 +1103,39 @@ 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) { \ + Ret.setValue(I, \ + -(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; \ + } +#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_RELLOGOP(RELLOGOP) \ vec operator RELLOGOP(const vec & Rhs) const { \ vec Ret{}; \ @@ -1103,7 +1166,28 @@ 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) { \ + Ret.setValue(I, -(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{}; \ @@ -1121,6 +1205,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(==) @@ -1137,6 +1222,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); \ @@ -1147,6 +1244,7 @@ template class vec { *this OPASSIGN vec_data::get(1); \ return Ret; \ } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) __SYCL_UOP(++, +=) __SYCL_UOP(--, -=) @@ -1156,120 +1254,73 @@ 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; + friend vec operator+(const vec &Lhs) { + if constexpr (IsUsingArrayOnDevice || IsUsingArrayOnHost) { + vec Ret{}; + for (size_t I = 0; I < NumElements; ++I) + Ret.setValue( + I, vec_data::get(+vec_data::get(Lhs.getValue(I)))); + return Ret; + } else { + return vec{+Lhs.m_Data}; + } } // operator - - template EnableIfNotUsingArray operator-() const { - vec Ret{-m_Data}; - if constexpr (std::is_same_v) { - Ret.ConvertToDataT(); + friend vec operator-(const vec &Lhs) { + if constexpr (IsUsingArrayOnDevice || IsUsingArrayOnHost) { + vec Ret{}; + 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{-Lhs.m_Data}; + if constexpr (std::is_same_v) { + Ret.ConvertToDataT(); + } + return Ret; } - return Ret; - } - - 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; } #endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) @@ -1691,6 +1742,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)); \ @@ -1701,6 +1765,7 @@ class SwizzleOp { operatorHelper(Rhs); \ return *this; \ } +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) __SYCL_OPASSIGN(+=, std::plus) __SYCL_OPASSIGN(-=, std::minus) @@ -1717,6 +1782,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); \ @@ -1727,11 +1804,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~() { @@ -1753,6 +1855,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 > @@ -2186,6 +2362,7 @@ class SwizzleOp { }; } // namespace detail +#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES) // scalar BINOP vec<> // scalar BINOP SwizzleOp // vec<> BINOP SwizzleOp @@ -2295,6 +2472,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..74e6ad775beac --- /dev/null +++ b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp @@ -0,0 +1,148 @@ +// REQUIRES: preview-breaking-changes-supported +// RUN: %{build} -fpreview-breaking-changes -o %t.out +// RUN: %{run} %t.out + +// 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; \ + int Failures = 0; \ + 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 cc07202f393db..36adbb11d14ce 100644 --- a/sycl/test/basic_tests/types.cpp +++ b/sycl/test/basic_tests/types.cpp @@ -134,10 +134,11 @@ 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"); } From 9ba0387bcd6d06b11ca092df5e0034aff4d23aec Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Jan 2024 01:01:27 -0800 Subject: [PATCH 2/6] Avoid SetValue for different vector types. Signed-off-by: Larsen, Steffen --- sycl/include/sycl/types.hpp | 26 ++++++++++++++++---------- 1 file changed, 16 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp index 74761c9ddb893..e2d874f5889bb 100644 --- a/sycl/include/sycl/types.hpp +++ b/sycl/include/sycl/types.hpp @@ -1114,9 +1114,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(Lhs.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( \ + Lhs.getValue(I)) RELLOGOP vec_data::get(Rhs.getValue(I)))); \ } \ } else { \ Ret = vec( \ @@ -1145,9 +1146,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( \ @@ -1174,8 +1176,10 @@ template class vec { const vec & Rhs) { \ vec Ret{}; \ for (size_t I = 0; I < NumElements; ++I) { \ - Ret.setValue(I, -(vec_data::get(Lhs.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( \ + Lhs.getValue(I)) RELLOGOP vec_data::get(Rhs.getValue(I)))); \ } \ return Ret; \ } \ @@ -1192,8 +1196,10 @@ template class vec { 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; \ } \ From 5e0224d27f2f9297c6c8c227bad9b0741970b764 Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Tue, 16 Jan 2024 07:43:35 -0800 Subject: [PATCH 3/6] Disable for HIP Signed-off-by: Larsen, Steffen --- sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp index 74e6ad775beac..0f1d42b1e64d1 100644 --- a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp +++ b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp @@ -2,6 +2,9 @@ // 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 From 2977692c839c227686312906472e4d952efa81db Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Tue, 30 Jan 2024 11:55:52 +0100 Subject: [PATCH 4/6] Update sycl/test/basic_tests/types.cpp --- sycl/test/basic_tests/types.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/basic_tests/types.cpp b/sycl/test/basic_tests/types.cpp index cb21cb81e3dbb..14a1070567274 100644 --- a/sycl/test/basic_tests/types.cpp +++ b/sycl/test/basic_tests/types.cpp @@ -140,7 +140,7 @@ template inline void checkVecNotReturnType() { 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, From 57ea6bf1d02df9d4ba7fa0254b05c75253f0957c Mon Sep 17 00:00:00 2001 From: Steffen Larsen Date: Tue, 30 Jan 2024 11:55:59 +0100 Subject: [PATCH 5/6] Update sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp --- sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp index 0f1d42b1e64d1..158b52ab5f27e 100644 --- a/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp +++ b/sycl/test-e2e/Basic/vector/vec_binary_scalar_order.cpp @@ -40,7 +40,6 @@ bool CheckResult(sycl::vec V, T2 Ref) { using VecT = sycl::vec; \ using ResT = sycl::vec, T>, N>; \ constexpr T RefVal = 2; \ - int Failures = 0; \ VecT InVec{static_cast(RefVal)}; \ { \ VecT OutVecsDevice[2]; \ From 4b97841d3499da460336465a7c2ecb24dfc3d85c Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 7 Feb 2024 01:53:15 -0800 Subject: [PATCH 6/6] Break up elses Signed-off-by: Larsen, Steffen --- sycl/include/sycl/types.hpp | 18 ++++++++++++------ 1 file changed, 12 insertions(+), 6 deletions(-) diff --git a/sycl/include/sycl/types.hpp b/sycl/include/sycl/types.hpp index dbedf34d1ff17..c12b9e6781eb9 100644 --- a/sycl/include/sycl/types.hpp +++ b/sycl/include/sycl/types.hpp @@ -1019,7 +1019,9 @@ template class vec { Lhs = Lhs BINOP vec(Rhs); \ return Lhs; \ } -#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#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 { \ @@ -1057,7 +1059,7 @@ template class vec { *this = *this BINOP vec(Rhs); \ return *this; \ } -#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES) #else // __SYCL_USE_EXT_VECTOR_TYPE__ @@ -1089,7 +1091,9 @@ template class vec { Lhs = Lhs BINOP vec(Rhs); \ return Lhs; \ } -#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) + +#if !defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT) \ vec operator BINOP(const vec &Rhs) const { \ vec Ret; \ @@ -1117,7 +1121,7 @@ template class vec { *this = *this BINOP vec(Rhs); \ return *this; \ } -#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES) #endif // __SYCL_USE_EXT_VECTOR_TYPE__ @@ -1183,7 +1187,9 @@ template class vec { const vec & Rhs) { \ return vec(Lhs) RELLOGOP Rhs; \ } -#else // defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#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{}; \ @@ -1215,7 +1221,7 @@ template class vec { operator RELLOGOP(const T & Rhs) const { \ return *this RELLOGOP vec(static_cast(Rhs)); \ } -#endif // defined(__INTEL_PREVIEW_BREAKING_CHANGES) +#endif // !defined(__INTEL_PREVIEW_BREAKING_CHANGES) #else #if defined(__INTEL_PREVIEW_BREAKING_CHANGES) #define __SYCL_RELLOGOP(RELLOGOP) \