19 #include <type_traits>
22 inline namespace _V1 {
31 #if defined(__SYCL_BINOP) || defined(BINOP_BASE)
32 #error "Undefine __SYCL_BINOP and BINOP_BASE macro"
35 #ifdef __SYCL_DEVICE_ONLY__
36 #define BINOP_BASE(BINOP, OPASSIGN, CONVERT, COND) \
37 template <typename T = DataT> \
38 friend std::enable_if_t<(COND), vec_t> operator BINOP(const vec_t & Lhs, \
39 const vec_t & Rhs) { \
41 if constexpr (vec_t::IsBfloat16) { \
42 for (size_t I = 0; I < NumElements; ++I) { \
43 Ret[I] = Lhs[I] BINOP Rhs[I]; \
46 auto ExtVecLhs = sycl::bit_cast<typename vec_t::vector_t>(Lhs); \
47 auto ExtVecRhs = sycl::bit_cast<typename vec_t::vector_t>(Rhs); \
48 Ret = vec<DataT, NumElements>(ExtVecLhs BINOP ExtVecRhs); \
49 if constexpr (std::is_same_v<DataT, bool> && CONVERT) { \
50 vec_arith_common<bool, NumElements>::ConvertToDataT(Ret); \
57 #define BINOP_BASE(BINOP, OPASSIGN, CONVERT, COND) \
58 template <typename T = DataT> \
59 friend std::enable_if_t<(COND), vec_t> operator BINOP(const vec_t & Lhs, \
60 const vec_t & Rhs) { \
62 for (size_t I = 0; I < NumElements; ++I) { \
63 Ret[I] = Lhs[I] BINOP Rhs[I]; \
69 #define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT, COND) \
70 BINOP_BASE(BINOP, OPASSIGN, CONVERT, COND) \
72 template <typename T = DataT> \
73 friend std::enable_if_t<(COND), vec_t> operator BINOP(const vec_t & Lhs, \
74 const DataT & Rhs) { \
75 return Lhs BINOP vec_t(Rhs); \
77 template <typename T = DataT> \
78 friend std::enable_if_t<(COND), vec_t> operator BINOP(const DataT & Lhs, \
79 const vec_t & Rhs) { \
80 return vec_t(Lhs) BINOP Rhs; \
82 template <typename T = DataT> \
83 friend std::enable_if_t<(COND), vec_t> &operator OPASSIGN( \
84 vec_t & Lhs, const vec_t & Rhs) { \
85 Lhs = Lhs BINOP Rhs; \
88 template <int Num = NumElements, typename T = DataT> \
89 friend std::enable_if_t<(Num != 1) && (COND), vec_t &> operator OPASSIGN( \
90 vec_t & Lhs, const DataT & Rhs) { \
91 Lhs = Lhs BINOP vec_t(Rhs); \
109 template <
typename DataT,
int NumElements>
class vec_arith_common;
112 template <
typename DataT,
int NumElements>
121 #ifdef __SYCL_DEVICE_ONLY__
123 auto extVec = sycl::bit_cast<typename vec_t::vector_t>(Rhs);
131 for (
size_t I = 0; I < NumElements; ++I) {
134 Ret[I] =
static_cast<ocl_t>(-1 * (!Rhs[I]));
142 #ifdef __SYCL_DEVICE_ONLY__
143 auto extVec = sycl::bit_cast<typename vec_t::vector_t>(Lhs);
144 return vec_t{+extVec};
147 for (
size_t I = 0; I < NumElements; ++I)
157 for (
size_t I = 0; I < NumElements; I++)
160 #ifndef __SYCL_DEVICE_ONLY__
161 for (
size_t I = 0; I < NumElements; ++I)
164 auto extVec = sycl::bit_cast<typename vec_t::vector_t>(Lhs);
165 Ret =
vec_t{-extVec};
166 if constexpr (std::is_same_v<DataT, bool>) {
178 #error "Undefine __SYCL_UOP macro"
180 #define __SYCL_UOP(UOP, OPASSIGN) \
181 friend vec_t &operator UOP(vec_t & Rhs) { \
182 Rhs OPASSIGN DataT{1}; \
185 friend vec_t operator UOP(vec_t &Lhs, int) { \
187 Lhs OPASSIGN DataT{1}; \
199 #if defined(__SYCL_RELLOGOP) || defined(RELLOGOP_BASE)
200 #error "Undefine __SYCL_RELLOGOP and RELLOGOP_BASE macro."
203 #ifdef __SYCL_DEVICE_ONLY__
204 #define RELLOGOP_BASE(RELLOGOP, COND) \
205 template <typename T = DataT> \
206 friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
207 const vec_t & Lhs, const vec_t & Rhs) { \
208 vec<ocl_t, NumElements> Ret{}; \
211 if constexpr (vec_t::IsBfloat16) { \
212 for (size_t I = 0; I < NumElements; ++I) { \
213 Ret[I] = static_cast<ocl_t>(-(Lhs[I] RELLOGOP Rhs[I])); \
216 auto ExtVecLhs = sycl::bit_cast<typename vec_t::vector_t>(Lhs); \
217 auto ExtVecRhs = sycl::bit_cast<typename vec_t::vector_t>(Rhs); \
220 Ret = vec<ocl_t, NumElements>( \
221 (typename vec<ocl_t, NumElements>::vector_t)( \
222 ExtVecLhs RELLOGOP ExtVecRhs)); \
224 if constexpr (NumElements == 1) { \
231 #define RELLOGOP_BASE(RELLOGOP, COND) \
232 template <typename T = DataT> \
233 friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
234 const vec_t & Lhs, const vec_t & Rhs) { \
235 vec<ocl_t, NumElements> Ret{}; \
236 for (size_t I = 0; I < NumElements; ++I) { \
237 Ret[I] = static_cast<ocl_t>(-(Lhs[I] RELLOGOP Rhs[I])); \
243 #define __SYCL_RELLOGOP(RELLOGOP, COND) \
244 RELLOGOP_BASE(RELLOGOP, COND) \
246 template <typename T = DataT> \
247 friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
248 const vec_t & Lhs, const DataT & Rhs) { \
249 return Lhs RELLOGOP vec_t(Rhs); \
251 template <typename T = DataT> \
252 friend std::enable_if_t<(COND), vec<ocl_t, NumElements>> operator RELLOGOP( \
253 const DataT & Lhs, const vec_t & Rhs) { \
254 return vec_t(Lhs) RELLOGOP Rhs; \
270 #undef __SYCL_RELLOGOP
293 #if (!defined(_HAS_STD_BYTE) || _HAS_STD_BYTE != 0)
294 template <
int NumElements>
300 using DataT =
typename std::conditional_t<NumElements == 0, int, std::byte>;
313 for (
size_t I = 0; I < NumElements; ++I) {
314 Ret[I] = Lhs[I] << shift;
328 for (
size_t I = 0; I < NumElements; ++I) {
329 Ret[I] = Lhs[I] >> shift;
352 std::is_same_v<DataT, sycl::ext::oneapi::bfloat16>;
356 template <
typename T = DataT>
357 friend std::enable_if_t<!detail::is_vgenfloat_v<T>,
vec_t>
359 #ifdef __SYCL_DEVICE_ONLY__
360 auto extVec = sycl::bit_cast<typename vec_t::vector_t>(Rhs);
362 if constexpr (std::is_same_v<DataT, bool>) {
368 for (
size_t I = 0; I < NumElements; ++I) {
375 #ifdef __SYCL_DEVICE_ONLY__
378 static void ConvertToDataT(vec_bool_t &Ret) {
379 for (
size_t I = 0; I < NumElements; ++I) {
380 Ret[I] = bit_cast<int8_t>(Ret[I]) != 0;
friend vec_t & operator>>=(vec_t &Lhs, int shift)
typename std::conditional_t< NumElements==0, int, std::byte > DataT
friend vec_t operator>>(const vec_t &Lhs, int shift)
friend vec_t operator<<(const vec_t &Lhs, int shift)
friend vec_t & operator<<=(vec_t &Lhs, int shift)
friend std::enable_if_t<!detail::is_vgenfloat_v< T >, vec_t > operator~(const vec_t &Rhs)
static constexpr bool IsBfloat16
friend vec< ocl_t, NumElements > operator!(const vec_t &Rhs)
detail::select_cl_scalar_integral_signed_t< DataT > ocl_t
friend vec_t operator+(const vec_t &Lhs)
friend vec_t operator-(const vec_t &Lhs)
constexpr bool is_vgenfloat_v
select_apply_cl_scalar_t< T, sycl::opencl::cl_char, sycl::opencl::cl_short, sycl::opencl::cl_int, sycl::opencl::cl_long > select_cl_scalar_integral_signed_t
#define __SYCL_UOP(UOP, OPASSIGN)
#define __SYCL_BINOP(BINOP, OPASSIGN, CONVERT, COND)
#define __SYCL_RELLOGOP(RELLOGOP, COND)