25 inline namespace _V1 {
26 namespace ext::intel::esimd {
68 template <
typename T0,
typename T1,
int SZ>
69 __ESIMD_API std::enable_if_t<!detail::is_generic_floating_point_v<T0> ||
70 std::is_same_v<T1, T0>,
73 if constexpr (detail::is_generic_floating_point_v<T0>)
74 return __esimd_sat<T0, T1, SZ>(src.
data());
75 else if constexpr (detail::is_generic_floating_point_v<T1>) {
76 if constexpr (std::is_unsigned_v<T0>)
77 return __esimd_fptoui_sat<T0, T1, SZ>(src.
data());
79 return __esimd_fptosi_sat<T0, T1, SZ>(src.
data());
80 }
else if constexpr (std::is_unsigned_v<T0>) {
81 if constexpr (std::is_unsigned_v<T1>)
82 return __esimd_uutrunc_sat<T0, T1, SZ>(src.
data());
84 return __esimd_ustrunc_sat<T0, T1, SZ>(src.
data());
86 if constexpr (std::is_signed_v<T1>)
87 return __esimd_sstrunc_sat<T0, T1, SZ>(src.
data());
89 return __esimd_sutrunc_sat<T0, T1, SZ>(src.
data());
97 template <
typename TRes,
typename TArg,
int SZ>
101 if constexpr (detail::is_generic_floating_point_v<TArg>)
105 return convert<TRes>(Result);
108 template <
typename TRes,
typename TArg>
110 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<TRes>::value &&
111 detail::is_esimd_scalar<TArg>::value,
113 __esimd_abs_common_internal(TArg
src0) {
115 simd<TArg, 1> Result = __esimd_abs_common_internal<TArg>(Src0);
116 return convert<TRes>(Result)[0];
127 template <
typename TRes,
typename TArg,
int SZ>
128 __ESIMD_API std::enable_if_t<
129 !std::is_same<std::remove_const_t<TRes>, std::remove_const_t<TArg>>::value,
132 return detail::__esimd_abs_common_internal<TRes, TArg, SZ>(
src0.data());
140 template <
typename TRes,
typename TArg>
141 __ESIMD_API std::enable_if_t<!std::is_same<std::remove_const_t<TRes>,
142 std::remove_const_t<TArg>>::value &&
143 detail::is_esimd_scalar<TRes>::value &&
144 detail::is_esimd_scalar<TArg>::value,
145 std::remove_const_t<TRes>>
147 return detail::__esimd_abs_common_internal<TRes, TArg>(
src0);
158 return detail::__esimd_abs_common_internal<T1, T1, SZ>(
src0.data());
167 template <
typename T1>
168 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T1>::value,
169 std::remove_const_t<T1>>
171 return detail::__esimd_abs_common_internal<T1, T1>(
src0);
183 template <
typename T,
int SZ,
class Sat = saturation_off_tag>
185 constexpr
bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
187 if constexpr (std::is_floating_point<T>::value) {
188 auto Result = __spirv_ocl_fmax<T, SZ>(
src0.data(),
src1.data());
189 if constexpr (is_sat)
190 Result = __esimd_sat<T, T, SZ>(Result);
192 }
else if constexpr (std::is_unsigned<T>::value) {
193 auto Result = __esimd_umax<T, SZ>(
src0.data(),
src1.data());
194 if constexpr (is_sat)
195 Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
198 auto Result = __esimd_smax<T, SZ>(
src0.data(),
src1.data());
199 if constexpr (is_sat)
200 Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
215 template <
typename T,
int SZ,
class Sat = saturation_off_tag>
216 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value,
simd<T, SZ>>(
233 template <
typename T,
int SZ,
class Sat = saturation_off_tag>
234 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value,
simd<T, SZ>>(
249 template <
typename T,
class Sat = saturation_off_tag>
250 ESIMD_NODEBUG ESIMD_INLINE
251 std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(
max)(T
src0, T
src1,
268 template <
typename T,
int SZ,
class Sat = saturation_off_tag>
270 constexpr
bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
272 if constexpr (std::is_floating_point<T>::value) {
273 auto Result = __spirv_ocl_fmin<T, SZ>(
src0.data(),
src1.data());
274 if constexpr (is_sat)
275 Result = __esimd_sat<T, T, SZ>(Result);
277 }
else if constexpr (std::is_unsigned<T>::value) {
278 auto Result = __esimd_umin<T, SZ>(
src0.data(),
src1.data());
279 if constexpr (is_sat)
280 Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
283 auto Result = __esimd_smin<T, SZ>(
src0.data(),
src1.data());
284 if constexpr (is_sat)
285 Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
300 template <
typename T,
int SZ,
class Sat = saturation_off_tag>
301 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value,
simd<T, SZ>>(
318 template <
typename T,
int SZ,
class Sat = saturation_off_tag>
319 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value,
simd<T, SZ>>(
334 template <
typename T,
class Sat = saturation_off_tag>
335 ESIMD_NODEBUG ESIMD_INLINE
336 std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(
min)(T
src0, T
src1,
349 #if defined(__SYCL_DEVICE_ONLY__)
350 #define __ESIMD_VECTOR_IMPL(T, name, iname) \
351 __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N> res = \
352 __spirv_ocl_native_##iname<__ESIMD_DNS::__raw_t<T>, N>(src.data()); \
353 if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
356 return esimd::saturate<T>(simd<T, N>(res));
357 #define __ESIMD_SCALAR_IMPL(T, name, iname) \
358 __ESIMD_DNS::__raw_t<T> res = \
359 __spirv_ocl_native_##iname<__ESIMD_DNS::__raw_t<T>>(src); \
360 if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
363 return esimd::saturate<T>(simd<T, 1>(res))[0];
365 #define __ESIMD_VECTOR_IMPL(T, name, iname) return 0;
366 #define __ESIMD_SCALAR_IMPL(T, name, iname) return 0;
369 #define __ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname) \
371 template <class T, int N, class Sat = saturation_off_tag, \
372 class = std::enable_if_t<COND>> \
373 __ESIMD_API simd<T, N> name(simd<T, N> src, Sat sat = {}) { \
374 __ESIMD_VECTOR_IMPL(T, name, iname) \
378 template <typename T, class Sat = saturation_off_tag, \
379 class = std::enable_if_t<COND>> \
380 __ESIMD_API T name(T src, Sat sat = {}) { \
381 __ESIMD_SCALAR_IMPL(T, name, iname) \
384 #define __ESIMD_EMATH_IEEE_COND \
385 detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4)
387 #define __ESIMD_EMATH_SPIRV_COND \
388 std::is_same_v<T, float> || std::is_same_v<T, sycl::half>
410 class = std::enable_if_t<__ESIMD_EMATH_IEEE_COND>>
412 __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N> res =
413 __esimd_ieee_sqrt<T, N>(src.
data());
414 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
421 template <
typename T,
class Sat = saturation_off_tag,
422 class = std::enable_if_t<__ESIMD_EMATH_IEEE_COND>>
442 template <
class T,
int N,
class Sat = saturation_off_tag>
446 __ESIMD_API std::enable_if_t<std::is_same_v<T, double>,
simd<double, N>>
448 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
451 return esimd::saturate<double>(
inv(
sqrt(src)));
455 template <
class T,
class Sat = saturation_off_tag>
456 __ESIMD_API std::enable_if_t<std::is_same_v<T, double>,
double>
458 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
461 return esimd::saturate<double>(
inv(
sqrt(src)));
464 #undef __ESIMD_UNARY_INTRINSIC_DEF
465 #undef __ESIMD_VECTOR_IMPL
466 #undef __ESIMD_SCALAR_IMPL
468 #define __ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname) \
470 template <class T, int N, class U, class Sat = saturation_off_tag, \
471 class = std::enable_if_t<COND>> \
472 __ESIMD_API simd<T, N> name(simd<T, N> src0, simd<U, N> src1, \
474 using RawVecT = __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N>; \
475 RawVecT src1_raw_conv = detail::convert_vector<T, U, N>(src1.data()); \
476 RawVecT res_raw = __esimd_##iname<T, N>(src0.data(), src1_raw_conv); \
477 if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
480 return esimd::saturate<T>(simd<T, N>(res_raw)); \
484 template <class T, int N, class U, class Sat = saturation_off_tag, \
485 class = std::enable_if_t<COND>> \
486 __ESIMD_API simd<T, N> name(simd<T, N> src0, U src1, Sat sat = {}) { \
487 return name<T, N, U>(src0, simd<U, N>(src1), sat); \
491 template <class T, class U, class Sat = saturation_off_tag, \
492 class = std::enable_if_t<COND>> \
493 __ESIMD_API T name(T src0, U src1, Sat sat = {}) { \
494 simd<T, 1> res = name<T, 1, U>(simd<T, 1>(src0), simd<U, 1>(src1), sat); \
500 template <
class T,
int N,
class U,
class Sat = saturation_off_tag,
501 class = std::enable_if_t<__ESIMD_EMATH_SPIRV_COND>>
503 #if defined(__SYCL_DEVICE_ONLY__)
504 using RawVecT = __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N>;
505 RawVecT src1_raw_conv = detail::convert_vector<T, U, N>(
src1.data());
506 RawVecT res_raw = __spirv_ocl_native_powr<__ESIMD_DNS::__raw_t<T>, N>(
507 src0.data(), src1_raw_conv);
508 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
511 return esimd::saturate<T>(
simd<T, N>(res_raw));
518 template <
class T,
int N,
class U,
class Sat = saturation_off_tag,
519 class = std::enable_if_t<__ESIMD_EMATH_SPIRV_COND>>
525 template <
class T,
class U,
class Sat = saturation_off_tag,
526 class = std::enable_if_t<__ESIMD_EMATH_SPIRV_COND>>
528 #if defined(__SYCL_DEVICE_ONLY__)
529 using ResT = __ESIMD_DNS::__raw_t<T>;
530 ResT src1_raw_conv = detail::convert_scalar<T, U>(
src1);
532 __spirv_ocl_native_powr<__ESIMD_DNS::__raw_t<T>>(
src0, src1_raw_conv);
533 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
536 return esimd::saturate<T>(
simd<T, 1>(res_raw))[0];
545 #undef __ESIMD_BINARY_INTRINSIC_DEF
546 #undef __ESIMD_EMATH_IEEE_COND
547 #undef __ESIMD_EMATH_SPIRV_COND
557 constexpr
float ln2 = 0.69314718f;
559 constexpr
float log2e = 1.442695f;
567 template <
class T,
int SZ,
class Sat = saturation_off_tag>
569 using CppT = __ESIMD_DNS::__cpp_t<T>;
571 esimd::log2<T, SZ, saturation_off_tag>(
src0) * detail::ln2;
573 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
576 return esimd::saturate<T>(Result);
579 template <
class T,
class Sat = saturation_off_tag>
581 return esimd::log<T, 1>(
src0,
sat)[0];
588 template <
class T,
int SZ,
class Sat = saturation_off_tag>
590 using CppT = __ESIMD_DNS::__cpp_t<T>;
591 return esimd::exp2<T, SZ>(
src0 * detail::log2e,
sat);
594 template <
class T,
class Sat = saturation_off_tag>
596 return esimd::exp<T, 1>(
src0,
sat)[0];
608 #define __ESIMD_INTRINSIC_DEF(name) \
616 template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag> \
617 __ESIMD_API __ESIMD_NS::simd<T, SZ> name(__ESIMD_NS::simd<float, SZ> src0, \
619 __ESIMD_NS::simd<float, SZ> Result = __esimd_##name<SZ>(src0.data()); \
620 if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>) \
622 else if constexpr (!std::is_same_v<float, T>) { \
623 auto RawRes = __ESIMD_NS::saturate<float>(Result).data(); \
624 return __ESIMD_DNS::convert_vector<T, float, SZ>(std::move(RawRes)); \
626 return __ESIMD_NS::saturate<T>(Result); \
630 template <typename T, class Sat = __ESIMD_NS::saturation_off_tag> \
631 __ESIMD_API T name(float src0, Sat sat = {}) { \
632 __ESIMD_NS::simd<float, 1> Src0 = src0; \
633 __ESIMD_NS::simd<T, 1> Result = name<T>(Src0, sat); \
671 #undef __ESIMD_INTRINSIC_DEF
678 template <
typename RT,
int SZ,
class Sat = __ESIMD_NS::saturation_off_tag>
679 ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
681 return esimd::rndd<RT, SZ>(
src0,
sat);
685 template <
typename RT,
class Sat = __ESIMD_NS::saturation_off_tag>
687 return esimd::rndd<RT, 1U>(
src0,
sat)[0];
691 template <
typename RT,
int SZ,
class Sat = __ESIMD_NS::saturation_off_tag>
692 ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
694 return esimd::rndu<RT, SZ>(
src0,
sat);
698 template <
typename RT,
class Sat = __ESIMD_NS::saturation_off_tag>
700 return esimd::rndu<RT, 1U>(
src0,
sat);
711 template <
typename RT,
int SZ,
class Sat = __ESIMD_NS::saturation_off_tag>
712 __ESIMD_API __ESIMD_NS::simd<RT, SZ>
714 return esimd::rndz<RT, SZ>(
src0,
sat);
724 template <
typename RT,
class Sat = __ESIMD_NS::saturation_off_tag>
726 return esimd::rndz<RT, 1U>(
src0,
sat)[0];
744 ESIMD_INLINE std::enable_if_t<(N == 8 || N == 16 || N == 32),
uint>
746 return __esimd_pack_mask<N>(
src0.data());
758 ESIMD_INLINE std::enable_if_t<(N == 8 || N == 16 || N == 32),
simd_mask<N>>
760 return __esimd_unpack_mask<N>(
src0);
766 __ESIMD_API std::enable_if_t<(N != 8 && N != 16 && N < 32), uint>
768 simd_mask<(N < 8 ? 8 : N < 16 ? 16 : 32)> src_0 = 0;
769 src_0.template select<N, 1>() =
src0.template bit_cast_view<ushort>();
779 template <
typename T,
int N>
781 std::enable_if_t<(std::is_same_v<T, ushort> || std::is_same_v<T, uint>) &&
786 if constexpr (N == 8 || N == 16 || N == 32) {
787 return __esimd_pack_mask<N>(cmp.
data());
789 constexpr
int N1 = (N <= 8 ? 8 : N <= 16 ? 16 : 32);
791 res.template select<N, 1>() = cmp.
data();
792 return __esimd_pack_mask<N1>(res.
data());
800 template <
typename T,
int N>
801 ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
804 return __esimd_cbit<T, N>(src.
data());
809 template <
typename T>
811 std::enable_if_t<std::is_integral<T>::value && (
sizeof(T) <= 4), uint32_t>
822 template <
typename BaseTy,
typename RegionTy>
823 __ESIMD_API std::enable_if_t<
843 template <
typename T,
int N>
845 std::enable_if_t<std::is_integral<T>::value && (
sizeof(T) == 4),
simd<T, N>>
847 return __esimd_fbl<T, N>(src.
data());
852 template <
typename T>
853 __ESIMD_API std::enable_if_t<std::is_integral<T>::value && (
sizeof(T) == 4), T>
865 template <
typename BaseTy,
typename RegionTy>
866 __ESIMD_API std::enable_if_t<
886 template <
typename T,
int N>
887 __ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
888 std::is_signed<T>::value && (
sizeof(T) == 4),
891 return __esimd_sfbh<T, N>(src.
data());
900 template <
typename T,
int N>
901 __ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
902 !std::is_signed<T>::value && (
sizeof(T) == 4),
905 return __esimd_ufbh<T, N>(src.
data());
910 template <
typename T>
911 __ESIMD_API std::enable_if_t<std::is_integral<T>::value && (
sizeof(T) == 4), T>
923 template <
typename BaseTy,
typename RegionTy>
924 __ESIMD_API std::enable_if_t<
947 template <
typename T0,
typename T1,
int SZ,
class Sat = saturation_off_tag>
949 std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
952 using ComputationTy =
953 __ESIMD_DNS::computation_type_t<decltype(
src0), int32_t>;
954 ComputationTy Src0 =
src0;
955 ComputationTy Src1 =
src1;
957 if constexpr (std::is_same_v<Sat, saturation_on_tag>) {
958 if constexpr (std::is_unsigned<T0>::value) {
959 if constexpr (std::is_unsigned<
960 typename ComputationTy::element_type>::value)
961 return __esimd_uushl_sat<T0, typename ComputationTy::element_type, SZ>(
962 Src0.data(), Src1.data());
964 return __esimd_usshl_sat<T0, typename ComputationTy::element_type, SZ>(
965 Src0.data(), Src1.data());
967 if constexpr (std::is_signed<typename ComputationTy::element_type>::value)
968 return __esimd_sushl_sat<T0, typename ComputationTy::element_type, SZ>(
969 Src0.data(), Src1.data());
971 return __esimd_ssshl_sat<T0, typename ComputationTy::element_type, SZ>(
972 Src0.data(), Src1.data());
975 if constexpr (std::is_unsigned<T0>::value) {
976 if constexpr (std::is_unsigned<
977 typename ComputationTy::element_type>::value)
978 return __esimd_uushl<T0, typename ComputationTy::element_type, SZ>(
979 Src0.data(), Src1.data());
981 return __esimd_usshl<T0, typename ComputationTy::element_type, SZ>(
982 Src0.data(), Src1.data());
984 if constexpr (std::is_signed<typename ComputationTy::element_type>::value)
985 return __esimd_sushl<T0, typename ComputationTy::element_type, SZ>(
986 Src0.data(), Src1.data());
988 return __esimd_ssshl<T0, typename ComputationTy::element_type, SZ>(
989 Src0.data(), Src1.data());
1003 template <
typename T0,
typename T1,
int SZ,
typename U,
1004 class Sat = saturation_off_tag>
1005 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
1006 std::is_integral<T1>::value &&
1007 std::is_integral<U>::value,
1011 return shl<T0, T1, SZ>(
src0, Src1,
sat);
1023 template <
typename T0,
typename T1,
typename T2,
class Sat = saturation_off_tag>
1024 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1025 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
1026 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
1027 std::is_integral<T0>::value &&
1028 std::is_integral<T1>::value &&
1029 std::is_integral<T2>::value,
1030 std::remove_const_t<T0>>
1047 template <
typename T0,
typename T1,
int SZ,
class Sat = saturation_off_tag>
1049 std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
1052 using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
1053 typedef typename std::make_unsigned<IntermedTy>::type ComputationTy;
1058 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
1061 return saturate<T0>(Result);
1074 template <
typename T0,
typename T1,
int SZ,
typename U,
1075 class Sat = saturation_off_tag>
1076 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
1077 std::is_integral<T1>::value &&
1078 std::is_integral<U>::value,
1082 return lsr<T0, T1, SZ>(
src0, Src1,
sat);
1095 template <
typename T0,
typename T1,
typename T2,
class Sat = saturation_off_tag>
1096 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1097 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
1098 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
1099 std::is_integral<T0>::value &&
1100 std::is_integral<T1>::value &&
1101 std::is_integral<T2>::value,
1102 std::remove_const_t<T0>>
1120 template <
typename T0,
typename T1,
int SZ,
class Sat = saturation_off_tag>
1122 std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
1125 using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
1126 typedef typename std::make_signed<IntermedTy>::type ComputationTy;
1130 if constexpr (std::is_same_v<Sat, saturation_off_tag>)
1133 return saturate<T0>(Result);
1146 template <
typename T0,
typename T1,
int SZ,
typename U,
1147 class Sat = saturation_off_tag>
1148 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
1149 std::is_integral<T1>::value &&
1150 std::is_integral<U>::value,
1154 return asr<T0, T1, SZ>(
src0, Src1,
sat);
1167 template <
typename T0,
typename T1,
typename T2,
class Sat = saturation_off_tag>
1168 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1169 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
1170 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
1171 std::is_integral<T0>::value &&
1172 std::is_integral<T1>::value &&
1173 std::is_integral<T2>::value,
1174 std::remove_const_t<T0>>
1191 template <
typename T0,
typename T1,
int SZ,
class Sat = saturation_off_tag>
1193 std::enable_if_t<std::is_integral<T0>::value && std::is_integral<T1>::value,
1196 if constexpr (std::is_unsigned<T1>::value) {
1213 template <
typename T0,
typename T1,
int SZ,
typename U,
1214 class Sat = saturation_off_tag>
1215 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
1216 std::is_integral<T1>::value &&
1217 std::is_integral<U>::value,
1221 return shr<T0, T1, SZ>(
src0, Src1,
sat);
1233 template <
typename T0,
typename T1,
typename T2,
class Sat = saturation_off_tag>
1234 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1235 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
1236 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
1237 std::is_integral<T0>::value &&
1238 std::is_integral<T1>::value &&
1239 std::is_integral<T2>::value,
1240 std::remove_const_t<T0>>
1255 template <
typename T0,
typename T1,
int SZ>
1256 __ESIMD_API std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t,
1257 uint32_t,
int64_t, uint64_t>() &&
1258 detail::is_type<T1, int16_t, uint16_t, int32_t,
1259 uint32_t,
int64_t, uint64_t>(),
1262 return __esimd_rol<T0, T1, SZ>(
src0.data(),
src1.data());
1273 template <
typename T0,
typename T1,
int SZ,
typename U>
1275 std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
1277 detail::is_type<T1, int16_t, uint16_t, int32_t,
1278 uint32_t,
int64_t, uint64_t>() &&
1279 detail::is_type<U, int16_t, uint16_t, int32_t,
1280 uint32_t,
int64_t, uint64_t>(),
1284 return rol<T0>(
src0, Src1);
1294 template <
typename T0,
typename T1,
typename T2>
1296 std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1297 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
1298 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
1299 detail::is_type<T0, int16_t, uint16_t, int32_t,
1300 uint32_t,
int64_t, uint64_t>() &&
1301 detail::is_type<T1, int16_t, uint16_t, int32_t,
1302 uint32_t,
int64_t, uint64_t>() &&
1303 detail::is_type<T2, int16_t, uint16_t, int32_t,
1304 uint32_t,
int64_t, uint64_t>(),
1305 std::remove_const_t<T0>>
1320 template <
typename T0,
typename T1,
int SZ>
1321 __ESIMD_API std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t,
1322 uint32_t,
int64_t, uint64_t>() &&
1323 detail::is_type<T1, int16_t, uint16_t, int32_t,
1324 uint32_t,
int64_t, uint64_t>(),
1327 return __esimd_ror<T0, T1, SZ>(
src0.data(),
src1.data());
1338 template <
typename T0,
typename T1,
int SZ,
typename U>
1340 std::enable_if_t<detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
1342 detail::is_type<T1, int16_t, uint16_t, int32_t,
1343 uint32_t,
int64_t, uint64_t>() &&
1344 detail::is_type<U, int16_t, uint16_t, int32_t,
1345 uint32_t,
int64_t, uint64_t>(),
1349 return esimd::ror<T0>(
src0, Src1);
1359 template <
typename T0,
typename T1,
typename T2>
1361 std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1362 __ESIMD_DNS::is_esimd_scalar<T1>::value &&
1363 __ESIMD_DNS::is_esimd_scalar<T2>::value &&
1364 detail::is_type<T0, int16_t, uint16_t, int32_t,
1365 uint32_t,
int64_t, uint64_t>() &&
1366 detail::is_type<T1, int16_t, uint16_t, int32_t,
1367 uint32_t,
int64_t, uint64_t>() &&
1368 detail::is_type<T2, int16_t, uint16_t, int32_t,
1369 uint32_t,
int64_t, uint64_t>(),
1370 std::remove_const_t<T0>>
1394 template <
typename T1,
typename T2,
typename T3,
typename T4,
int N,
1395 class Sat = saturation_off_tag>
1396 __ESIMD_API std::enable_if_t<
1397 detail::is_dword_type<T1>::value && detail::is_dword_type<T2>::value &&
1398 detail::is_dword_type<T3>::value && detail::is_dword_type<T4>::value,
1401 #if defined(__SYCL_DEVICE_ONLY__)
1406 if constexpr (std::is_same_v<Sat, saturation_off_tag>) {
1407 if constexpr (std::is_unsigned<T1>::value) {
1408 if constexpr (std::is_unsigned<T2>::value) {
1409 Result = __esimd_uudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1412 Result = __esimd_usdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1416 if constexpr (std::is_unsigned<T2>::value) {
1417 Result = __esimd_sudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1420 Result = __esimd_ssdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1425 if constexpr (std::is_unsigned<T1>::value) {
1426 if constexpr (std::is_unsigned<T2>::value) {
1427 Result = __esimd_uudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1430 Result = __esimd_usdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1434 if constexpr (std::is_unsigned<T2>::value) {
1435 Result = __esimd_sudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1438 Result = __esimd_ssdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
1445 __ESIMD_UNSUPPORTED_ON_HOST;
1452 template <
typename... T>
1459 template <
typename... T>
1466 template <
typename... T>
1468 if constexpr (std::is_floating_point<T1>::value) {
1469 return __spirv_ocl_fmax<T1, SZ>(v1.
data(), v2.
data());
1470 }
else if constexpr (std::is_unsigned<T1>::value) {
1471 return __esimd_umax<T1, SZ>(v1.
data(), v2.
data());
1473 return __esimd_smax<T1, SZ>(v1.
data(), v2.
data());
1479 template <
typename... T>
1481 if constexpr (std::is_floating_point<T1>::value) {
1482 return __spirv_ocl_fmin<T1, SZ>(v1.
data(), v2.
data());
1483 }
else if constexpr (std::is_unsigned<T1>::value) {
1484 return __esimd_umin<T1, SZ>(v1.
data(), v2.
data());
1486 return __esimd_smin<T1, SZ>(v1.
data(), v2.
data());
1491 template <
typename T0,
typename T1,
int SZ,
1492 template <
typename RT,
typename T,
int N>
class OpType>
1494 if constexpr (
SZ == 1) {
1498 "Invaid input for reduce_single - the vector size must "
1499 "be power of two.");
1500 constexpr
int N =
SZ / 2;
1501 simd<T0, N> tmp = OpType<T0, T1, N>()(v.template select<N, 1>(0),
1502 v.template select<N, 1>(N));
1503 return reduce_single<T0, T0, N, OpType>(tmp);
1507 template <
typename T0,
typename T1,
int N1,
int N2,
1508 template <
typename RT,
typename T,
int N>
class OpType>
1510 if constexpr (N1 == N2) {
1512 return reduce_single<T0, T0, N1, OpType>(tmp);
1513 }
else if constexpr (N1 < N2) {
1514 simd<T0, N1> tmp1 = OpType<T0, T1, N1>()(v1, v2.template select<N1, 1>(0));
1515 constexpr
int N = N2 - N1;
1517 NT tmp2 = convert<T0>(v2.template select<N, 1>(N1).
read());
1518 return reduce_pair<T0, T0, N1, N, OpType>(tmp1, tmp2);
1521 "Invaid input for reduce_pair - N1 must be power of two.");
1522 constexpr
int N = N1 / 2;
1523 simd<T0, N> tmp = OpType<T0, T1, N>()(v1.template select<N, 1>(0),
1524 v1.template select<N, 1>(N));
1526 NT tmp2 = convert<T0>(v2);
1527 return reduce_pair<T0, T0, N, N2, OpType>(tmp, tmp2);
1531 template <
typename T0,
typename T1,
int SZ,
1532 template <
typename RT,
typename T,
int N>
class OpType>
1536 return reduce_single<T0, T1, SZ, OpType>(v);
1538 constexpr
unsigned N1 = 1u << detail::log2<SZ>();
1539 constexpr
unsigned N2 =
SZ - N1;
1543 return reduce_pair<T0, T1, N1, N2, OpType>(v1, v2);
1547 template <
typename T0,
typename T1,
int SZ>
1549 using TT = detail::computation_type_t<simd<T1, SZ>>;
1550 using RT =
typename TT::element_type;
1551 T0 retv = reduce<RT, T1, SZ, esimd_apply_sum>(v);
1555 template <
typename T0,
typename T1,
int SZ>
1557 using TT = detail::computation_type_t<simd<T1, SZ>>;
1558 using RT =
typename TT::element_type;
1559 T0 retv = reduce<RT, T1, SZ, esimd_apply_prod>(v);
1572 template <
typename T0,
typename T1,
int SZ>
1574 T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_max>(v);
1585 template <
typename T0,
typename T1,
int SZ>
1587 T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_min>(v);
1605 template <
typename T0,
typename T1,
int SZ,
typename BinaryOperation>
1607 if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
1608 std::plus<>>::value) {
1609 T0 retv = detail::sum<T0>(v);
1611 }
else if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
1612 std::multiplies<>>::value) {
1613 T0 retv = detail::prod<T0>(v);
1625 enum class bfn_t : uint8_t {
x = 0xAA,
y = 0xCC,
z = 0xF0 };
1628 uint8_t val =
static_cast<uint8_t
>(
x);
1630 return static_cast<bfn_t>(res);
1634 uint8_t arg0 =
static_cast<uint8_t
>(
x);
1635 uint8_t arg1 =
static_cast<uint8_t
>(
y);
1636 uint8_t res = arg0 | arg1;
1637 return static_cast<bfn_t>(res);
1641 uint8_t arg0 =
static_cast<uint8_t
>(
x);
1642 uint8_t arg1 =
static_cast<uint8_t
>(
y);
1643 uint8_t res = arg0 & arg1;
1644 return static_cast<bfn_t>(res);
1648 uint8_t arg0 =
static_cast<uint8_t
>(
x);
1649 uint8_t arg1 =
static_cast<uint8_t
>(
y);
1650 uint8_t res = arg0 ^ arg1;
1651 return static_cast<bfn_t>(res);
1662 template <bfn_t FuncControl,
typename T,
int N>
1663 __ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
1664 bfn(__ESIMD_NS::simd<T, N>
src0, __ESIMD_NS::simd<T, N>
src1,
1665 __ESIMD_NS::simd<T, N> src2) {
1666 if constexpr ((
sizeof(T) == 8) || ((
sizeof(T) == 1) && (N % 4 == 0)) ||
1667 ((
sizeof(T) == 2) && (N % 2 == 0))) {
1671 auto Result = __ESIMD_NS::bfn<FuncControl>(
1672 src0.template bit_cast_view<int32_t>().read(),
1673 src1.template bit_cast_view<int32_t>().read(),
1674 src2.template bit_cast_view<int32_t>().read());
1675 return Result.template bit_cast_view<T>();
1676 }
else if constexpr (
sizeof(T) == 2 ||
sizeof(T) == 4) {
1677 constexpr uint8_t FC =
static_cast<uint8_t
>(FuncControl);
1678 return __esimd_bfn<FC, T, N>(
src0.data(),
src1.data(), src2.data());
1679 }
else if constexpr (N % 2 == 0) {
1681 auto Result = __ESIMD_NS::bfn<FuncControl>(
1682 src0.template bit_cast_view<int16_t>().read(),
1683 src1.template bit_cast_view<int16_t>().read(),
1684 src2.template bit_cast_view<int16_t>().read());
1685 return Result.template bit_cast_view<T>();
1688 __ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
1689 Src0.template select<N, 1>() =
src0;
1690 Src1.template select<N, 1>() =
src1;
1691 Src2.template select<N, 1>() = src2;
1692 auto Result = __ESIMD_NS::bfn<FuncControl>(Src0, Src1, Src2);
1693 return Result.template select<N, 1>();
1704 template <bfn_t FuncControl,
typename T>
1705 ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
1706 __ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
1708 __ESIMD_NS::simd<T, 1> Src0 =
src0;
1709 __ESIMD_NS::simd<T, 1> Src1 =
src1;
1710 __ESIMD_NS::simd<T, 1> Src2 = src2;
1711 __ESIMD_NS::simd<T, 1> Result =
1712 esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
1725 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1726 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N>
src0,
1727 __ESIMD_NS::simd<uint32_t, N>
src1) {
1728 std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1729 __ESIMD_DNS::vector_type_t<uint32_t, N>>
1730 Result = __esimd_addc<uint32_t, N>(
src0.data(),
src1.data());
1732 carry = Result.first;
1733 return Result.second;
1743 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1744 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N>
src0,
1746 __ESIMD_NS::simd<uint32_t, N> Src1V =
src1;
1757 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1758 addc(__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t
src0,
1759 __ESIMD_NS::simd<uint32_t, N>
src1) {
1760 __ESIMD_NS::simd<uint32_t, N> Src0V =
src0;
1770 __ESIMD_API uint32_t
addc(uint32_t &carry, uint32_t
src0, uint32_t
src1) {
1771 __ESIMD_NS::simd<uint32_t, 1> CarryV = carry;
1772 __ESIMD_NS::simd<uint32_t, 1> Src0V =
src0;
1773 __ESIMD_NS::simd<uint32_t, 1> Src1V =
src1;
1774 __ESIMD_NS::simd<uint32_t, 1> Res =
addc(CarryV, Src0V, Src1V);
1787 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1788 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N>
src0,
1789 __ESIMD_NS::simd<uint32_t, N>
src1) {
1790 std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1791 __ESIMD_DNS::vector_type_t<uint32_t, N>>
1792 Result = __esimd_subb<uint32_t, N>(
src0.data(),
src1.data());
1794 borrow = Result.first;
1795 return Result.second;
1806 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1807 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N>
src0,
1809 __ESIMD_NS::simd<uint32_t, N> Src1V =
src1;
1821 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1822 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t
src0,
1823 __ESIMD_NS::simd<uint32_t, N>
src1) {
1824 __ESIMD_NS::simd<uint32_t, N> Src0V =
src0;
1835 __ESIMD_API uint32_t
subb(uint32_t &borrow, uint32_t
src0, uint32_t
src1) {
1836 __ESIMD_NS::simd<uint32_t, 1> BorrowV = borrow;
1837 __ESIMD_NS::simd<uint32_t, 1> Src0V =
src0;
1838 __ESIMD_NS::simd<uint32_t, 1> Src1V =
src1;
1839 __ESIMD_NS::simd<uint32_t, 1> Res =
subb(BorrowV, Src0V, Src1V);
1840 borrow = BorrowV[0];
1846 __ESIMD_API uint64_t
rdtsc() {
1847 #ifdef __SYCL_DEVICE_ONLY__
1848 return __spirv_ReadClockKHR(0);
1850 __ESIMD_UNSUPPORTED_ON_HOST;
const auto & data() const noexcept
raw_vector_type data() const
This class represents a reference to a sub-region of a base simd object.
static constexpr int length
typename ShapeTy::element_type element_type
The element type of this class, which could be different from the element type of the base object typ...
The main simd vector class.
__ESIMD_API std::enable_if_t< detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), simd< T0, SZ > > ror(simd< T1, SZ > src0, simd< T1, SZ > src1)
Rotate right operation with two vector inputs.
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value, simd< T0, SZ > > asr(simd< T1, SZ > src0, simd< T1, SZ > src1, Sat sat={})
Arithmetical Shift Right (vector version)
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value, simd< T0, SZ > > lsr(simd< T1, SZ > src0, simd< T1, SZ > src1, Sat sat={})
Logical Shift Right (vector version)
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<(N==8||N==16||N==32), uint > pack_mask(simd_mask< N > src0)
Pack a simd_mask into a single unsigned 32-bit integer value.
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value, simd< T0, SZ > > shr(simd< T1, SZ > src0, simd< T1, SZ > src1, Sat sat={})
Shift right operation (vector version)
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&(sizeof(T)==4), simd< T, N > > fbl(simd< T, N > src)
Find the per element number of the first bit set in the source operand starting from the least signif...
__ESIMD_API std::enable_if_t< std::is_integral< T0 >::value &&std::is_integral< T1 >::value, simd< T0, SZ > > shl(simd< T1, SZ > src0, simd< T1, SZ > src1, Sat sat={})
Shift left operation (vector version)
__ESIMD_API std::enable_if_t<(std::is_same_v< T, ushort >||std::is_same_v< T, uint >) &&(N > 0 &&N<=32), uint > ballot(simd< T, N > mask)
Compare source vector elements against zero and return a bitfield combining the comparison result.
__ESIMD_API std::enable_if_t< detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), simd< T0, SZ > > rol(simd< T1, SZ > src0, simd< T1, SZ > src1)
Rotate left operation with two vector inputs.
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< std::is_integral< T >::value &&(sizeof(T)<=4), simd< uint32_t, N > > cbit(simd< T, N > src)
Count number of bits set in the source operand per element.
__ESIMD_API std::enable_if_t< std::is_integral< T >::value &&std::is_signed< T >::value &&(sizeof(T)==4), simd< T, N > > fbh(simd< T, N > src)
Find the per element number of the first bit set in the source operand starting from the most signifi...
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<(N==8||N==16||N==32), simd_mask< N > > unpack_mask(uint src0)
Unpack an unsigned 32-bit integer value into a simd_mask.
#define __ESIMD_INTRINSIC_DEF(name)
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ > rndu(sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
Round-up (also known as ceil).
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ > rndd(sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
Round-down (also known as floor).
__ESIMD_API sycl::ext::intel::esimd::simd< RT, SZ > trunc(const sycl::ext::intel::esimd::simd< float, SZ > &src0, Sat sat={})
Round to integral value using the round to zero rounding mode (vector version).
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ > floor(const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
"Floor" operation, vector version - alias of rndd.
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ > rnde(sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
Round-to-even (also known as round).
ESIMD_INLINE sycl::ext::intel::esimd::simd< RT, SZ > ceil(const sycl::ext::intel::esimd::simd< float, SZ > src0, Sat sat={})
"Ceiling" operation, vector version - alias of rndu.
__ESIMD_API sycl::ext::intel::esimd ::simd< T, SZ > rndz(sycl::ext::intel::esimd ::simd< float, SZ > src0, Sat sat={})
Round-to-zero (also known as trunc).
__ESIMD_API std::enable_if_t< std::is_integral_v< T >, sycl::ext::intel::esimd::simd< T, N > > bfn(sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T, N > src1, sycl::ext::intel::esimd::simd< T, N > src2)
Performs binary function computation with three vector operands.
static constexpr bfn_t operator&(bfn_t x, bfn_t y)
bfn_t
This enum is used to encode all possible logical operations performed on the 3 input operands.
static constexpr bfn_t operator|(bfn_t x, bfn_t y)
static constexpr bfn_t operator~(bfn_t x)
static constexpr bfn_t operator^(bfn_t x, bfn_t y)
__ESIMD_API simd< T, N > cos(simd< T, N > src, Sat sat={})
Cosine.
__ESIMD_API simd< T, N > sqrt(simd< T, N > src, Sat sat={})
Square root.
__ESIMD_API simd< T, N > rsqrt(simd< T, N > src, Sat sat={})
Square root reciprocal - calculates 1/sqrt(x).
#define __ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname)
__ESIMD_API simd< T, N > pow(simd< T, N > src0, simd< U, N > src1, Sat sat={})
Power - calculates src0 in power of src1.
#define __ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname)
__ESIMD_API simd< T, N > inv(simd< T, N > src, Sat sat={})
Inversion - calculates (1/x).
__ESIMD_API simd< T, N > sqrt_ieee(simd< T, N > src, Sat sat={})
IEEE754-compliant square root. Supports float and double.
#define __ESIMD_EMATH_SPIRV_COND
__ESIMD_API simd< T, N > div_ieee(simd< T, N > src0, simd< U, N > src1, Sat sat={})
IEEE754-compliant floating-point division. Supports float and double.
__ESIMD_API simd< T, N > sin(simd< T, N > src, Sat sat={})
Sine.
#define __ESIMD_EMATH_IEEE_COND
__ESIMD_API simd< T, N > log2(simd< T, N > src, Sat sat={})
Logarithm base 2.
__ESIMD_API simd< T, N > exp2(simd< T, N > src, Sat sat={})
Exponent base 2.
ESIMD_INLINE ESIMD_NODEBUG T0 hmax(simd< T1, SZ > v)
ESIMD_DETAIL.
__ESIMD_API SZ simd< T, SZ > src1
__ESIMD_API simd< T, SZ >(max)(simd< T
Selects component-wise the maximum of the two vectors.
ESIMD_DETAIL ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ > log(simd< T, SZ > src0, Sat sat={})
Computes the natural logarithm of the given argument.
__ESIMD_API SZ simd< T, SZ > Sat sat
ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd< T1, SZ > v, BinaryOperation op)
Performs reduction over elements of the input vector.
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >() min(simd< T, SZ > src0, T src1, Sat sat={})
__ESIMD_API std::enable_if_t<!detail::is_generic_floating_point_v< T0 >||std::is_same_v< T1, T0 >, simd< T0, SZ > > saturate(simd< T1, SZ > src)
Conversion of input vector elements of type T1 into vector of elements of type T0 with saturation.
__ESIMD_API SZ simd< T, SZ > Sat int SZ
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N > addc(sycl::ext::intel::esimd::simd< uint32_t, N > &carry, sycl::ext::intel::esimd::simd< uint32_t, N > src0, sycl::ext::intel::esimd::simd< uint32_t, N > src1)
Performs add with carry of 2 unsigned 32-bit vectors.
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >() max(simd< T, SZ > src0, T src1, Sat sat={})
__ESIMD_API uint64_t rdtsc()
rdtsc - get the value of timestamp counter.
ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd< T1, SZ > v)
Performs 'minimum' operation reduction over elements of the input vector, that is,...
ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ > exp(simd< T, SZ > src0, Sat sat={})
Computes e raised to the power of the given argument.
__ESIMD_API SZ simd< T, SZ > Sat int class Sat
__ESIMD_API sycl::ext::intel::esimd::simd< uint32_t, N > subb(sycl::ext::intel::esimd::simd< uint32_t, N > &borrow, sycl::ext::intel::esimd::simd< uint32_t, N > src0, sycl::ext::intel::esimd::simd< uint32_t, N > src1)
Performs substraction with borrow of 2 unsigned 32-bit vectors.
__ESIMD_API std::enable_if_t< detail::is_dword_type< T1 >::value &&detail::is_dword_type< T2 >::value &&detail::is_dword_type< T3 >::value &&detail::is_dword_type< T4 >::value, simd< T1, N > > dp4a(simd< T2, N > src0, simd< T3, N > src1, simd< T4, N > src2, Sat sat={})
DP4A.
ESIMD_DETAIL __ESIMD_API std::enable_if_t< !std::is_same< std::remove_const_t< TRes >, std::remove_const_t< TArg > >::value, simd< TRes, SZ > > abs(simd< TArg, SZ > src0)
Get absolute value (vector version)
conditional< sizeof(long)==8, long, long long >::type int64_t
T0 reduce(simd< T1, SZ > v)
ESIMD_INLINE ESIMD_NODEBUG T0 prod(simd< T1, SZ > v)
T0 reduce_single(simd< T1, SZ > v)
ESIMD_INLINE ESIMD_NODEBUG T0 sum(simd< T1, SZ > v)
T0 reduce_pair(simd< T1, N1 > v1, simd< T1, N2 > v2)
constexpr ESIMD_INLINE bool isPowerOf2(unsigned int n)
Check if a given 32 bit positive integer is a power of 2 at compile time.
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
This type tag represents "saturation off" behavior.