91 namespace __ESIMD_DNS {
111 enum class CmpOp { lt, lte, gte, gt, eq, ne };
113 enum class UnaryOp { minus,
plus, bit_not, log_not };
117 static inline constexpr
bool is_wrapper_elem_type_v =
118 std::is_same_v<T, sycl::half>;
121 static inline constexpr
bool is_valid_simd_elem_type_v =
122 (is_vectorizable_v<T> || is_wrapper_elem_type_v<T>);
124 struct invalid_raw_element_type;
127 template <
class T,
class SFINAE>
struct element_type_traits {
130 using RawT = invalid_raw_element_type;
134 using EnclosingCppT = void;
137 static inline constexpr
bool use_native_cpp_ops =
true;
141 static inline constexpr
bool is_floating_point =
false;
146 struct element_type_traits<
T,
std::
enable_if_t<is_vectorizable_v<T>>> {
148 using EnclosingCppT =
T;
149 static inline constexpr
bool use_native_cpp_ops =
true;
150 static inline constexpr
bool is_floating_point = std::is_floating_point_v<T>;
161 template <
class WrapperTy,
class StdTy,
int N>
162 ESIMD_INLINE vector_type_t<__raw_t<WrapperTy>, N>
163 __esimd_convertvector_to(vector_type_t<StdTy, N> Val)
164 #ifdef __SYCL_DEVICE_ONLY__
170 __ESIMD_UNSUPPORTED_ON_HOST;
172 #endif // __SYCL_DEVICE_ONLY__
174 template <
class WrapperTy,
class StdTy,
int N>
175 ESIMD_INLINE vector_type_t<StdTy, N>
176 __esimd_convertvector_from(vector_type_t<__raw_t<WrapperTy>, N> Val)
177 #ifdef __SYCL_DEVICE_ONLY__
183 __ESIMD_UNSUPPORTED_ON_HOST;
185 #endif // __SYCL_DEVICE_ONLY__
188 template <
class WrapperTy>
189 WrapperTy __esimd_wrapper_type_bitcast_to(__raw_t<WrapperTy> Val);
190 template <
class WrapperTy>
191 __raw_t<WrapperTy> __esimd_wrapper_type_bitcast_from(WrapperTy Val);
193 template <
class WrapperTy,
class StdTy>
struct wrapper_type_converter {
194 using RawTy = __raw_t<WrapperTy>;
197 ESIMD_INLINE
static vector_type_t<RawTy, N>
198 to_vector(vector_type_t<StdTy, N> Val) {
199 if constexpr (element_type_traits<WrapperTy>::use_native_cpp_ops) {
200 return __builtin_convertvector(Val, vector_type_t<RawTy, N>);
202 return __esimd_convertvector_to<WrapperTy, StdTy, N>(Val);
207 ESIMD_INLINE
static vector_type_t<StdTy, N>
208 from_vector(vector_type_t<RawTy, N> Val) {
209 if constexpr (element_type_traits<WrapperTy>::use_native_cpp_ops) {
210 return __builtin_convertvector(Val, vector_type_t<StdTy, N>);
212 return __esimd_convertvector_from<WrapperTy, StdTy, N>(Val);
220 template <
class DstWrapperTy,
class SrcWrapperTy,
int N,
221 class DstRawVecTy = vector_type_t<__raw_t<DstWrapperTy>, N>,
222 class SrcRawVecTy = vector_type_t<__raw_t<SrcWrapperTy>, N>>
223 ESIMD_INLINE DstRawVecTy convert_vector(SrcRawVecTy Val) {
224 if constexpr (std::is_same_v<SrcWrapperTy, DstWrapperTy>) {
226 }
else if constexpr (!is_wrapper_elem_type_v<SrcWrapperTy> &&
227 !is_wrapper_elem_type_v<DstWrapperTy>) {
228 return __builtin_convertvector(Val, DstRawVecTy);
242 using DstStdT =
typename element_type_traits<DstWrapperTy>::EnclosingCppT;
243 using SrcStdT =
typename element_type_traits<SrcWrapperTy>::EnclosingCppT;
244 using SrcConv = wrapper_type_converter<SrcWrapperTy, SrcStdT>;
245 using DstConv = wrapper_type_converter<DstWrapperTy, DstStdT>;
246 using DstStdVecT = vector_type_t<DstStdT, N>;
247 using SrcStdVecT = vector_type_t<SrcStdT, N>;
248 SrcStdVecT TmpSrcVal;
250 if constexpr (std::is_same_v<SrcStdT, SrcWrapperTy>) {
251 TmpSrcVal = std::move(Val);
253 TmpSrcVal = SrcConv::template from_vector<N>(Val);
255 if constexpr (std::is_same_v<SrcStdT, DstWrapperTy>) {
258 DstStdVecT TmpDstVal;
260 if constexpr (std::is_same_v<SrcStdT, DstStdVecT>) {
261 TmpDstVal = std::move(TmpSrcVal);
263 TmpDstVal = __builtin_convertvector(TmpSrcVal, DstStdVecT);
265 if constexpr (std::is_same_v<DstStdT, DstWrapperTy>) {
268 return DstConv::template to_vector<N>(TmpDstVal);
274 template <
class Ty> ESIMD_INLINE __raw_t<Ty> bitcast_to_raw_type(Ty Val) {
275 if constexpr (!is_wrapper_elem_type_v<Ty>) {
278 return __esimd_wrapper_type_bitcast_from<Ty>(Val);
282 template <
class Ty> ESIMD_INLINE Ty bitcast_to_wrapper_type(__raw_t<Ty> Val) {
283 if constexpr (!is_wrapper_elem_type_v<Ty>) {
286 return __esimd_wrapper_type_bitcast_to<Ty>(Val);
295 template <
class DstWrapperTy,
class SrcWrapperTy,
296 class DstRawTy = __raw_t<DstWrapperTy>,
297 class SrcRawTy = __raw_t<SrcWrapperTy>>
298 ESIMD_INLINE DstWrapperTy convert_scalar(SrcWrapperTy Val) {
299 if constexpr (std::is_same_v<SrcWrapperTy, DstWrapperTy>) {
301 }
else if constexpr (!is_wrapper_elem_type_v<SrcWrapperTy> &&
302 !is_wrapper_elem_type_v<DstWrapperTy>) {
303 return static_cast<DstRawTy
>(Val);
305 vector_type_t<SrcRawTy, 1> V0 = bitcast_to_raw_type<SrcWrapperTy>(Val);
306 vector_type_t<DstRawTy, 1> V1 =
307 convert_vector<DstWrapperTy, SrcWrapperTy, 1>(V0);
308 return bitcast_to_wrapper_type<DstWrapperTy>(V1[0]);
312 template <BinOp Op,
class T>
T binary_op_default_impl(
T X,
T Y) {
316 else if constexpr (Op == BinOp::sub)
318 else if constexpr (Op == BinOp::mul)
322 else if constexpr (Op == BinOp::rem)
334 else if constexpr (Op == BinOp::log_or)
336 else if constexpr (Op == BinOp::log_and)
341 template <CmpOp Op,
class T>
auto comparison_op_default_impl(
T X,
T Y) {
342 decltype(X < Y) Res{};
343 if constexpr (Op == CmpOp::lt)
345 else if constexpr (Op == CmpOp::lte)
347 else if constexpr (Op == CmpOp::eq)
349 else if constexpr (Op == CmpOp::ne)
351 else if constexpr (Op == CmpOp::gte)
353 else if constexpr (Op == CmpOp::gt)
358 template <UnaryOp Op,
class T>
auto unary_op_default_impl(
T X) {
359 if constexpr (Op == UnaryOp::minus)
363 else if constexpr (Op == UnaryOp::bit_not)
365 else if constexpr (Op == UnaryOp::log_not)
369 template <
class ElemT,
int N>
struct __hlp {
370 using RawElemT = __raw_t<ElemT>;
371 using RawVecT = vector_type_t<RawElemT, N>;
372 using BinopT = decltype(std::declval<RawVecT>() + std::declval<RawVecT>());
373 using CmpT = decltype(std::declval<RawVecT>() < std::declval<RawVecT>());
376 template <
class Hlp>
using __re_t =
typename Hlp::RawElemT;
377 template <
class Hlp>
using __rv_t =
typename Hlp::RawVecT;
378 template <
class Hlp>
using __cmp_t =
typename Hlp::CmpT;
382 template <BinOp Op,
class T> ESIMD_INLINE
T __esimd_binary_op(
T X,
T Y);
384 template <BinOp Op,
class T,
385 class = std::enable_if_t<is_valid_simd_elem_type_v<T>>>
386 ESIMD_INLINE
T binary_op_default(
T X,
T Y) {
387 static_assert(element_type_traits<T>::use_native_cpp_ops);
388 using T1 = __raw_t<T>;
389 T1 X1 = bitcast_to_raw_type(X);
390 T1 Y1 = bitcast_to_raw_type(Y);
391 T1 Res = binary_op_default_impl<Op>(X1, Y1);
392 return bitcast_to_wrapper_type<T>(Res);
398 template <BinOp Op,
class T> ESIMD_INLINE
T __esimd_binary_op(
T X,
T Y) {
399 using T1 =
typename element_type_traits<T>::EnclosingCppT;
400 T1 X1 = convert_scalar<T1, T>(X);
401 T1 Y1 = convert_scalar<T1, T>(Y);
402 return convert_scalar<T>(binary_op_default<Op, T1>(X1, Y1));
405 template <BinOp Op,
class T,
406 class = std::enable_if_t<is_valid_simd_elem_type_v<T>>>
407 ESIMD_INLINE
T binary_op(
T X,
T Y) {
408 if constexpr (element_type_traits<T>::use_native_cpp_ops) {
409 return binary_op_default<Op>(X, Y);
411 return __esimd_binary_op<Op>(X, Y);
417 template <BinOp Op,
class ElemT,
int N,
class RawVecT = __rv_t<__hlp<ElemT, N>>>
418 ESIMD_INLINE RawVecT vector_binary_op_default(RawVecT X, RawVecT Y) {
419 static_assert(element_type_traits<ElemT>::use_native_cpp_ops);
420 return binary_op_default_impl<Op, RawVecT>(X, Y);
426 template <BinOp Op,
class ElemT,
int N,
class RawVecT = __rv_t<__hlp<ElemT, N>>>
427 ESIMD_INLINE RawVecT __esimd_vector_binary_op(RawVecT X, RawVecT Y) {
428 using T1 =
typename element_type_traits<ElemT>::EnclosingCppT;
429 using VecT1 = vector_type_t<T1, N>;
430 VecT1 X1 = convert_vector<T1, ElemT, N>(X);
431 VecT1 Y1 = convert_vector<T1, ElemT, N>(Y);
432 return convert_vector<ElemT, T1, N>(
433 vector_binary_op_default<Op, T1, N>(X1, Y1));
436 template <BinOp Op,
class ElemT,
int N,
class RawVecT = __rv_t<__hlp<ElemT, N>>>
437 ESIMD_INLINE RawVecT vector_binary_op(RawVecT X, RawVecT Y) {
438 if constexpr (element_type_traits<ElemT>::use_native_cpp_ops) {
439 return vector_binary_op_default<Op, ElemT, N>(X, Y);
441 return __esimd_vector_binary_op<Op, ElemT, N>(X, Y);
447 template <UnaryOp Op,
class T> ESIMD_INLINE
T __esimd_unary_op(
T X);
449 template <UnaryOp Op,
class T,
450 class = std::enable_if_t<is_valid_simd_elem_type_v<T>>>
451 ESIMD_INLINE
T unary_op_default(
T X) {
452 static_assert(element_type_traits<T>::use_native_cpp_ops);
453 using T1 = __raw_t<T>;
454 T1 X1 = bitcast_to_raw_type(X);
455 T1 Res = unary_op_default_impl<Op>(X1);
456 return bitcast_to_wrapper_type<T>(Res);
462 template <UnaryOp Op,
class T> ESIMD_INLINE
T __esimd_unary_op(
T X) {
463 using T1 =
typename element_type_traits<T>::EnclosingCppT;
464 T1 X1 = convert_scalar<T1, T>(X);
465 return convert_scalar<T>(unary_op_default<Op, T1>(X1));
468 template <UnaryOp Op,
class T,
469 class = std::enable_if_t<is_valid_simd_elem_type_v<T>>>
470 ESIMD_INLINE
T unary_op(
T X) {
471 if constexpr (element_type_traits<T>::use_native_cpp_ops) {
472 return unary_op_default<Op>(X);
474 return __esimd_unary_op<Op>(X);
480 template <UnaryOp Op,
class ElemT,
int N,
481 class RawVecT = __rv_t<__hlp<ElemT, N>>>
482 ESIMD_INLINE RawVecT vector_unary_op_default(RawVecT X) {
483 static_assert(element_type_traits<ElemT>::use_native_cpp_ops);
484 return unary_op_default_impl<Op, RawVecT>(X);
490 template <UnaryOp Op,
class ElemT,
int N,
491 class RawVecT = __rv_t<__hlp<ElemT, N>>>
492 ESIMD_INLINE RawVecT __esimd_vector_unary_op(RawVecT X) {
493 using T1 =
typename element_type_traits<ElemT>::EnclosingCppT;
494 using VecT1 = vector_type_t<T1, N>;
495 VecT1 X1 = convert_vector<T1, ElemT, N>(X);
496 return convert_vector<ElemT, T1, N>(vector_unary_op_default<Op, T1, N>(X1));
499 template <UnaryOp Op,
class ElemT,
int N,
500 class RawVecT = __rv_t<__hlp<ElemT, N>>>
501 ESIMD_INLINE RawVecT vector_unary_op(RawVecT X) {
502 if constexpr (element_type_traits<ElemT>::use_native_cpp_ops) {
503 return vector_unary_op_default<Op, ElemT, N>(X);
505 return __esimd_vector_unary_op<Op, ElemT, N>(X);
511 template <CmpOp Op,
class ElemT,
int N,
class H = __hlp<ElemT, N>,
512 class RetT = __cmp_t<H>,
class RawVecT = __rv_t<H>>
513 ESIMD_INLINE RetT vector_comparison_op_default(RawVecT X, RawVecT Y) {
514 static_assert(element_type_traits<ElemT>::use_native_cpp_ops);
515 return comparison_op_default_impl<Op, RawVecT>(X, Y);
521 template <CmpOp Op,
class ElemT,
int N,
class H = __hlp<ElemT, N>,
522 class RetT = __cmp_t<H>,
class RawVecT = __rv_t<H>>
523 ESIMD_INLINE RetT __esimd_vector_comparison_op(RawVecT X, RawVecT Y) {
524 using T1 =
typename element_type_traits<ElemT>::EnclosingCppT;
525 using VecT1 = vector_type_t<T1, N>;
526 VecT1 X1 = convert_vector<T1, ElemT, N>(X);
527 VecT1 Y1 = convert_vector<T1, ElemT, N>(Y);
528 return convert_vector<element_type_t<RetT>, T1, N>(
529 vector_comparison_op_default<Op, T1, N>(X1, Y1));
532 template <CmpOp Op,
class ElemT,
int N,
class H = __hlp<ElemT, N>,
533 class RetT = __cmp_t<H>,
class RawVecT = __rv_t<H>>
534 ESIMD_INLINE RetT vector_comparison_op(RawVecT X, RawVecT Y) {
535 if constexpr (element_type_traits<ElemT>::use_native_cpp_ops) {
536 return vector_comparison_op_default<Op, ElemT, N>(X, Y);
538 return __esimd_vector_comparison_op<Op, ElemT, N>(X, Y);
546 class WrapperElementTypeProxy {
548 template <
class T = sycl::half>
549 static inline __raw_t<T> bitcast_from_half(
T Val) {
550 #ifdef __SYCL_DEVICE_ONLY__
554 #endif // __SYCL_DEVICE_ONLY__
557 template <
class T = sycl::half>
558 static inline T bitcast_to_half(__raw_t<T> Bits) {
559 #ifndef __SYCL_DEVICE_ONLY__
565 #endif // __SYCL_DEVICE_ONLY__
571 template <
typename T>
572 static inline constexpr
bool is_generic_floating_point_v =
573 element_type_traits<T>::is_floating_point;
584 struct invalid_computation_type;
586 template <
class T1,
class T2,
class SFINAE =
void>
struct computation_type {
587 using type = invalid_computation_type;
590 template <
class T1,
class T2>
591 struct computation_type<T1, T2,
593 is_valid_simd_elem_type_v<T2>>> {
595 template <
class T>
using tr = element_type_traits<T>;
598 std::conditional_t<tr<T>::use_native_cpp_ops,
typename tr<T>::RawT,
599 typename tr<T>::EnclosingCppT>;
600 static inline constexpr
bool is_wr1 = is_wrapper_elem_type_v<T1>;
601 static inline constexpr
bool is_wr2 = is_wrapper_elem_type_v<T2>;
602 static inline constexpr
bool is_fp1 = is_generic_floating_point_v<T1>;
603 static inline constexpr
bool is_fp2 = is_generic_floating_point_v<T2>;
609 decltype(std::declval<T1>() + std::declval<T2>()),
611 std::is_same_v<T1, T2>,
618 std::conditional_t<is_fp1, T1, T2>,
622 decltype(std::declval<native_t<T1>>() +
623 std::declval<native_t<T2>>())>>>;
626 template <
class T1,
class T2>
627 struct computation_type<
629 std::
enable_if_t<is_simd_like_type_v<T1> || is_simd_like_type_v<T2>>> {
631 using Ty1 = element_type_t<T1>;
632 using Ty2 = element_type_t<T2>;
633 using EltTy =
typename computation_type<Ty1, Ty2>::type;
634 static constexpr
int N1 = is_simd_like_type_v<T1> ?
T1::length : 0;
635 static constexpr
int N2 = is_simd_like_type_v<T2> ?
T2::length : 0;
636 static_assert((N1 == N2) || ((N1 & N2) == 0),
"size mismatch");
637 static constexpr
int N = N1 ? N1 : N2;
643 template <
class T1,
class T2 = T1>
644 using computation_type_t =
645 typename computation_type<remove_cvref_t<T1>, remove_cvref_t<T2>>::type;
652 struct element_type_traits<
T,
std::
enable_if_t<std::is_same_v<T, sycl::half>>> {
657 #ifdef __SYCL_DEVICE_ONLY__
660 using EnclosingCppT = RawT;
663 static inline constexpr
bool use_native_cpp_ops =
true;
665 using RawT = uint16_t;
666 using EnclosingCppT = float;
669 static inline constexpr
bool use_native_cpp_ops =
false;
670 #endif // __SYCL_DEVICE_ONLY__
672 static inline constexpr
bool is_floating_point =
true;
675 using half_raw = __raw_t<sycl::half>;
679 __esimd_wrapper_type_bitcast_to<sycl::half>(half_raw Val) {
680 return WrapperElementTypeProxy::bitcast_to_half(Val);
684 ESIMD_INLINE half_raw
685 __esimd_wrapper_type_bitcast_from<sycl::half>(
sycl::half Val) {
686 return WrapperElementTypeProxy::bitcast_from_half(Val);
690 struct is_esimd_arithmetic_type<__raw_t<
sycl::
half>, void> : std::true_type {};
694 O << static_cast<float>(rhs);
699 float ValFloat = 0.0f;