DPC++ Runtime
Runtime libraries for oneAPI DPC++
math.hpp
Go to the documentation of this file.
1 //==-------------- math.hpp - DPC++ Explicit SIMD API --------------------==//
2 //
3 // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4 // See https://llvm.org/LICENSE.txt for license information.
5 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6 //
7 //===----------------------------------------------------------------------===//
8 // Implement Explicit SIMD math APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
21 
22 #include <cstdint>
23 
24 namespace sycl {
25 inline namespace _V1 {
26 namespace ext::intel::esimd {
27 
40 
43 
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());
78  else
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());
83  else
84  return __esimd_ustrunc_sat<T0, T1, SZ>(src.data());
85  } else {
86  if constexpr (std::is_signed_v<T1>)
87  return __esimd_sstrunc_sat<T0, T1, SZ>(src.data());
88  else
89  return __esimd_sutrunc_sat<T0, T1, SZ>(src.data());
90  }
91 }
92 
94 // abs
95 namespace detail {
96 
97 template <typename TRes, typename TArg, int SZ>
98 ESIMD_NODEBUG ESIMD_INLINE simd<TRes, SZ>
99 __esimd_abs_common_internal(simd<TArg, SZ> src0) {
100  simd<TArg, SZ> Result = simd<TArg, SZ>(__esimd_abs<TArg, SZ>(src0.data()));
101  return convert<TRes>(Result);
102 }
103 
104 template <typename TRes, typename TArg>
105 ESIMD_NODEBUG ESIMD_INLINE
106  std::enable_if_t<detail::is_esimd_scalar<TRes>::value &&
107  detail::is_esimd_scalar<TArg>::value,
108  TRes>
109  __esimd_abs_common_internal(TArg src0) {
110  simd<TArg, 1> Src0 = src0;
111  simd<TArg, 1> Result = __esimd_abs_common_internal<TArg>(Src0);
112  return convert<TRes>(Result)[0];
113 }
114 } // namespace detail
116 
123 template <typename TRes, typename TArg, int SZ>
124 __ESIMD_API std::enable_if_t<
125  !std::is_same<std::remove_const_t<TRes>, std::remove_const_t<TArg>>::value,
128  return detail::__esimd_abs_common_internal<TRes, TArg, SZ>(src0.data());
129 }
130 
136 template <typename TRes, typename TArg>
137 __ESIMD_API std::enable_if_t<!std::is_same<std::remove_const_t<TRes>,
138  std::remove_const_t<TArg>>::value &&
139  detail::is_esimd_scalar<TRes>::value &&
140  detail::is_esimd_scalar<TArg>::value,
141  std::remove_const_t<TRes>>
142 abs(TArg src0) {
143  return detail::__esimd_abs_common_internal<TRes, TArg>(src0);
144 }
145 
153 template <typename T1, int SZ> __ESIMD_API simd<T1, SZ> abs(simd<T1, SZ> src0) {
154  return detail::__esimd_abs_common_internal<T1, T1, SZ>(src0.data());
155 }
156 
163 template <typename T1>
164 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T1>::value,
165  std::remove_const_t<T1>>
166 abs(T1 src0) {
167  return detail::__esimd_abs_common_internal<T1, T1>(src0);
168 }
169 
179 template <typename T, int SZ, class Sat = saturation_off_tag>
181  constexpr bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
182 
183  if constexpr (std::is_floating_point<T>::value) {
184  auto Result = __esimd_fmax<T, SZ>(src0.data(), src1.data());
185  if constexpr (is_sat)
186  Result = __esimd_sat<T, T, SZ>(Result);
187  return simd<T, SZ>(Result);
188  } else if constexpr (std::is_unsigned<T>::value) {
189  auto Result = __esimd_umax<T, SZ>(src0.data(), src1.data());
190  if constexpr (is_sat)
191  Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
192  return simd<T, SZ>(Result);
193  } else {
194  auto Result = __esimd_smax<T, SZ>(src0.data(), src1.data());
195  if constexpr (is_sat)
196  Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
197  return simd<T, SZ>(Result);
198  }
199 }
200 
211 template <typename T, int SZ, class Sat = saturation_off_tag>
212 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
213  max)(simd<T, SZ> src0, T src1, Sat sat = {}) {
214  simd<T, SZ> Src1 = src1;
215  simd<T, SZ> Result = (esimd::max)(src0, Src1, sat);
216  return Result;
217 }
218 
229 template <typename T, int SZ, class Sat = saturation_off_tag>
230 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
231  max)(T src0, simd<T, SZ> src1, Sat sat = {}) {
232  simd<T, SZ> Src0 = src0;
233  simd<T, SZ> Result = (esimd::max)(Src0, src1, sat);
234  return Result;
235 }
236 
245 template <typename T, class Sat = saturation_off_tag>
246 ESIMD_NODEBUG ESIMD_INLINE
247 std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(max)(T src0, T src1,
248  Sat sat = {}) {
249  simd<T, 1> Src0 = src0;
250  simd<T, 1> Src1 = src1;
251  simd<T, 1> Result = (esimd::max)(Src0, Src1, sat);
252  return Result[0];
253 }
254 
264 template <typename T, int SZ, class Sat = saturation_off_tag>
266  constexpr bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
267 
268  if constexpr (std::is_floating_point<T>::value) {
269  auto Result = __esimd_fmin<T, SZ>(src0.data(), src1.data());
270  if constexpr (is_sat)
271  Result = __esimd_sat<T, T, SZ>(Result);
272  return simd<T, SZ>(Result);
273  } else if constexpr (std::is_unsigned<T>::value) {
274  auto Result = __esimd_umin<T, SZ>(src0.data(), src1.data());
275  if constexpr (is_sat)
276  Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
277  return simd<T, SZ>(Result);
278  } else {
279  auto Result = __esimd_smin<T, SZ>(src0.data(), src1.data());
280  if constexpr (is_sat)
281  Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
282  return simd<T, SZ>(Result);
283  }
284 }
285 
296 template <typename T, int SZ, class Sat = saturation_off_tag>
297 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
298  min)(simd<T, SZ> src0, T src1, Sat sat = {}) {
299  simd<T, SZ> Src1 = src1;
300  simd<T, SZ> Result = (esimd::min)(src0, Src1, sat);
301  return Result;
302 }
303 
314 template <typename T, int SZ, class Sat = saturation_off_tag>
315 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
316  min)(T src0, simd<T, SZ> src1, Sat sat = {}) {
317  simd<T, SZ> Src0 = src0;
318  simd<T, SZ> Result = (esimd::min)(Src0, src1, sat);
319  return Result;
320 }
321 
330 template <typename T, class Sat = saturation_off_tag>
331 ESIMD_NODEBUG ESIMD_INLINE
332 std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(min)(T src0, T src1,
333  Sat sat = {}) {
334  simd<T, 1> Src0 = src0;
335  simd<T, 1> Src1 = src1;
336  simd<T, 1> Result = (esimd::min)(Src0, Src1, sat);
337  return Result[0];
338 }
339 
341 
344 
345 #define __ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname) \
346  \
347  template <class T, int N, class Sat = saturation_off_tag, \
348  class = std::enable_if_t<COND>> \
349  __ESIMD_API simd<T, N> name(simd<T, N> src, Sat sat = {}) { \
350  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N> res = \
351  __esimd_##iname<T, N>(src.data()); \
352  if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
353  return res; \
354  else \
355  return esimd::saturate<T>(simd<T, N>(res)); \
356  } \
357  \
358  \
359  template <typename T, class Sat = saturation_off_tag, \
360  class = std::enable_if_t<COND>> \
361  __ESIMD_API T name(T src, Sat sat = {}) { \
362  simd<T, 1> src_vec = src; \
363  simd<T, 1> res = name<T, 1>(src_vec, sat); \
364  return res[0]; \
365  }
366 
367 #define __ESIMD_EMATH_COND \
368  detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4)
369 
370 #define __ESIMD_EMATH_IEEE_COND \
371  detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4)
372 
376 
382 
386 
390 
393 
398 
402 
406 
407 #undef __ESIMD_UNARY_INTRINSIC_DEF
408 
409 #define __ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname) \
410  \
411  template <class T, int N, class U, class Sat = saturation_off_tag, \
412  class = std::enable_if_t<COND>> \
413  __ESIMD_API simd<T, N> name(simd<T, N> src0, simd<U, N> src1, \
414  Sat sat = {}) { \
415  using RawVecT = __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N>; \
416  RawVecT src1_raw_conv = detail::convert_vector<T, U, N>(src1.data()); \
417  RawVecT res_raw = __esimd_##iname<T, N>(src0.data(), src1_raw_conv); \
418  if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
419  return res_raw; \
420  else \
421  return esimd::saturate<T>(simd<T, N>(res_raw)); \
422  } \
423  \
424  \
425  template <class T, int N, class U, class Sat = saturation_off_tag, \
426  class = std::enable_if_t<COND>> \
427  __ESIMD_API simd<T, N> name(simd<T, N> src0, U src1, Sat sat = {}) { \
428  return name<T, N, U>(src0, simd<U, N>(src1), sat); \
429  } \
430  \
431  \
432  template <class T, class U, class Sat = saturation_off_tag, \
433  class = std::enable_if_t<COND>> \
434  __ESIMD_API T name(T src0, U src1, Sat sat = {}) { \
435  simd<T, 1> res = name<T, 1, U>(simd<T, 1>(src0), simd<U, 1>(src1), sat); \
436  return res[0]; \
437  }
438 
443 
446 
447 #undef __ESIMD_BINARY_INTRINSIC_DEF
448 #undef __ESIMD_EMATH_COND
449 #undef __ESIMD_EMATH_IEEE_COND
450 
452 
455 
457 namespace detail {
458 // std::numbers::ln2_v<float> in c++20
459 constexpr float ln2 = 0.69314718f;
460 // std::numbers::log2e_v<float> in c++20
461 constexpr float log2e = 1.442695f;
462 } // namespace detail
464 
469 template <class T, int SZ, class Sat = saturation_off_tag>
470 ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> log(simd<T, SZ> src0, Sat sat = {}) {
471  using CppT = __ESIMD_DNS::__cpp_t<T>;
472  simd<T, SZ> Result =
473  esimd::log2<T, SZ, saturation_off_tag>(src0) * detail::ln2;
474 
475  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
476  return Result;
477  else
478  return esimd::saturate<T>(Result);
479 }
480 
481 template <class T, class Sat = saturation_off_tag>
482 ESIMD_NODEBUG ESIMD_INLINE T log(T src0, Sat sat = {}) {
483  return esimd::log<T, 1>(src0, sat)[0];
484 }
485 
490 template <class T, int SZ, class Sat = saturation_off_tag>
491 ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> exp(simd<T, SZ> src0, Sat sat = {}) {
492  using CppT = __ESIMD_DNS::__cpp_t<T>;
493  return esimd::exp2<T, SZ>(src0 * detail::log2e, sat);
494 }
495 
496 template <class T, class Sat = saturation_off_tag>
497 ESIMD_NODEBUG ESIMD_INLINE T exp(T src0, Sat sat = {}) {
498  return esimd::exp<T, 1>(src0, sat)[0];
499 }
500 
502 
505 
507 // Rounding intrinsics.
509 
510 #define __ESIMD_INTRINSIC_DEF(name) \
511  \
512  \
513  \
515  \
516  \
517  \
518  template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag> \
519  __ESIMD_API __ESIMD_NS::simd<T, SZ> name(__ESIMD_NS::simd<float, SZ> src0, \
520  Sat sat = {}) { \
521  __ESIMD_NS::simd<float, SZ> Result = __esimd_##name<SZ>(src0.data()); \
522  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>) \
523  return Result; \
524  else if constexpr (!std::is_same_v<float, T>) { \
525  auto RawRes = __ESIMD_NS::saturate<float>(Result).data(); \
526  return __ESIMD_DNS::convert_vector<T, float, SZ>(std::move(RawRes)); \
527  } else { \
528  return __ESIMD_NS::saturate<T>(Result); \
529  } \
530  } \
531  \
532  template <typename T, class Sat = __ESIMD_NS::saturation_off_tag> \
533  __ESIMD_API T name(float src0, Sat sat = {}) { \
534  __ESIMD_NS::simd<float, 1> Src0 = src0; \
535  __ESIMD_NS::simd<T, 1> Result = name<T>(Src0, sat); \
536  return Result[0]; \
537  }
538 
547 
556 
564 
572 
573 #undef __ESIMD_INTRINSIC_DEF
575 
578 
580 template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
581 ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
582 floor(const __ESIMD_NS::simd<float, SZ> src0, Sat sat = {}) {
583  return esimd::rndd<RT, SZ>(src0, sat);
584 }
585 
587 template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
588 ESIMD_INLINE RT floor(float src0, Sat sat = {}) {
589  return esimd::rndd<RT, 1U>(src0, sat)[0];
590 }
591 
593 template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
594 ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
595 ceil(const __ESIMD_NS::simd<float, SZ> src0, Sat sat = {}) {
596  return esimd::rndu<RT, SZ>(src0, sat);
597 }
598 
600 template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
601 ESIMD_INLINE RT ceil(float src0, Sat sat = {}) {
602  return esimd::rndu<RT, 1U>(src0, sat);
603 }
604 
613 template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
614 __ESIMD_API __ESIMD_NS::simd<RT, SZ>
615 trunc(const __ESIMD_NS::simd<float, SZ> &src0, Sat sat = {}) {
616  return esimd::rndz<RT, SZ>(src0, sat);
617 }
618 
626 template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
627 __ESIMD_API RT trunc(float src0, Sat sat = {}) {
628  return esimd::rndz<RT, 1U>(src0, sat)[0];
629 }
630 
632 
635 
644 template <int N>
645 ESIMD_NODEBUG ESIMD_INLINE
646  std::enable_if_t<(N == 8 || N == 16 || N == 32), uint>
648  return __esimd_pack_mask<N>(src0.data());
649 }
650 
658 template <int N>
659 ESIMD_NODEBUG ESIMD_INLINE
660  std::enable_if_t<(N == 8 || N == 16 || N == 32), simd_mask<N>>
662  return __esimd_unpack_mask<N>(src0);
663 }
664 
667 template <int N>
668 __ESIMD_API std::enable_if_t<(N != 8 && N != 16 && N < 32), uint>
670  simd_mask<(N < 8 ? 8 : N < 16 ? 16 : 32)> src_0 = 0;
671  src_0.template select<N, 1>() = src0.template bit_cast_view<ushort>();
672  return esimd::pack_mask(src_0);
673 }
674 
681 template <typename T, int N>
682 __ESIMD_API std::enable_if_t<(std::is_same_v<T, ushort> ||
683  std::is_same_v<T, uint>)&&(N > 0 && N <= 32),
685 ballot(simd<T, N> mask) {
686  simd_mask<N> cmp = (mask != 0);
687  if constexpr (N == 8 || N == 16 || N == 32) {
688  return __esimd_pack_mask<N>(cmp.data());
689  } else {
690  constexpr int N1 = (N <= 8 ? 8 : N <= 16 ? 16 : 32);
691  simd<uint16_t, N1> res = 0;
692  res.template select<N, 1>() = cmp.data();
693  return __esimd_pack_mask<N1>(res.data());
694  }
695 }
696 
701 template <typename T, int N>
702 ESIMD_NODEBUG ESIMD_INLINE
703  std::enable_if_t<std::is_integral<T>::value && (sizeof(T) <= 4),
705  cbit(simd<T, N> src) {
706  return __esimd_cbit<T, N>(src.data());
707 }
708 
711 template <typename T>
712 __ESIMD_API
713  std::enable_if_t<std::is_integral<T>::value && (sizeof(T) <= 4), uint32_t>
714  cbit(T src) {
715  simd<T, 1> Src = src;
716  simd<uint32_t, 1> Result = esimd::cbit(Src);
717  return Result[0];
718 }
719 
724 template <typename BaseTy, typename RegionTy>
725 __ESIMD_API std::enable_if_t<
726  std::is_integral<
728  (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) <= 4) &&
730  uint32_t>
732  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
733  simd<Ty, 1> Src = src;
734  simd<uint32_t, 1> Result = esimd::cbit(Src);
735  return Result[0];
736 }
737 
745 template <typename T, int N>
746 __ESIMD_API
747  std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), simd<T, N>>
748  fbl(simd<T, N> src) {
749  return __esimd_fbl<T, N>(src.data());
750 }
751 
754 template <typename T>
755 __ESIMD_API std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), T>
756 fbl(T src) {
757  simd<T, 1> Src = src;
758  simd<T, 1> Result = esimd::fbl(Src);
759  return Result[0];
760 }
761 
767 template <typename BaseTy, typename RegionTy>
768 __ESIMD_API std::enable_if_t<
769  std::is_integral<
771  (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) == 4) &&
775  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
776  simd<Ty, 1> Src = src;
777  simd<Ty, 1> Result = esimd::fbl(Src);
778  return Result[0];
779 }
780 
788 template <typename T, int N>
789 __ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
790  std::is_signed<T>::value && (sizeof(T) == 4),
792 fbh(simd<T, N> src) {
793  return __esimd_sfbh<T, N>(src.data());
794 }
795 
802 template <typename T, int N>
803 __ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
804  !std::is_signed<T>::value && (sizeof(T) == 4),
806 fbh(simd<T, N> src) {
807  return __esimd_ufbh<T, N>(src.data());
808 }
809 
812 template <typename T>
813 __ESIMD_API std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), T>
814 fbh(T src) {
815  simd<T, 1> Src = src;
816  simd<T, 1> Result = esimd::fbh(Src);
817  return Result[0];
818 }
819 
825 template <typename BaseTy, typename RegionTy>
826 __ESIMD_API std::enable_if_t<
827  std::is_integral<
829  (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) == 4) &&
833  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
834  simd<Ty, 1> Src = src;
835  simd<Ty, 1> Result = esimd::fbh(Src);
836  return Result[0];
837 }
838 
840 
843 
856 template <typename T1, typename T2, typename T3, typename T4, int N,
857  class Sat = saturation_off_tag>
858 __ESIMD_API std::enable_if_t<
859  detail::is_dword_type<T1>::value && detail::is_dword_type<T2>::value &&
860  detail::is_dword_type<T3>::value && detail::is_dword_type<T4>::value,
863 #if defined(__SYCL_DEVICE_ONLY__)
864  simd<T1, N> Result;
865  simd<T2, N> Src0 = src0;
866  simd<T3, N> Src1 = src1;
867  simd<T4, N> Src2 = src2;
868  if constexpr (std::is_same_v<Sat, saturation_off_tag>) {
869  if constexpr (std::is_unsigned<T1>::value) {
870  if constexpr (std::is_unsigned<T2>::value) {
871  Result = __esimd_uudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
872  Src2.data());
873  } else {
874  Result = __esimd_usdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
875  Src2.data());
876  }
877  } else {
878  if constexpr (std::is_unsigned<T2>::value) {
879  Result = __esimd_sudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
880  Src2.data());
881  } else {
882  Result = __esimd_ssdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
883  Src2.data());
884  }
885  }
886  } else {
887  if constexpr (std::is_unsigned<T1>::value) {
888  if constexpr (std::is_unsigned<T2>::value) {
889  Result = __esimd_uudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
890  Src2.data());
891  } else {
892  Result = __esimd_usdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
893  Src2.data());
894  }
895  } else {
896  if constexpr (std::is_unsigned<T2>::value) {
897  Result = __esimd_sudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
898  Src2.data());
899  } else {
900  Result = __esimd_ssdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
901  Src2.data());
902  }
903  }
904  }
905  return Result;
906 #else
907  __ESIMD_UNSUPPORTED_ON_HOST;
908 #endif // __SYCL_DEVICE_ONLY__
909 }
910 
911 // reduction functions
912 namespace detail {
913 template <typename T0, typename T1, int SZ> struct esimd_apply_sum {
914  template <typename... T>
916  return v1 + v2;
917  }
918 };
919 
920 template <typename T0, typename T1, int SZ> struct esimd_apply_prod {
921  template <typename... T>
923  return v1 * v2;
924  }
925 };
926 
927 template <typename T0, typename T1, int SZ> struct esimd_apply_reduced_max {
928  template <typename... T>
930  if constexpr (std::is_floating_point<T1>::value) {
931  return __esimd_fmax<T1, SZ>(v1.data(), v2.data());
932  } else if constexpr (std::is_unsigned<T1>::value) {
933  return __esimd_umax<T1, SZ>(v1.data(), v2.data());
934  } else {
935  return __esimd_smax<T1, SZ>(v1.data(), v2.data());
936  }
937  }
938 };
939 
940 template <typename T0, typename T1, int SZ> struct esimd_apply_reduced_min {
941  template <typename... T>
943  if constexpr (std::is_floating_point<T1>::value) {
944  return __esimd_fmin<T1, SZ>(v1.data(), v2.data());
945  } else if constexpr (std::is_unsigned<T1>::value) {
946  return __esimd_umin<T1, SZ>(v1.data(), v2.data());
947  } else {
948  return __esimd_smin<T1, SZ>(v1.data(), v2.data());
949  }
950  }
951 };
952 
953 template <typename T0, typename T1, int SZ,
954  template <typename RT, typename T, int N> class OpType>
956  if constexpr (SZ == 1) {
957  return v[0];
958  } else {
959  static_assert(detail::isPowerOf2(SZ),
960  "Invaid input for reduce_single - the vector size must "
961  "be power of two.");
962  constexpr int N = SZ / 2;
963  simd<T0, N> tmp = OpType<T0, T1, N>()(v.template select<N, 1>(0),
964  v.template select<N, 1>(N));
965  return reduce_single<T0, T0, N, OpType>(tmp);
966  }
967 }
968 
969 template <typename T0, typename T1, int N1, int N2,
970  template <typename RT, typename T, int N> class OpType>
972  if constexpr (N1 == N2) {
973  simd<T0, N1> tmp = OpType<T0, T1, N1>()(v1, v2);
974  return reduce_single<T0, T0, N1, OpType>(tmp);
975  } else if constexpr (N1 < N2) {
976  simd<T0, N1> tmp1 = OpType<T0, T1, N1>()(v1, v2.template select<N1, 1>(0));
977  constexpr int N = N2 - N1;
978  using NT = simd<T0, N>;
979  NT tmp2 = convert<T0>(v2.template select<N, 1>(N1).read());
980  return reduce_pair<T0, T0, N1, N, OpType>(tmp1, tmp2);
981  } else {
982  static_assert(detail::isPowerOf2(N1),
983  "Invaid input for reduce_pair - N1 must be power of two.");
984  constexpr int N = N1 / 2;
985  simd<T0, N> tmp = OpType<T0, T1, N>()(v1.template select<N, 1>(0),
986  v1.template select<N, 1>(N));
987  using NT = simd<T0, N2>;
988  NT tmp2 = convert<T0>(v2);
989  return reduce_pair<T0, T0, N, N2, OpType>(tmp, tmp2);
990  }
991 }
992 
993 template <typename T0, typename T1, int SZ,
994  template <typename RT, typename T, int N> class OpType>
995 T0 reduce(simd<T1, SZ> v) {
996  constexpr bool isPowerOf2 = detail::isPowerOf2(SZ);
997  if constexpr (isPowerOf2) {
998  return reduce_single<T0, T1, SZ, OpType>(v);
999  } else {
1000  constexpr unsigned N1 = 1u << detail::log2<SZ>();
1001  constexpr unsigned N2 = SZ - N1;
1002 
1003  simd<T1, N1> v1 = v.template select<N1, 1>(0);
1004  simd<T1, N2> v2 = v.template select<N2, 1>(N1);
1005  return reduce_pair<T0, T1, N1, N2, OpType>(v1, v2);
1006  }
1007 };
1008 
1009 template <typename T0, typename T1, int SZ>
1010 ESIMD_INLINE ESIMD_NODEBUG T0 sum(simd<T1, SZ> v) {
1011  using TT = detail::computation_type_t<simd<T1, SZ>>;
1012  using RT = typename TT::element_type;
1013  T0 retv = reduce<RT, T1, SZ, esimd_apply_sum>(v);
1014  return retv;
1015 }
1016 
1017 template <typename T0, typename T1, int SZ>
1018 ESIMD_INLINE ESIMD_NODEBUG T0 prod(simd<T1, SZ> v) {
1019  using TT = detail::computation_type_t<simd<T1, SZ>>;
1020  using RT = typename TT::element_type;
1021  T0 retv = reduce<RT, T1, SZ, esimd_apply_prod>(v);
1022  return retv;
1023 }
1024 } // namespace detail
1026 
1034 template <typename T0, typename T1, int SZ>
1035 ESIMD_INLINE ESIMD_NODEBUG T0 hmax(simd<T1, SZ> v) {
1036  T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_max>(v);
1037  return retv;
1038 }
1039 
1047 template <typename T0, typename T1, int SZ>
1048 ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd<T1, SZ> v) {
1049  T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_min>(v);
1050  return retv;
1051 }
1052 
1066 // TODO 1) enforce BinaryOperation constraints 2) support std::minimum/maximum
1067 template <typename T0, typename T1, int SZ, typename BinaryOperation>
1068 ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
1069  if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
1070  std::plus<>>::value) {
1071  T0 retv = detail::sum<T0>(v);
1072  return retv;
1073  } else if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
1074  std::multiplies<>>::value) {
1075  T0 retv = detail::prod<T0>(v);
1076  return retv;
1077  }
1078 }
1079 
1082 
1087 enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };
1089 static constexpr bfn_t operator~(bfn_t x) {
1090  uint8_t val = static_cast<uint8_t>(x);
1091  uint8_t res = ~val;
1092  return static_cast<bfn_t>(res);
1093 }
1095 static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
1096  uint8_t arg0 = static_cast<uint8_t>(x);
1097  uint8_t arg1 = static_cast<uint8_t>(y);
1098  uint8_t res = arg0 | arg1;
1099  return static_cast<bfn_t>(res);
1100 }
1102 static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
1103  uint8_t arg0 = static_cast<uint8_t>(x);
1104  uint8_t arg1 = static_cast<uint8_t>(y);
1105  uint8_t res = arg0 & arg1;
1106  return static_cast<bfn_t>(res);
1107 }
1109 static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
1110  uint8_t arg0 = static_cast<uint8_t>(x);
1111  uint8_t arg1 = static_cast<uint8_t>(y);
1112  uint8_t res = arg0 ^ arg1;
1113  return static_cast<bfn_t>(res);
1114 }
1115 
1124 template <bfn_t FuncControl, typename T, int N>
1125 __ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
1126 bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
1127  __ESIMD_NS::simd<T, N> src2) {
1128  if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
1129  ((sizeof(T) == 2) && (N % 2 == 0))) {
1130  // Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
1131  // Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
1132  // Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
1133  auto Result = __ESIMD_NS::bfn<FuncControl>(
1134  src0.template bit_cast_view<int32_t>().read(),
1135  src1.template bit_cast_view<int32_t>().read(),
1136  src2.template bit_cast_view<int32_t>().read());
1137  return Result.template bit_cast_view<T>();
1138  } else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
1139  constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
1140  return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
1141  } else if constexpr (N % 2 == 0) {
1142  // Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
1143  auto Result = __ESIMD_NS::bfn<FuncControl>(
1144  src0.template bit_cast_view<int16_t>().read(),
1145  src1.template bit_cast_view<int16_t>().read(),
1146  src2.template bit_cast_view<int16_t>().read());
1147  return Result.template bit_cast_view<T>();
1148  } else {
1149  // Odd number of 1-byte elements.
1150  __ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
1151  Src0.template select<N, 1>() = src0;
1152  Src1.template select<N, 1>() = src1;
1153  Src2.template select<N, 1>() = src2;
1154  auto Result = __ESIMD_NS::bfn<FuncControl>(Src0, Src1, Src2);
1155  return Result.template select<N, 1>();
1156  }
1157 }
1158 
1166 template <bfn_t FuncControl, typename T>
1167 ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
1168  __ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
1169 bfn(T src0, T src1, T src2) {
1170  __ESIMD_NS::simd<T, 1> Src0 = src0;
1171  __ESIMD_NS::simd<T, 1> Src1 = src1;
1172  __ESIMD_NS::simd<T, 1> Src2 = src2;
1173  __ESIMD_NS::simd<T, 1> Result =
1174  esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
1175  return Result[0];
1176 }
1177 
1179 
1186 template <int N>
1187 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1188 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
1189  __ESIMD_NS::simd<uint32_t, N> src1) {
1190  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1191  __ESIMD_DNS::vector_type_t<uint32_t, N>>
1192  Result = __esimd_addc<uint32_t, N>(src0.data(), src1.data());
1193 
1194  carry = Result.first;
1195  return Result.second;
1196 }
1197 
1204 template <int N>
1205 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1206 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
1207  uint32_t src1) {
1208  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
1209  return addc(carry, src0, Src1V);
1210 }
1211 
1218 template <int N>
1219 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1220 addc(__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t src0,
1221  __ESIMD_NS::simd<uint32_t, N> src1) {
1222  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
1223  return addc(carry, Src0V, src1);
1224 }
1225 
1232 __ESIMD_API uint32_t addc(uint32_t &carry, uint32_t src0, uint32_t src1) {
1233  __ESIMD_NS::simd<uint32_t, 1> CarryV = carry;
1234  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
1235  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
1236  __ESIMD_NS::simd<uint32_t, 1> Res = addc(CarryV, Src0V, Src1V);
1237  carry = CarryV[0];
1238  return Res[0];
1239 }
1240 
1248 template <int N>
1249 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1250 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
1251  __ESIMD_NS::simd<uint32_t, N> src1) {
1252  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1253  __ESIMD_DNS::vector_type_t<uint32_t, N>>
1254  Result = __esimd_subb<uint32_t, N>(src0.data(), src1.data());
1255 
1256  borrow = Result.first;
1257  return Result.second;
1258 }
1259 
1267 template <int N>
1268 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1269 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
1270  uint32_t src1) {
1271  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
1272  return subb(borrow, src0, Src1V);
1273 }
1274 
1282 template <int N>
1283 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1284 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t src0,
1285  __ESIMD_NS::simd<uint32_t, N> src1) {
1286  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
1287  return subb(borrow, Src0V, src1);
1288 }
1289 
1297 __ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1) {
1298  __ESIMD_NS::simd<uint32_t, 1> BorrowV = borrow;
1299  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
1300  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
1301  __ESIMD_NS::simd<uint32_t, 1> Res = subb(BorrowV, Src0V, Src1V);
1302  borrow = BorrowV[0];
1303  return Res[0];
1304 }
1305 
1308 __ESIMD_API uint64_t rdtsc() {
1309  __ESIMD_NS::simd<uint32_t, 4> retv = __esimd_timestamp();
1310  return retv.template bit_cast_view<uint64_t>()[0];
1311 }
1312 
1314 
1315 } // namespace ext::intel::esimd
1316 } // namespace _V1
1317 } // namespace sycl
const auto & data() const noexcept
Definition: simd.hpp:1673
Definition: simd.hpp:1387
This class represents a reference to a sub-region of a base simd object.
Definition: simd_view.hpp:37
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...
Definition: simd_view.hpp:64
The main simd vector class.
Definition: simd.hpp:53
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.
Definition: math.hpp:704
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.
Definition: math.hpp:646
__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...
Definition: math.hpp:747
__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.
Definition: math.hpp:684
__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...
Definition: math.hpp:791
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.
Definition: math.hpp:660
#define __ESIMD_INTRINSIC_DEF(name)
Definition: math.hpp:510
__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).
Definition: math.hpp:554
__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).
Definition: math.hpp:545
__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).
Definition: math.hpp:614
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.
Definition: math.hpp:581
__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).
Definition: math.hpp:562
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.
Definition: math.hpp:594
__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).
Definition: math.hpp:570
__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.
Definition: math.hpp:1125
static constexpr bfn_t operator&(bfn_t x, bfn_t y)
Definition: math.hpp:1101
bfn_t
This enum is used to encode all possible logical operations performed on the 3 input operands.
Definition: math.hpp:1086
static constexpr bfn_t operator|(bfn_t x, bfn_t y)
Definition: math.hpp:1094
static constexpr bfn_t operator~(bfn_t x)
Definition: math.hpp:1088
static constexpr bfn_t operator^(bfn_t x, bfn_t y)
Definition: math.hpp:1108
__ESIMD_API simd< T, N > cos(simd< T, N > src, Sat sat={})
Cosine.
Definition: math.hpp:405
__ESIMD_API simd< T, N > sqrt(simd< T, N > src, Sat sat={})
Square root.
Definition: math.hpp:389
__ESIMD_API simd< T, N > rsqrt(simd< T, N > src, Sat sat={})
Square root reciprocal - calculates 1/sqrt(x).
Definition: math.hpp:397
#define __ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname)
Definition: math.hpp:345
__ESIMD_API simd< T, N > pow(simd< T, N > src0, simd< U, N > src1, Sat sat={})
Power - calculates src0 in power of src1.
Definition: math.hpp:442
#define __ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname)
Definition: math.hpp:409
__ESIMD_API simd< T, N > inv(simd< T, N > src, Sat sat={})
Inversion - calculates (1/x).
Definition: math.hpp:375
__ESIMD_API simd< T, N > sqrt_ieee(simd< T, N > src, Sat sat={})
IEEE754-compliant square root. Supports float and double.
Definition: math.hpp:392
#define __ESIMD_EMATH_COND
Definition: math.hpp:367
__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.
Definition: math.hpp:445
__ESIMD_API simd< T, N > sin(simd< T, N > src, Sat sat={})
Sine.
Definition: math.hpp:401
#define __ESIMD_EMATH_IEEE_COND
Definition: math.hpp:370
__ESIMD_API simd< T, N > log2(simd< T, N > src, Sat sat={})
Logarithm base 2.
Definition: math.hpp:381
__ESIMD_API simd< T, N > exp2(simd< T, N > src, Sat sat={})
Exponent base 2.
Definition: math.hpp:385
ESIMD_INLINE ESIMD_NODEBUG T0 hmax(simd< T1, SZ > v)
ESIMD_DETAIL.
Definition: math.hpp:1034
__ESIMD_API SZ simd< T, SZ > src1
Definition: math.hpp:180
__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.
Definition: math.hpp:470
__ESIMD_API SZ simd< T, SZ > Sat sat
Definition: math.hpp:180
ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd< T1, SZ > v, BinaryOperation op)
Performs reduction over elements of the input vector.
Definition: math.hpp:1067
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >() min(simd< T, SZ > src0, T src1, Sat sat={})
Definition: math.hpp:298
__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.
Definition: math.hpp:72
__ESIMD_API SZ simd< T, SZ > Sat int SZ
Definition: math.hpp:211
__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.
Definition: math.hpp:1187
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >() max(simd< T, SZ > src0, T src1, Sat sat={})
Definition: math.hpp:213
__ESIMD_API uint64_t rdtsc()
rdtsc - get the value of timestamp counter.
Definition: math.hpp:1307
ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd< T1, SZ > v)
Performs 'minimum' operation reduction over elements of the input vector, that is,...
Definition: math.hpp:1047
ESIMD_NODEBUG ESIMD_INLINE simd< T, SZ > exp(simd< T, SZ > src0, Sat sat={})
Computes e raised to the power of the given argument.
Definition: math.hpp:491
__ESIMD_API SZ simd< T, SZ > Sat int class Sat
Definition: math.hpp:211
__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.
Definition: math.hpp:1249
__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.
Definition: math.hpp:861
__ESIMD_API SZ src0
Definition: math.hpp:180
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)
Definition: math.hpp:127
T0 reduce(simd< T1, SZ > v)
Definition: math.hpp:994
ESIMD_INLINE ESIMD_NODEBUG T0 prod(simd< T1, SZ > v)
Definition: math.hpp:1017
T0 reduce_single(simd< T1, SZ > v)
Definition: math.hpp:954
ESIMD_INLINE ESIMD_NODEBUG T0 sum(simd< T1, SZ > v)
Definition: math.hpp:1009
T0 reduce_pair(simd< T1, N1 > v1, simd< T1, N2 > v2)
Definition: math.hpp:970
constexpr ESIMD_INLINE bool isPowerOf2(unsigned int n)
Check if a given 32 bit positive integer is a power of 2 at compile time.
Definition: common.hpp:96
autodecltype(x) x
Definition: access.hpp:18
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:921
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:928
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:941
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:914