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 
20 
21 #include <cstdint>
22 
23 namespace sycl {
24 inline namespace _V1 {
25 namespace ext::intel::esimd {
26 
39 
42 
67 template <typename T0, typename T1, int SZ>
68 __ESIMD_API std::enable_if_t<!detail::is_generic_floating_point_v<T0> ||
69  std::is_same_v<T1, T0>,
72  if constexpr (detail::is_generic_floating_point_v<T0>)
73  return __esimd_sat<T0, T1, SZ>(src.data());
74  else if constexpr (detail::is_generic_floating_point_v<T1>) {
75  if constexpr (std::is_unsigned_v<T0>)
76  return __esimd_fptoui_sat<T0, T1, SZ>(src.data());
77  else
78  return __esimd_fptosi_sat<T0, T1, SZ>(src.data());
79  } else if constexpr (std::is_unsigned_v<T0>) {
80  if constexpr (std::is_unsigned_v<T1>)
81  return __esimd_uutrunc_sat<T0, T1, SZ>(src.data());
82  else
83  return __esimd_ustrunc_sat<T0, T1, SZ>(src.data());
84  } else {
85  if constexpr (std::is_signed_v<T1>)
86  return __esimd_sstrunc_sat<T0, T1, SZ>(src.data());
87  else
88  return __esimd_sutrunc_sat<T0, T1, SZ>(src.data());
89  }
90 }
91 
93 // abs
94 namespace detail {
95 
96 template <typename TRes, typename TArg, int SZ>
97 ESIMD_NODEBUG ESIMD_INLINE simd<TRes, SZ>
98 __esimd_abs_common_internal(simd<TArg, SZ> src0) {
99  simd<TArg, SZ> Result = simd<TArg, SZ>(__esimd_abs<TArg, SZ>(src0.data()));
100  return convert<TRes>(Result);
101 }
102 
103 template <typename TRes, typename TArg>
104 ESIMD_NODEBUG ESIMD_INLINE
105  std::enable_if_t<detail::is_esimd_scalar<TRes>::value &&
106  detail::is_esimd_scalar<TArg>::value,
107  TRes>
108  __esimd_abs_common_internal(TArg src0) {
109  simd<TArg, 1> Src0 = src0;
110  simd<TArg, 1> Result = __esimd_abs_common_internal<TArg>(Src0);
111  return convert<TRes>(Result)[0];
112 }
113 } // namespace detail
115 
122 template <typename TRes, typename TArg, int SZ>
123 __ESIMD_API std::enable_if_t<
124  !std::is_same<std::remove_const_t<TRes>, std::remove_const_t<TArg>>::value,
127  return detail::__esimd_abs_common_internal<TRes, TArg, SZ>(src0.data());
128 }
129 
135 template <typename TRes, typename TArg>
136 __ESIMD_API std::enable_if_t<!std::is_same<std::remove_const_t<TRes>,
137  std::remove_const_t<TArg>>::value &&
138  detail::is_esimd_scalar<TRes>::value &&
139  detail::is_esimd_scalar<TArg>::value,
140  std::remove_const_t<TRes>>
141 abs(TArg src0) {
142  return detail::__esimd_abs_common_internal<TRes, TArg>(src0);
143 }
144 
152 template <typename T1, int SZ> __ESIMD_API simd<T1, SZ> abs(simd<T1, SZ> src0) {
153  return detail::__esimd_abs_common_internal<T1, T1, SZ>(src0.data());
154 }
155 
162 template <typename T1>
163 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T1>::value,
164  std::remove_const_t<T1>>
165 abs(T1 src0) {
166  return detail::__esimd_abs_common_internal<T1, T1>(src0);
167 }
168 
178 template <typename T, int SZ, class Sat = saturation_off_tag>
180  constexpr bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
181 
182  if constexpr (std::is_floating_point<T>::value) {
183  auto Result = __esimd_fmax<T, SZ>(src0.data(), src1.data());
184  if constexpr (is_sat)
185  Result = __esimd_sat<T, T, SZ>(Result);
186  return simd<T, SZ>(Result);
187  } else if constexpr (std::is_unsigned<T>::value) {
188  auto Result = __esimd_umax<T, SZ>(src0.data(), src1.data());
189  if constexpr (is_sat)
190  Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
191  return simd<T, SZ>(Result);
192  } else {
193  auto Result = __esimd_smax<T, SZ>(src0.data(), src1.data());
194  if constexpr (is_sat)
195  Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
196  return simd<T, SZ>(Result);
197  }
198 }
199 
210 template <typename T, int SZ, class Sat = saturation_off_tag>
211 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
212  max)(simd<T, SZ> src0, T src1, Sat sat = {}) {
213  simd<T, SZ> Src1 = src1;
214  simd<T, SZ> Result = (esimd::max)(src0, Src1, sat);
215  return Result;
216 }
217 
228 template <typename T, int SZ, class Sat = saturation_off_tag>
229 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
230  max)(T src0, simd<T, SZ> src1, Sat sat = {}) {
231  simd<T, SZ> Src0 = src0;
232  simd<T, SZ> Result = (esimd::max)(Src0, src1, sat);
233  return Result;
234 }
235 
244 template <typename T, class Sat = saturation_off_tag>
245 ESIMD_NODEBUG ESIMD_INLINE
246 std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(max)(T src0, T src1,
247  Sat sat = {}) {
248  simd<T, 1> Src0 = src0;
249  simd<T, 1> Src1 = src1;
250  simd<T, 1> Result = (esimd::max)(Src0, Src1, sat);
251  return Result[0];
252 }
253 
263 template <typename T, int SZ, class Sat = saturation_off_tag>
265  constexpr bool is_sat = std::is_same_v<Sat, saturation_on_tag>;
266 
267  if constexpr (std::is_floating_point<T>::value) {
268  auto Result = __esimd_fmin<T, SZ>(src0.data(), src1.data());
269  if constexpr (is_sat)
270  Result = __esimd_sat<T, T, SZ>(Result);
271  return simd<T, SZ>(Result);
272  } else if constexpr (std::is_unsigned<T>::value) {
273  auto Result = __esimd_umin<T, SZ>(src0.data(), src1.data());
274  if constexpr (is_sat)
275  Result = __esimd_uutrunc_sat<T, T, SZ>(Result);
276  return simd<T, SZ>(Result);
277  } else {
278  auto Result = __esimd_smin<T, SZ>(src0.data(), src1.data());
279  if constexpr (is_sat)
280  Result = __esimd_sstrunc_sat<T, T, SZ>(Result);
281  return simd<T, SZ>(Result);
282  }
283 }
284 
295 template <typename T, int SZ, class Sat = saturation_off_tag>
296 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
297  min)(simd<T, SZ> src0, T src1, Sat sat = {}) {
298  simd<T, SZ> Src1 = src1;
299  simd<T, SZ> Result = (esimd::min)(src0, Src1, sat);
300  return Result;
301 }
302 
313 template <typename T, int SZ, class Sat = saturation_off_tag>
314 __ESIMD_API std::enable_if_t<detail::is_esimd_scalar<T>::value, simd<T, SZ>>(
315  min)(T src0, simd<T, SZ> src1, Sat sat = {}) {
316  simd<T, SZ> Src0 = src0;
317  simd<T, SZ> Result = (esimd::min)(Src0, src1, sat);
318  return Result;
319 }
320 
329 template <typename T, class Sat = saturation_off_tag>
330 ESIMD_NODEBUG ESIMD_INLINE
331 std::enable_if_t<detail::is_esimd_scalar<T>::value, T>(min)(T src0, T src1,
332  Sat sat = {}) {
333  simd<T, 1> Src0 = src0;
334  simd<T, 1> Src1 = src1;
335  simd<T, 1> Result = (esimd::min)(Src0, Src1, sat);
336  return Result[0];
337 }
338 
340 
343 
344 #define __ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname) \
345  \
346  template <class T, int N, class Sat = saturation_off_tag, \
347  class = std::enable_if_t<COND>> \
348  __ESIMD_API simd<T, N> name(simd<T, N> src, Sat sat = {}) { \
349  __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N> res = \
350  __esimd_##iname<T, N>(src.data()); \
351  if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
352  return res; \
353  else \
354  return esimd::saturate<T>(simd<T, N>(res)); \
355  } \
356  \
357  \
358  template <typename T, class Sat = saturation_off_tag, \
359  class = std::enable_if_t<COND>> \
360  __ESIMD_API T name(T src, Sat sat = {}) { \
361  simd<T, 1> src_vec = src; \
362  simd<T, 1> res = name<T, 1>(src_vec, sat); \
363  return res[0]; \
364  }
365 
366 #define __ESIMD_EMATH_COND \
367  detail::is_generic_floating_point_v<T> && (sizeof(T) <= 4)
368 
369 #define __ESIMD_EMATH_IEEE_COND \
370  detail::is_generic_floating_point_v<T> && (sizeof(T) >= 4)
371 
375 
381 
385 
389 
392 
397 
401 
405 
406 #undef __ESIMD_UNARY_INTRINSIC_DEF
407 
408 #define __ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname) \
409  \
410  template <class T, int N, class U, class Sat = saturation_off_tag, \
411  class = std::enable_if_t<COND>> \
412  __ESIMD_API simd<T, N> name(simd<T, N> src0, simd<U, N> src1, \
413  Sat sat = {}) { \
414  using RawVecT = __ESIMD_DNS::vector_type_t<__ESIMD_DNS::__raw_t<T>, N>; \
415  RawVecT src1_raw_conv = detail::convert_vector<T, U, N>(src1.data()); \
416  RawVecT res_raw = __esimd_##iname<T, N>(src0.data(), src1_raw_conv); \
417  if constexpr (std::is_same_v<Sat, saturation_off_tag>) \
418  return res_raw; \
419  else \
420  return esimd::saturate<T>(simd<T, N>(res_raw)); \
421  } \
422  \
423  \
424  template <class T, int N, class U, class Sat = saturation_off_tag, \
425  class = std::enable_if_t<COND>> \
426  __ESIMD_API simd<T, N> name(simd<T, N> src0, U src1, Sat sat = {}) { \
427  return name<T, N, U>(src0, simd<U, N>(src1), sat); \
428  } \
429  \
430  \
431  template <class T, class U, class Sat = saturation_off_tag, \
432  class = std::enable_if_t<COND>> \
433  __ESIMD_API T name(T src0, U src1, Sat sat = {}) { \
434  simd<T, 1> res = name<T, 1, U>(simd<T, 1>(src0), simd<U, 1>(src1), sat); \
435  return res[0]; \
436  }
437 
442 
445 
446 #undef __ESIMD_BINARY_INTRINSIC_DEF
447 #undef __ESIMD_EMATH_COND
448 #undef __ESIMD_EMATH_IEEE_COND
449 
451 
454 
456 namespace detail {
457 // std::numbers::ln2_v<float> in c++20
458 constexpr float ln2 = 0.69314718f;
459 // std::numbers::log2e_v<float> in c++20
460 constexpr float log2e = 1.442695f;
461 } // namespace detail
463 
468 template <class T, int SZ, class Sat = saturation_off_tag>
469 ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> log(simd<T, SZ> src0, Sat sat = {}) {
470  using CppT = __ESIMD_DNS::__cpp_t<T>;
471  simd<T, SZ> Result =
472  esimd::log2<T, SZ, saturation_off_tag>(src0) * detail::ln2;
473 
474  if constexpr (std::is_same_v<Sat, saturation_off_tag>)
475  return Result;
476  else
477  return esimd::saturate<T>(Result);
478 }
479 
480 template <class T, class Sat = saturation_off_tag>
481 ESIMD_NODEBUG ESIMD_INLINE T log(T src0, Sat sat = {}) {
482  return esimd::log<T, 1>(src0, sat)[0];
483 }
484 
489 template <class T, int SZ, class Sat = saturation_off_tag>
490 ESIMD_NODEBUG ESIMD_INLINE simd<T, SZ> exp(simd<T, SZ> src0, Sat sat = {}) {
491  using CppT = __ESIMD_DNS::__cpp_t<T>;
492  return esimd::exp2<T, SZ>(src0 * detail::log2e, sat);
493 }
494 
495 template <class T, class Sat = saturation_off_tag>
496 ESIMD_NODEBUG ESIMD_INLINE T exp(T src0, Sat sat = {}) {
497  return esimd::exp<T, 1>(src0, sat)[0];
498 }
499 
501 
504 
506 // Rounding intrinsics.
508 
509 #define __ESIMD_INTRINSIC_DEF(name) \
510  \
511  \
512  \
514  \
515  \
516  \
517  template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag> \
518  __ESIMD_API __ESIMD_NS::simd<T, SZ> name(__ESIMD_NS::simd<float, SZ> src0, \
519  Sat sat = {}) { \
520  __ESIMD_NS::simd<float, SZ> Result = __esimd_##name<SZ>(src0.data()); \
521  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>) \
522  return Result; \
523  else if constexpr (!std::is_same_v<float, T>) { \
524  auto RawRes = __ESIMD_NS::saturate<float>(Result).data(); \
525  return __ESIMD_DNS::convert_vector<T, float, SZ>(std::move(RawRes)); \
526  } else { \
527  return __ESIMD_NS::saturate<T>(Result); \
528  } \
529  } \
530  \
531  template <typename T, class Sat = __ESIMD_NS::saturation_off_tag> \
532  __ESIMD_API T name(float src0, Sat sat = {}) { \
533  __ESIMD_NS::simd<float, 1> Src0 = src0; \
534  __ESIMD_NS::simd<T, 1> Result = name<T>(Src0, sat); \
535  return Result[0]; \
536  }
537 
546 
555 
563 
571 
572 #undef __ESIMD_INTRINSIC_DEF
574 
577 
579 template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
580 ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
581 floor(const __ESIMD_NS::simd<float, SZ> src0, Sat sat = {}) {
582  return esimd::rndd<RT, SZ>(src0, sat);
583 }
584 
586 template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
587 ESIMD_INLINE RT floor(float src0, Sat sat = {}) {
588  return esimd::rndd<RT, 1U>(src0, sat)[0];
589 }
590 
592 template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
593 ESIMD_INLINE __ESIMD_NS::simd<RT, SZ>
594 ceil(const __ESIMD_NS::simd<float, SZ> src0, Sat sat = {}) {
595  return esimd::rndu<RT, SZ>(src0, sat);
596 }
597 
599 template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
600 ESIMD_INLINE RT ceil(float src0, Sat sat = {}) {
601  return esimd::rndu<RT, 1U>(src0, sat);
602 }
603 
612 template <typename RT, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
613 __ESIMD_API __ESIMD_NS::simd<RT, SZ>
614 trunc(const __ESIMD_NS::simd<float, SZ> &src0, Sat sat = {}) {
615  return esimd::rndz<RT, SZ>(src0, sat);
616 }
617 
625 template <typename RT, class Sat = __ESIMD_NS::saturation_off_tag>
626 __ESIMD_API RT trunc(float src0, Sat sat = {}) {
627  return esimd::rndz<RT, 1U>(src0, sat)[0];
628 }
629 
631 
634 
643 template <int N>
644 ESIMD_NODEBUG ESIMD_INLINE
645  std::enable_if_t<(N == 8 || N == 16 || N == 32), uint>
647  return __esimd_pack_mask<N>(src0.data());
648 }
649 
657 template <int N>
658 ESIMD_NODEBUG ESIMD_INLINE
659  std::enable_if_t<(N == 8 || N == 16 || N == 32), simd_mask<N>>
661  return __esimd_unpack_mask<N>(src0);
662 }
663 
666 template <int N>
667 __ESIMD_API std::enable_if_t<(N != 8 && N != 16 && N < 32), uint>
669  simd_mask<(N < 8 ? 8 : N < 16 ? 16 : 32)> src_0 = 0;
670  src_0.template select<N, 1>() = src0.template bit_cast_view<ushort>();
671  return esimd::pack_mask(src_0);
672 }
673 
680 template <typename T, int N>
681 __ESIMD_API std::enable_if_t<(std::is_same_v<T, ushort> ||
682  std::is_same_v<T, uint>)&&(N > 0 && N <= 32),
684 ballot(simd<T, N> mask) {
685  simd_mask<N> cmp = (mask != 0);
686  if constexpr (N == 8 || N == 16 || N == 32) {
687  return __esimd_pack_mask<N>(cmp.data());
688  } else {
689  constexpr int N1 = (N <= 8 ? 8 : N <= 16 ? 16 : 32);
690  simd<uint16_t, N1> res = 0;
691  res.template select<N, 1>() = cmp.data();
692  return __esimd_pack_mask<N1>(res.data());
693  }
694 }
695 
700 template <typename T, int N>
701 ESIMD_NODEBUG ESIMD_INLINE
702  std::enable_if_t<std::is_integral<T>::value && (sizeof(T) <= 4),
704  cbit(simd<T, N> src) {
705  return __esimd_cbit<T, N>(src.data());
706 }
707 
710 template <typename T>
711 __ESIMD_API
712  std::enable_if_t<std::is_integral<T>::value && (sizeof(T) <= 4), uint32_t>
713  cbit(T src) {
714  simd<T, 1> Src = src;
715  simd<uint32_t, 1> Result = esimd::cbit(Src);
716  return Result[0];
717 }
718 
723 template <typename BaseTy, typename RegionTy>
724 __ESIMD_API std::enable_if_t<
725  std::is_integral<
727  (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) <= 4) &&
729  uint32_t>
731  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
732  simd<Ty, 1> Src = src;
733  simd<uint32_t, 1> Result = esimd::cbit(Src);
734  return Result[0];
735 }
736 
744 template <typename T, int N>
745 __ESIMD_API
746  std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), simd<T, N>>
747  fbl(simd<T, N> src) {
748  return __esimd_fbl<T, N>(src.data());
749 }
750 
753 template <typename T>
754 __ESIMD_API std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), T>
755 fbl(T src) {
756  simd<T, 1> Src = src;
757  simd<T, 1> Result = esimd::fbl(Src);
758  return Result[0];
759 }
760 
766 template <typename BaseTy, typename RegionTy>
767 __ESIMD_API std::enable_if_t<
768  std::is_integral<
770  (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) == 4) &&
774  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
775  simd<Ty, 1> Src = src;
776  simd<Ty, 1> Result = esimd::fbl(Src);
777  return Result[0];
778 }
779 
787 template <typename T, int N>
788 __ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
789  std::is_signed<T>::value && (sizeof(T) == 4),
791 fbh(simd<T, N> src) {
792  return __esimd_sfbh<T, N>(src.data());
793 }
794 
801 template <typename T, int N>
802 __ESIMD_API std::enable_if_t<std::is_integral<T>::value &&
803  !std::is_signed<T>::value && (sizeof(T) == 4),
805 fbh(simd<T, N> src) {
806  return __esimd_ufbh<T, N>(src.data());
807 }
808 
811 template <typename T>
812 __ESIMD_API std::enable_if_t<std::is_integral<T>::value && (sizeof(T) == 4), T>
813 fbh(T src) {
814  simd<T, 1> Src = src;
815  simd<T, 1> Result = esimd::fbh(Src);
816  return Result[0];
817 }
818 
824 template <typename BaseTy, typename RegionTy>
825 __ESIMD_API std::enable_if_t<
826  std::is_integral<
828  (sizeof(typename simd_view<BaseTy, RegionTy>::element_type) == 4) &&
832  using Ty = typename simd_view<BaseTy, RegionTy>::element_type;
833  simd<Ty, 1> Src = src;
834  simd<Ty, 1> Result = esimd::fbh(Src);
835  return Result[0];
836 }
837 
839 
842 
855 template <typename T1, typename T2, typename T3, typename T4, int N,
856  class Sat = saturation_off_tag>
857 __ESIMD_API std::enable_if_t<
858  detail::is_dword_type<T1>::value && detail::is_dword_type<T2>::value &&
859  detail::is_dword_type<T3>::value && detail::is_dword_type<T4>::value,
862 #if defined(__SYCL_DEVICE_ONLY__)
863  simd<T1, N> Result;
864  simd<T2, N> Src0 = src0;
865  simd<T3, N> Src1 = src1;
866  simd<T4, N> Src2 = src2;
867  if constexpr (std::is_same_v<Sat, saturation_off_tag>) {
868  if constexpr (std::is_unsigned<T1>::value) {
869  if constexpr (std::is_unsigned<T2>::value) {
870  Result = __esimd_uudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
871  Src2.data());
872  } else {
873  Result = __esimd_usdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
874  Src2.data());
875  }
876  } else {
877  if constexpr (std::is_unsigned<T2>::value) {
878  Result = __esimd_sudp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
879  Src2.data());
880  } else {
881  Result = __esimd_ssdp4a<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
882  Src2.data());
883  }
884  }
885  } else {
886  if constexpr (std::is_unsigned<T1>::value) {
887  if constexpr (std::is_unsigned<T2>::value) {
888  Result = __esimd_uudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
889  Src2.data());
890  } else {
891  Result = __esimd_usdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
892  Src2.data());
893  }
894  } else {
895  if constexpr (std::is_unsigned<T2>::value) {
896  Result = __esimd_sudp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
897  Src2.data());
898  } else {
899  Result = __esimd_ssdp4a_sat<T1, T2, T3, T4, N>(Src0.data(), Src1.data(),
900  Src2.data());
901  }
902  }
903  }
904  return Result;
905 #else
906  __ESIMD_UNSUPPORTED_ON_HOST;
907 #endif // __SYCL_DEVICE_ONLY__
908 }
909 
910 // reduction functions
911 namespace detail {
912 template <typename T0, typename T1, int SZ> struct esimd_apply_sum {
913  template <typename... T>
915  return v1 + v2;
916  }
917 };
918 
919 template <typename T0, typename T1, int SZ> struct esimd_apply_prod {
920  template <typename... T>
922  return v1 * v2;
923  }
924 };
925 
926 template <typename T0, typename T1, int SZ> struct esimd_apply_reduced_max {
927  template <typename... T>
929  if constexpr (std::is_floating_point<T1>::value) {
930  return __esimd_fmax<T1, SZ>(v1.data(), v2.data());
931  } else if constexpr (std::is_unsigned<T1>::value) {
932  return __esimd_umax<T1, SZ>(v1.data(), v2.data());
933  } else {
934  return __esimd_smax<T1, SZ>(v1.data(), v2.data());
935  }
936  }
937 };
938 
939 template <typename T0, typename T1, int SZ> struct esimd_apply_reduced_min {
940  template <typename... T>
942  if constexpr (std::is_floating_point<T1>::value) {
943  return __esimd_fmin<T1, SZ>(v1.data(), v2.data());
944  } else if constexpr (std::is_unsigned<T1>::value) {
945  return __esimd_umin<T1, SZ>(v1.data(), v2.data());
946  } else {
947  return __esimd_smin<T1, SZ>(v1.data(), v2.data());
948  }
949  }
950 };
951 
952 template <typename T0, typename T1, int SZ,
953  template <typename RT, typename T, int N> class OpType>
955  if constexpr (SZ == 1) {
956  return v[0];
957  } else {
958  static_assert(detail::isPowerOf2(SZ),
959  "Invaid input for reduce_single - the vector size must "
960  "be power of two.");
961  constexpr int N = SZ / 2;
962  simd<T0, N> tmp = OpType<T0, T1, N>()(v.template select<N, 1>(0),
963  v.template select<N, 1>(N));
964  return reduce_single<T0, T0, N, OpType>(tmp);
965  }
966 }
967 
968 template <typename T0, typename T1, int N1, int N2,
969  template <typename RT, typename T, int N> class OpType>
971  if constexpr (N1 == N2) {
972  simd<T0, N1> tmp = OpType<T0, T1, N1>()(v1, v2);
973  return reduce_single<T0, T0, N1, OpType>(tmp);
974  } else if constexpr (N1 < N2) {
975  simd<T0, N1> tmp1 = OpType<T0, T1, N1>()(v1, v2.template select<N1, 1>(0));
976  constexpr int N = N2 - N1;
977  using NT = simd<T0, N>;
978  NT tmp2 = convert<T0>(v2.template select<N, 1>(N1).read());
979  return reduce_pair<T0, T0, N1, N, OpType>(tmp1, tmp2);
980  } else {
981  static_assert(detail::isPowerOf2(N1),
982  "Invaid input for reduce_pair - N1 must be power of two.");
983  constexpr int N = N1 / 2;
984  simd<T0, N> tmp = OpType<T0, T1, N>()(v1.template select<N, 1>(0),
985  v1.template select<N, 1>(N));
986  using NT = simd<T0, N2>;
987  NT tmp2 = convert<T0>(v2);
988  return reduce_pair<T0, T0, N, N2, OpType>(tmp, tmp2);
989  }
990 }
991 
992 template <typename T0, typename T1, int SZ,
993  template <typename RT, typename T, int N> class OpType>
994 T0 reduce(simd<T1, SZ> v) {
995  constexpr bool isPowerOf2 = detail::isPowerOf2(SZ);
996  if constexpr (isPowerOf2) {
997  return reduce_single<T0, T1, SZ, OpType>(v);
998  } else {
999  constexpr unsigned N1 = 1u << detail::log2<SZ>();
1000  constexpr unsigned N2 = SZ - N1;
1001 
1002  simd<T1, N1> v1 = v.template select<N1, 1>(0);
1003  simd<T1, N2> v2 = v.template select<N2, 1>(N1);
1004  return reduce_pair<T0, T1, N1, N2, OpType>(v1, v2);
1005  }
1006 };
1007 
1008 template <typename T0, typename T1, int SZ>
1009 ESIMD_INLINE ESIMD_NODEBUG T0 sum(simd<T1, SZ> v) {
1010  using TT = detail::computation_type_t<simd<T1, SZ>>;
1011  using RT = typename TT::element_type;
1012  T0 retv = reduce<RT, T1, SZ, esimd_apply_sum>(v);
1013  return retv;
1014 }
1015 
1016 template <typename T0, typename T1, int SZ>
1017 ESIMD_INLINE ESIMD_NODEBUG T0 prod(simd<T1, SZ> v) {
1018  using TT = detail::computation_type_t<simd<T1, SZ>>;
1019  using RT = typename TT::element_type;
1020  T0 retv = reduce<RT, T1, SZ, esimd_apply_prod>(v);
1021  return retv;
1022 }
1023 } // namespace detail
1025 
1033 template <typename T0, typename T1, int SZ>
1034 ESIMD_INLINE ESIMD_NODEBUG T0 hmax(simd<T1, SZ> v) {
1035  T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_max>(v);
1036  return retv;
1037 }
1038 
1046 template <typename T0, typename T1, int SZ>
1047 ESIMD_INLINE ESIMD_NODEBUG T0 hmin(simd<T1, SZ> v) {
1048  T0 retv = detail::reduce<T1, T1, SZ, detail::esimd_apply_reduced_min>(v);
1049  return retv;
1050 }
1051 
1065 // TODO 1) enforce BinaryOperation constraints 2) support std::minimum/maximum
1066 template <typename T0, typename T1, int SZ, typename BinaryOperation>
1067 ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd<T1, SZ> v, BinaryOperation op) {
1068  if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
1069  std::plus<>>::value) {
1070  T0 retv = detail::sum<T0>(v);
1071  return retv;
1072  } else if constexpr (std::is_same<detail::remove_cvref_t<BinaryOperation>,
1073  std::multiplies<>>::value) {
1074  T0 retv = detail::prod<T0>(v);
1075  return retv;
1076  }
1077 }
1078 
1081 
1086 enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };
1088 static constexpr bfn_t operator~(bfn_t x) {
1089  uint8_t val = static_cast<uint8_t>(x);
1090  uint8_t res = ~val;
1091  return static_cast<bfn_t>(res);
1092 }
1094 static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
1095  uint8_t arg0 = static_cast<uint8_t>(x);
1096  uint8_t arg1 = static_cast<uint8_t>(y);
1097  uint8_t res = arg0 | arg1;
1098  return static_cast<bfn_t>(res);
1099 }
1101 static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
1102  uint8_t arg0 = static_cast<uint8_t>(x);
1103  uint8_t arg1 = static_cast<uint8_t>(y);
1104  uint8_t res = arg0 & arg1;
1105  return static_cast<bfn_t>(res);
1106 }
1108 static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
1109  uint8_t arg0 = static_cast<uint8_t>(x);
1110  uint8_t arg1 = static_cast<uint8_t>(y);
1111  uint8_t res = arg0 ^ arg1;
1112  return static_cast<bfn_t>(res);
1113 }
1114 
1123 template <bfn_t FuncControl, typename T, int N>
1124 __ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
1125 bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
1126  __ESIMD_NS::simd<T, N> src2) {
1127  if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
1128  ((sizeof(T) == 2) && (N % 2 == 0))) {
1129  // Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
1130  // Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
1131  // Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
1132  auto Result = __ESIMD_NS::bfn<FuncControl>(
1133  src0.template bit_cast_view<int32_t>().read(),
1134  src1.template bit_cast_view<int32_t>().read(),
1135  src2.template bit_cast_view<int32_t>().read());
1136  return Result.template bit_cast_view<T>();
1137  } else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
1138  constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
1139  return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
1140  } else if constexpr (N % 2 == 0) {
1141  // Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
1142  auto Result = __ESIMD_NS::bfn<FuncControl>(
1143  src0.template bit_cast_view<int16_t>().read(),
1144  src1.template bit_cast_view<int16_t>().read(),
1145  src2.template bit_cast_view<int16_t>().read());
1146  return Result.template bit_cast_view<T>();
1147  } else {
1148  // Odd number of 1-byte elements.
1149  __ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
1150  Src0.template select<N, 1>() = src0;
1151  Src1.template select<N, 1>() = src1;
1152  Src2.template select<N, 1>() = src2;
1153  auto Result = __ESIMD_NS::bfn<FuncControl>(Src0, Src1, Src2);
1154  return Result.template select<N, 1>();
1155  }
1156 }
1157 
1165 template <bfn_t FuncControl, typename T>
1166 ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
1167  __ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
1168 bfn(T src0, T src1, T src2) {
1169  __ESIMD_NS::simd<T, 1> Src0 = src0;
1170  __ESIMD_NS::simd<T, 1> Src1 = src1;
1171  __ESIMD_NS::simd<T, 1> Src2 = src2;
1172  __ESIMD_NS::simd<T, 1> Result =
1173  esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
1174  return Result[0];
1175 }
1176 
1178 
1185 template <int N>
1186 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1187 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
1188  __ESIMD_NS::simd<uint32_t, N> src1) {
1189  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1190  __ESIMD_DNS::vector_type_t<uint32_t, N>>
1191  Result = __esimd_addc<uint32_t, N>(src0.data(), src1.data());
1192 
1193  carry = Result.first;
1194  return Result.second;
1195 }
1196 
1203 template <int N>
1204 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1205 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
1206  uint32_t src1) {
1207  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
1208  return addc(carry, src0, Src1V);
1209 }
1210 
1217 template <int N>
1218 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1219 addc(__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t src0,
1220  __ESIMD_NS::simd<uint32_t, N> src1) {
1221  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
1222  return addc(carry, Src0V, src1);
1223 }
1224 
1231 __ESIMD_API uint32_t addc(uint32_t &carry, uint32_t src0, uint32_t src1) {
1232  __ESIMD_NS::simd<uint32_t, 1> CarryV = carry;
1233  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
1234  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
1235  __ESIMD_NS::simd<uint32_t, 1> Res = addc(CarryV, Src0V, Src1V);
1236  carry = CarryV[0];
1237  return Res[0];
1238 }
1239 
1247 template <int N>
1248 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1249 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
1250  __ESIMD_NS::simd<uint32_t, N> src1) {
1251  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
1252  __ESIMD_DNS::vector_type_t<uint32_t, N>>
1253  Result = __esimd_subb<uint32_t, N>(src0.data(), src1.data());
1254 
1255  borrow = Result.first;
1256  return Result.second;
1257 }
1258 
1266 template <int N>
1267 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1268 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
1269  uint32_t src1) {
1270  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
1271  return subb(borrow, src0, Src1V);
1272 }
1273 
1281 template <int N>
1282 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
1283 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t src0,
1284  __ESIMD_NS::simd<uint32_t, N> src1) {
1285  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
1286  return subb(borrow, Src0V, src1);
1287 }
1288 
1296 __ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1) {
1297  __ESIMD_NS::simd<uint32_t, 1> BorrowV = borrow;
1298  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
1299  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
1300  __ESIMD_NS::simd<uint32_t, 1> Res = subb(BorrowV, Src0V, Src1V);
1301  borrow = BorrowV[0];
1302  return Res[0];
1303 }
1304 
1306 
1307 } // namespace ext::intel::esimd
1308 } // namespace _V1
1309 } // 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:703
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:645
__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:746
__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:683
__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:790
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:659
#define __ESIMD_INTRINSIC_DEF(name)
Definition: math.hpp:509
__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:553
__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:544
__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:613
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:580
__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:561
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:593
__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:569
__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:1124
static constexpr bfn_t operator&(bfn_t x, bfn_t y)
Definition: math.hpp:1100
bfn_t
This enum is used to encode all possible logical operations performed on the 3 input operands.
Definition: math.hpp:1085
static constexpr bfn_t operator|(bfn_t x, bfn_t y)
Definition: math.hpp:1093
static constexpr bfn_t operator~(bfn_t x)
Definition: math.hpp:1087
static constexpr bfn_t operator^(bfn_t x, bfn_t y)
Definition: math.hpp:1107
__ESIMD_API simd< T, N > cos(simd< T, N > src, Sat sat={})
Cosine.
Definition: math.hpp:404
__ESIMD_API simd< T, N > sqrt(simd< T, N > src, Sat sat={})
Square root.
Definition: math.hpp:388
__ESIMD_API simd< T, N > rsqrt(simd< T, N > src, Sat sat={})
Square root reciprocal - calculates 1/sqrt(x).
Definition: math.hpp:396
#define __ESIMD_UNARY_INTRINSIC_DEF(COND, name, iname)
Definition: math.hpp:344
__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:441
#define __ESIMD_BINARY_INTRINSIC_DEF(COND, name, iname)
Definition: math.hpp:408
__ESIMD_API simd< T, N > inv(simd< T, N > src, Sat sat={})
Inversion - calculates (1/x).
Definition: math.hpp:374
__ESIMD_API simd< T, N > sqrt_ieee(simd< T, N > src, Sat sat={})
IEEE754-compliant square root. Supports float and double.
Definition: math.hpp:391
#define __ESIMD_EMATH_COND
Definition: math.hpp:366
__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:444
__ESIMD_API simd< T, N > sin(simd< T, N > src, Sat sat={})
Sine.
Definition: math.hpp:400
#define __ESIMD_EMATH_IEEE_COND
Definition: math.hpp:369
__ESIMD_API simd< T, N > log2(simd< T, N > src, Sat sat={})
Logarithm base 2.
Definition: math.hpp:380
__ESIMD_API simd< T, N > exp2(simd< T, N > src, Sat sat={})
Exponent base 2.
Definition: math.hpp:384
ESIMD_INLINE ESIMD_NODEBUG T0 hmax(simd< T1, SZ > v)
ESIMD_DETAIL.
Definition: math.hpp:1033
__ESIMD_API SZ simd< T, SZ > src1
Definition: math.hpp:179
__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:469
__ESIMD_API SZ simd< T, SZ > Sat sat
Definition: math.hpp:179
ESIMD_INLINE ESIMD_NODEBUG T0 reduce(simd< T1, SZ > v, BinaryOperation op)
Performs reduction over elements of the input vector.
Definition: math.hpp:1066
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >() min(simd< T, SZ > src0, T src1, Sat sat={})
Definition: math.hpp:297
__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:71
__ESIMD_API SZ simd< T, SZ > Sat int SZ
Definition: math.hpp:210
__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:1186
__ESIMD_API SZ simd< T, SZ > Sat int class simd< T, SZ >() max(simd< T, SZ > src0, T src1, Sat sat={})
Definition: math.hpp:212
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:1046
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:490
__ESIMD_API SZ simd< T, SZ > Sat int class Sat
Definition: math.hpp:210
__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:1248
__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:860
__ESIMD_API SZ src0
Definition: math.hpp:179
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:126
T0 reduce(simd< T1, SZ > v)
Definition: math.hpp:993
ESIMD_INLINE ESIMD_NODEBUG T0 prod(simd< T1, SZ > v)
Definition: math.hpp:1016
T0 reduce_single(simd< T1, SZ > v)
Definition: math.hpp:953
ESIMD_INLINE ESIMD_NODEBUG T0 sum(simd< T1, SZ > v)
Definition: math.hpp:1008
T0 reduce_pair(simd< T1, N1 > v1, simd< T1, N2 > v2)
Definition: math.hpp:969
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:95
std::conditional_t< std::is_same_v< ElementType, half >, sycl::detail::half_impl::BIsRepresentationT, ElementType > element_type
Definition: multi_ptr.hpp:752
Definition: access.hpp:18
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:920
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:927
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:940
simd< T0, SZ > operator()(simd< T1, SZ > v1, simd< T1, SZ > v2)
Definition: math.hpp:913