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 experimental Explicit SIMD math APIs.
9 //===----------------------------------------------------------------------===//
10 
11 #pragma once
12 
13 #include <sycl/ext/intel/esimd/detail/defines_elementary.hpp>
14 #include <sycl/ext/intel/esimd/detail/types.hpp>
15 #include <sycl/ext/intel/esimd/math.hpp>
16 #include <sycl/ext/intel/experimental/esimd/common.hpp>
17 #include <sycl/ext/intel/experimental/esimd/detail/math_intrin.hpp>
18 #include <sycl/ext/intel/experimental/esimd/detail/util.hpp>
19 
20 namespace sycl {
21 __SYCL_INLINE_VER_NAMESPACE(_V1) {
22 namespace ext::intel::experimental::esimd {
23 
26 
37 template <typename T0, typename T1, int SZ, typename U,
38  class Sat = __ESIMD_NS::saturation_off_tag>
39 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
40  std::is_integral<T1>::value &&
41  std::is_integral<U>::value,
42  __ESIMD_NS::simd<T0, SZ>>
43 shl(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
44  using ComputationTy =
45  __ESIMD_DNS::computation_type_t<decltype(src0), int32_t>;
46  ComputationTy Src0 = src0;
47  ComputationTy Src1 = src1;
48 
49  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_on_tag>) {
50  if constexpr (std::is_unsigned<T0>::value) {
51  if constexpr (std::is_unsigned<
52  typename ComputationTy::element_type>::value)
53  return __esimd_uushl_sat<T0, typename ComputationTy::element_type, SZ>(
54  Src0.data(), Src1.data());
55  else
56  return __esimd_usshl_sat<T0, typename ComputationTy::element_type, SZ>(
57  Src0.data(), Src1.data());
58  } else {
59  if constexpr (std::is_signed<typename ComputationTy::element_type>::value)
60  return __esimd_sushl_sat<T0, typename ComputationTy::element_type, SZ>(
61  Src0.data(), Src1.data());
62  else
63  return __esimd_ssshl_sat<T0, typename ComputationTy::element_type, SZ>(
64  Src0.data(), Src1.data());
65  }
66  } else {
67  if constexpr (std::is_unsigned<T0>::value) {
68  if constexpr (std::is_unsigned<
69  typename ComputationTy::element_type>::value)
70  return __esimd_uushl<T0, typename ComputationTy::element_type, SZ>(
71  Src0.data(), Src1.data());
72  else
73  return __esimd_usshl<T0, typename ComputationTy::element_type, SZ>(
74  Src0.data(), Src1.data());
75  } else {
76  if constexpr (std::is_signed<typename ComputationTy::element_type>::value)
77  return __esimd_sushl<T0, typename ComputationTy::element_type, SZ>(
78  Src0.data(), Src1.data());
79  else
80  return __esimd_ssshl<T0, typename ComputationTy::element_type, SZ>(
81  Src0.data(), Src1.data());
82  }
83  }
84 }
85 
95 template <typename T0, typename T1, typename T2,
96  class Sat = __ESIMD_NS::saturation_off_tag>
97 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
98  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
99  __ESIMD_DNS::is_esimd_scalar<T2>::value &&
100  std::is_integral<T0>::value &&
101  std::is_integral<T1>::value &&
102  std::is_integral<T2>::value,
103  std::remove_const_t<T0>>
104 shl(T1 src0, T2 src1, Sat sat = {}) {
105  __ESIMD_NS::simd<T1, 1> Src0 = src0;
106  __ESIMD_NS::simd<T0, 1> Result =
107  esimd::shl<T0, T1, 1, T2, Sat>(Src0, src1, sat);
108  return Result[0];
109 }
110 
121 template <typename T0, typename T1, int SZ, typename U,
122  class Sat = __ESIMD_NS::saturation_off_tag>
123 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
124  std::is_integral<T1>::value &&
125  std::is_integral<U>::value,
126  __ESIMD_NS::simd<T0, SZ>>
127 lsr(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
128  using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
129  typedef typename std::make_unsigned<IntermedTy>::type ComputationTy;
130  __ESIMD_NS::simd<ComputationTy, SZ> Src0 = src0;
131  __ESIMD_NS::simd<ComputationTy, SZ> Src1 = src1;
132  // TODO H/W supports saturation with this op - map to more efficient version.
133  __ESIMD_NS::simd<ComputationTy, SZ> Result = Src0.data() >> Src1.data();
134 
135  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
136  return Result;
137  else
138  return __ESIMD_NS::saturate<T0>(Result);
139 }
140 
151 template <typename T0, typename T1, typename T2,
152  class Sat = __ESIMD_NS::saturation_off_tag>
153 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
154  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
155  __ESIMD_DNS::is_esimd_scalar<T2>::value &&
156  std::is_integral<T0>::value &&
157  std::is_integral<T1>::value &&
158  std::is_integral<T2>::value,
159  std::remove_const_t<T0>>
160 lsr(T1 src0, T2 src1, Sat sat = {}) {
161  __ESIMD_NS::simd<T1, 1> Src0 = src0;
162  __ESIMD_NS::simd<T0, 1> Result =
163  esimd::lsr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
164 
165  return Result[0];
166 }
167 
178 template <typename T0, typename T1, int SZ, typename U,
179  class Sat = __ESIMD_NS::saturation_off_tag>
180 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
181  std::is_integral<T1>::value &&
182  std::is_integral<U>::value,
183  __ESIMD_NS::simd<T0, SZ>>
184 asr(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
185  using IntermedTy = __ESIMD_DNS::computation_type_t<T1, T1>;
186  typedef typename std::make_signed<IntermedTy>::type ComputationTy;
187  __ESIMD_NS::simd<ComputationTy, SZ> Src0 = src0;
188  // TODO H/W supports saturation with this op - map to more efficient version.
189  __ESIMD_NS::simd<ComputationTy, SZ> Result = Src0 >> src1;
190 
191  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
192  return Result;
193  else
194  return __ESIMD_NS::saturate<T0>(Result);
195 }
196 
207 template <typename T0, typename T1, typename T2,
208  class Sat = __ESIMD_NS::saturation_off_tag>
209 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
210  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
211  __ESIMD_DNS::is_esimd_scalar<T2>::value &&
212  std::is_integral<T0>::value &&
213  std::is_integral<T1>::value &&
214  std::is_integral<T2>::value,
215  std::remove_const_t<T0>>
216 asr(T1 src0, T2 src1, Sat sat = {}) {
217  __ESIMD_NS::simd<T1, 1> Src0 = src0;
218  __ESIMD_NS::simd<T0, 1> Result =
219  esimd::asr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
220  return Result[0];
221 }
222 
233 template <typename T0, typename T1, int SZ, typename U,
234  class Sat = __ESIMD_NS::saturation_off_tag>
235 __ESIMD_API std::enable_if_t<std::is_integral<T0>::value &&
236  std::is_integral<T1>::value &&
237  std::is_integral<U>::value,
238  __ESIMD_NS::simd<T0, SZ>>
239 shr(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
240  if constexpr (std::is_unsigned<T1>::value) {
241  return esimd::lsr<T0, T1, SZ, U, Sat>(src0, src1, sat);
242  } else {
243  return esimd::asr<T0, T1, SZ, U, Sat>(src0, src1, sat);
244  }
245 }
246 
256 template <typename T0, typename T1, typename T2,
257  class Sat = __ESIMD_NS::saturation_off_tag>
258 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
259  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
260  __ESIMD_DNS::is_esimd_scalar<T2>::value &&
261  std::is_integral<T0>::value &&
262  std::is_integral<T1>::value &&
263  std::is_integral<T2>::value,
264  std::remove_const_t<T0>>
265 shr(T1 src0, T2 src1, Sat sat = {}) {
266  __ESIMD_NS::simd<T1, 1> Src0 = src0;
267  __ESIMD_NS::simd<T0, 1> Result =
268  esimd::shr<T0, T1, 1, T2, Sat>(Src0, src1, sat);
269  return Result[0];
270 }
271 
280 template <typename T0, typename T1, int SZ>
281 __ESIMD_API std::enable_if_t<
282  __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
283  int64_t, uint64_t>() &&
284  __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
285  int64_t, uint64_t>(),
286  __ESIMD_NS::simd<T0, SZ>>
287 rol(__ESIMD_NS::simd<T1, SZ> src0, __ESIMD_NS::simd<T1, SZ> src1) {
288  return __esimd_rol<T0, T1, SZ>(src0.data(), src1.data());
289 }
290 
299 template <typename T0, typename T1, int SZ, typename U>
300 __ESIMD_API std::enable_if_t<
301  __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
302  int64_t, uint64_t>() &&
303  __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
304  int64_t, uint64_t>() &&
305  __ESIMD_NS::detail::is_type<U, int16_t, uint16_t, int32_t, uint32_t,
306  int64_t, uint64_t>(),
307  __ESIMD_NS::simd<T0, SZ>>
308 rol(__ESIMD_NS::simd<T1, SZ> src0, U src1) {
309  __ESIMD_NS::simd<T1, SZ> Src1 = src1;
310  return esimd::rol<T0>(src0, Src1);
311 }
312 
320 template <typename T0, typename T1, typename T2>
321 __ESIMD_API std::enable_if_t<
322  __ESIMD_DNS::is_esimd_scalar<T0>::value &&
323  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
324  __ESIMD_DNS::is_esimd_scalar<T2>::value &&
325  __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
326  int64_t, uint64_t>() &&
327  __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
328  int64_t, uint64_t>() &&
329  __ESIMD_NS::detail::is_type<T2, int16_t, uint16_t, int32_t, uint32_t,
330  int64_t, uint64_t>(),
331  std::remove_const_t<T0>>
332 rol(T1 src0, T2 src1) {
333  __ESIMD_NS::simd<T1, 1> Src0 = src0;
334  __ESIMD_NS::simd<T0, 1> Result = esimd::rol<T0, T1, 1, T2>(Src0, src1);
335  return Result[0];
336 }
337 
346 template <typename T0, typename T1, int SZ>
347 __ESIMD_API std::enable_if_t<
348  __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
349  int64_t, uint64_t>() &&
350  __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
351  int64_t, uint64_t>(),
352  __ESIMD_NS::simd<T0, SZ>>
353 ror(__ESIMD_NS::simd<T1, SZ> src0, __ESIMD_NS::simd<T1, SZ> src1) {
354  return __esimd_ror<T0, T1, SZ>(src0.data(), src1.data());
355 }
356 
365 template <typename T0, typename T1, int SZ, typename U>
366 __ESIMD_API std::enable_if_t<
367  __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
368  int64_t, uint64_t>() &&
369  __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
370  int64_t, uint64_t>() &&
371  __ESIMD_NS::detail::is_type<U, int16_t, uint16_t, int32_t, uint32_t,
372  int64_t, uint64_t>(),
373  __ESIMD_NS::simd<T0, SZ>>
374 ror(__ESIMD_NS::simd<T1, SZ> src0, U src1) {
375  __ESIMD_NS::simd<T1, SZ> Src1 = src1;
376  return esimd::ror<T0>(src0, Src1);
377 }
378 
386 template <typename T0, typename T1, typename T2>
387 __ESIMD_API std::enable_if_t<
388  __ESIMD_DNS::is_esimd_scalar<T0>::value &&
389  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
390  __ESIMD_DNS::is_esimd_scalar<T2>::value &&
391  __ESIMD_NS::detail::is_type<T0, int16_t, uint16_t, int32_t, uint32_t,
392  int64_t, uint64_t>() &&
393  __ESIMD_NS::detail::is_type<T1, int16_t, uint16_t, int32_t, uint32_t,
394  int64_t, uint64_t>() &&
395  __ESIMD_NS::detail::is_type<T2, int16_t, uint16_t, int32_t, uint32_t,
396  int64_t, uint64_t>(),
397  std::remove_const_t<T0>>
398 ror(T1 src0, T2 src1) {
399  __ESIMD_NS::simd<T1, 1> Src0 = src0;
400  __ESIMD_NS::simd<T0, 1> Result = esimd::ror<T0, T1, 1, T2>(Src0, src1);
401  return Result[0];
402 }
403 
405 
408 
414 template <typename T, typename T0, typename T1, int N>
415 __ESIMD_API __ESIMD_NS::simd<T, N> imul_impl(__ESIMD_NS::simd<T, N> &rmd,
416  __ESIMD_NS::simd<T0, N> src0,
417  __ESIMD_NS::simd<T1, N> src1) {
418  static_assert(__ESIMD_DNS::is_dword_type<T>::value &&
419  __ESIMD_DNS::is_dword_type<T0>::value &&
420  __ESIMD_DNS::is_dword_type<T1>::value,
421  "expected 32-bit integer vector operands.");
422  using Comp32T = __ESIMD_DNS::computation_type_t<T0, T1>;
423  auto Src0 = src0.template bit_cast_view<Comp32T>();
424  auto Src1 = src1.template bit_cast_view<Comp32T>();
425 
426  // Compute the result using 64-bit multiplication operation.
427  using Comp64T =
428  std::conditional_t<std::is_signed_v<Comp32T>, int64_t, uint64_t>;
429  __ESIMD_NS::simd<Comp64T, N> Product64 = Src0;
430  Product64 *= Src1;
431 
432  // Split the 32-bit high and low parts to return them from this function.
433  auto Product32 = Product64.template bit_cast_view<T>();
434  if constexpr (N == 1) {
435  rmd = Product32[0];
436  return Product32[1];
437  } else {
438  rmd = Product32.template select<N, 2>(0);
439  return Product32.template select<N, 2>(1);
440  }
441 }
442 
447 template <typename T, typename T0, typename T1, int N>
448 __ESIMD_API __ESIMD_NS::simd<T, N> imul(__ESIMD_NS::simd<T, N> &rmd,
449  __ESIMD_NS::simd<T0, N> src0,
450  __ESIMD_NS::simd<T1, N> src1) {
451  return imul_impl<T, T0, T1, N>(rmd, src0, src1);
452 }
453 
459 template <typename T, typename T0, typename T1, int N>
460 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_dword_type<T1>::value,
461  __ESIMD_NS::simd<T, N>>
462 imul(__ESIMD_NS::simd<T, N> &rmd, __ESIMD_NS::simd<T0, N> src0, T1 src1) {
463  __ESIMD_NS::simd<T1, N> Src1V = src1;
464  return esimd::imul_impl<T, T0, T1, N>(rmd, src0, Src1V);
465 }
466 
472 template <typename T, typename T0, typename T1, int N>
473 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_dword_type<T0>::value,
474  __ESIMD_NS::simd<T, N>>
475 imul(__ESIMD_NS::simd<T, N> &rmd, T0 src0, __ESIMD_NS::simd<T1, N> src1) {
476  __ESIMD_NS::simd<T0, N> Src0V = src0;
477  return esimd::imul_impl<T, T0, T1, N>(rmd, Src0V, src1);
478 }
479 
484 template <typename T, typename T0, typename T1>
485 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_dword_type<T>::value &&
486  __ESIMD_DNS::is_dword_type<T0>::value &&
487  __ESIMD_DNS::is_dword_type<T1>::value,
488  T>
489 imul(T &rmd, T0 src0, T1 src1) {
490  __ESIMD_NS::simd<T, 1> RmdV = rmd;
491  __ESIMD_NS::simd<T0, 1> Src0V = src0;
492  __ESIMD_NS::simd<T1, 1> Src1V = src1;
493  __ESIMD_NS::simd<T, 1> Res =
494  esimd::imul_impl<T, T0, T1, 1>(RmdV, Src0V, Src1V);
495  rmd = RmdV[0];
496  return Res[0];
497 }
498 
499 template <int N>
500 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
501 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
502  __ESIMD_NS::simd<uint32_t, N> src1) {
503  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
504  __ESIMD_DNS::vector_type_t<uint32_t, N>>
505  Result = __esimd_addc<uint32_t, N>(src0.data(), src1.data());
506 
507  carry = Result.first;
508  return Result.second;
509 }
510 
511 template <int N>
512 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
513 addc(__ESIMD_NS::simd<uint32_t, N> &carry, __ESIMD_NS::simd<uint32_t, N> src0,
514  uint32_t src1) {
515  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
516  return addc(carry, src0, Src1V);
517 }
518 
519 template <int N>
520 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
521 addc(__ESIMD_NS::simd<uint32_t, N> &carry, uint32_t src0,
522  __ESIMD_NS::simd<uint32_t, N> src1) {
523  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
524  return addc(carry, Src0V, src1);
525 }
526 
527 __ESIMD_API uint32_t addc(uint32_t &carry, uint32_t src0, uint32_t src1) {
528  __ESIMD_NS::simd<uint32_t, 1> CarryV = carry;
529  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
530  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
531  __ESIMD_NS::simd<uint32_t, 1> Res = addc(CarryV, Src0V, Src1V);
532  carry = CarryV[0];
533  return Res[0];
534 }
535 
536 template <int N>
537 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
538 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
539  __ESIMD_NS::simd<uint32_t, N> src1) {
540  std::pair<__ESIMD_DNS::vector_type_t<uint32_t, N>,
541  __ESIMD_DNS::vector_type_t<uint32_t, N>>
542  Result = __esimd_subb<uint32_t, N>(src0.data(), src1.data());
543 
544  borrow = Result.first;
545  return Result.second;
546 }
547 
548 template <int N>
549 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
550 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, __ESIMD_NS::simd<uint32_t, N> src0,
551  uint32_t src1) {
552  __ESIMD_NS::simd<uint32_t, N> Src1V = src1;
553  return subb(borrow, src0, Src1V);
554 }
555 
556 template <int N>
557 __ESIMD_API __ESIMD_NS::simd<uint32_t, N>
558 subb(__ESIMD_NS::simd<uint32_t, N> &borrow, uint32_t src0,
559  __ESIMD_NS::simd<uint32_t, N> src1) {
560  __ESIMD_NS::simd<uint32_t, N> Src0V = src0;
561  return subb(borrow, Src0V, src1);
562 }
563 
564 __ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1) {
565  __ESIMD_NS::simd<uint32_t, 1> BorrowV = borrow;
566  __ESIMD_NS::simd<uint32_t, 1> Src0V = src0;
567  __ESIMD_NS::simd<uint32_t, 1> Src1V = src1;
568  __ESIMD_NS::simd<uint32_t, 1> Res = subb(BorrowV, Src0V, Src1V);
569  borrow = BorrowV[0];
570  return Res[0];
571 }
572 
580 template <typename T, int SZ, typename U>
581 __ESIMD_API
582  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
583  __ESIMD_NS::simd<T, SZ>>
584  quot(__ESIMD_NS::simd<T, SZ> src0, U src1) {
585  return src0 / src1;
586 }
587 
594 template <typename T0, typename T1>
595 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
596  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
597  std::is_integral<T0>::value &&
598  std::is_integral<T1>::value,
599  std::remove_const_t<T0>>
600 quot(T0 src0, T1 src1) {
601  return src0 / src1;
602 }
603 
611 template <typename T, int SZ, typename U>
612 __ESIMD_API
613  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
614  __ESIMD_NS::simd<T, SZ>>
615  mod(__ESIMD_NS::simd<T, SZ> src0, U src1) {
616  return src0 % src1;
617 }
618 
625 template <typename T0, typename T1>
626 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
627  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
628  std::is_integral<T0>::value &&
629  std::is_integral<T1>::value,
630  std::remove_const_t<T0>>
631 mod(T0 src0, T1 src1) {
632  return src0 % src1;
633 }
634 
644 template <typename T, int SZ, typename U>
645 __ESIMD_API
646  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
647  __ESIMD_NS::simd<T, SZ>>
648  div(__ESIMD_NS::simd<T, SZ> &remainder, __ESIMD_NS::simd<T, SZ> src0,
649  U src1) {
650  remainder = src0 % src1;
651  return src0 / src1;
652 }
653 
663 template <typename T, int SZ, typename U>
664 __ESIMD_API
665  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value &&
666  __ESIMD_DNS::is_esimd_scalar<U>::value,
667  __ESIMD_NS::simd<T, SZ>>
668  div(__ESIMD_NS::simd<T, SZ> &remainder, U src0,
669  __ESIMD_NS::simd<T, SZ> src1) {
670  remainder = src0 % src1;
671  return src0 / src1;
672 }
673 
683 template <typename RT, typename T0, typename T1>
684 ESIMD_NODEBUG
685  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<RT>::value &&
686  __ESIMD_DNS::is_esimd_scalar<T0>::value &&
687  __ESIMD_DNS::is_esimd_scalar<T1>::value,
688  std::remove_const_t<RT>>
689  div(__ESIMD_NS::simd<std::remove_const_t<RT>, 1> &remainder, T0 src0,
690  T1 src1) {
691  remainder[0] = src0 % src1;
692  return src0 / src1;
693 }
694 
695 // Dot product builtins
696 #if defined(ESIMD_GEN7_5) || defined(ESIMD_GEN8) || defined(ESIMD_GEN8_5) || \
697  defined(ESIMD_GEN9) || defined(ESIMD_GEN9_5)
698 
709 template <typename T0, typename T1, int SZ, typename U,
710  class Sat = __ESIMD_NS::saturation_off_tag>
711 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dp2(__ESIMD_NS::simd<T1, SZ> src0, U src1,
712  Sat sat = {}) {
713  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
714  __ESIMD_NS::simd<float, SZ> Src0 = src0;
715  __ESIMD_NS::simd<float, SZ> Src1 = src1;
716  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp2(Src0.data(), Src1.data());
717  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
718  return Result;
719  else
720  return __ESIMD_NS::saturate<T0>(Result);
721 }
722 
733 template <typename T0, typename T1, int SZ, typename U,
734  class Sat = __ESIMD_NS::saturation_off_tag>
735 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dp3(__ESIMD_NS::simd<T1, SZ> src0, U src1,
736  Sat sat = {}) {
737  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
738  __ESIMD_NS::simd<float, SZ> Src0 = src0;
739  __ESIMD_NS::simd<float, SZ> Src1 = src1;
740  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp3(Src0.data(), Src1.data());
741  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
742  return Result;
743  else
744  return __ESIMD_NS::saturate<T0>(Result);
745 }
746 
757 template <typename T0, typename T1, int SZ, typename U,
758  class Sat = __ESIMD_NS::saturation_off_tag>
759 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dp4(__ESIMD_NS::simd<T1, SZ> src0, U src1,
760  Sat sat = {}) {
761  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
762  __ESIMD_NS::simd<float, SZ> Src0 = src0;
763  __ESIMD_NS::simd<float, SZ> Src1 = src1;
764  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp4(Src0.data(), Src1.data());
765  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
766  return Result;
767  else
768  return __ESIMD_NS::saturate<T0>(Result);
769 }
770 
781 template <typename T0, typename T1, typename U, int SZ,
782  class Sat = __ESIMD_NS::saturation_off_tag>
783 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dph(__ESIMD_NS::simd<T1, SZ> src0, U src1,
784  Sat sat = {}) {
785  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
786  __ESIMD_NS::simd<float, SZ> Src0 = src0;
787  __ESIMD_NS::simd<float, SZ> Src1 = src1;
788  __ESIMD_NS::simd<float, SZ> Result = __esimd_dph(Src0.data(), Src1.data());
789  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
790  return Result;
791  else
792  return __ESIMD_NS::saturate<T0>(Result);
793 }
794 
806 template <typename RT, typename T1, typename T2, int SZ,
807  class Sat = __ESIMD_NS::saturation_off_tag>
808 __ESIMD_API __ESIMD_NS::simd<RT, SZ> line(__ESIMD_NS::simd<T1, 4> src0,
809  __ESIMD_NS::simd<T2, SZ> src1,
810  Sat sat = {}) {
811  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
812 
813  __ESIMD_NS::simd<float, 4> Src0 = src0;
814  __ESIMD_NS::simd<float, SZ> Src1 = src1;
815  __ESIMD_NS::simd<float, SZ> Result = __esimd_line(Src0.data(), Src1.data());
816 
817  __ESIMD_NS::simd<RT, SZ> Result;
818  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
819  return Result;
820  else
821  return __ESIMD_NS::saturate<RT>(Result);
822 }
823 
835 template <typename RT, typename T, int SZ,
836  class Sat = __ESIMD_NS::saturation_off_tag>
837 __ESIMD_API __ESIMD_NS::simd<RT, SZ>
838 line(float P, float Q, __ESIMD_NS::simd<T, SZ> src1, Sat sat = {}) {
839  __ESIMD_NS::simd<float, 4> Src0 = P;
840  Src0(3) = Q;
841  return esimd::line<RT>(Src0, src1, sat);
842 }
843 
844 #else
845 // The old implementation is to generate vISA IRs for dp2/dp3/dp4/dph/line.
846 // Now We change to use direct mul/add, and hope to generate mad instructions
847 // at the end, to still get the performance as good as HW solution.
848 // We rely on "pragma unroll" to get better code.
849 // The only input and return types for these APIs are floats.
850 // In order to be able to use the old emu code, we keep the template argument
851 // for the type, although the type "T" can only be float.
852 // We use std::enable_if to force the float type only.
853 // If the gen is not specified we warn the programmer that they are potentially
854 // using a less efficient implementation if not on GEN10 or above.
855 
866 template <typename T0, typename T1, int SZ, typename U,
867  class Sat = __ESIMD_NS::saturation_off_tag>
868 ESIMD_NODEBUG ESIMD_INLINE
869  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
870  std::is_floating_point<T1>::value &&
871  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
872  std::is_floating_point<U>::value,
873  __ESIMD_NS::simd<T0, SZ>>
874  dp2(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
875  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
876 
877  __ESIMD_NS::simd<float, SZ> Src1 = src1;
878  __ESIMD_NS::simd<float, SZ> Result;
879 #pragma unroll
880  for (int i = 0; i < SZ; i += 4) {
881  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1];
882  }
883  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
884  return Result;
885  else
886  return __ESIMD_NS::saturate<T1>(Result);
887 }
888 
899 template <typename T0, typename T1, int SZ, typename U,
900  class Sat = __ESIMD_NS::saturation_off_tag>
901 ESIMD_NODEBUG ESIMD_INLINE
902  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
903  std::is_floating_point<T1>::value &&
904  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
905  std::is_floating_point<U>::value,
906  __ESIMD_NS::simd<T0, SZ>>
907  dp3(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
908  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
909 
910  __ESIMD_NS::simd<float, SZ> Src1 = src1;
911  __ESIMD_NS::simd<float, SZ> Result;
912 #pragma unroll
913  for (int i = 0; i < SZ; i += 4) {
914  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
915  src0[i + 2] * Src1[i + 2];
916  }
917  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
918  return Result;
919  else
920  return __ESIMD_NS::saturate<T1>(Result);
921 }
922 
933 template <typename T0, typename T1, int SZ, typename U,
934  class Sat = __ESIMD_NS::saturation_off_tag>
935 ESIMD_NODEBUG ESIMD_INLINE
936  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
937  std::is_floating_point<T1>::value &&
938  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
939  std::is_floating_point<U>::value,
940  __ESIMD_NS::simd<T0, SZ>>
941  dp4(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
942  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
943 
944  __ESIMD_NS::simd<T1, SZ> Src1 = src1;
945  __ESIMD_NS::simd<float, SZ> Result;
946 #pragma unroll
947  for (int i = 0; i < SZ; i += 4) {
948  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
949  src0[i + 2] * Src1[i + 2] +
950  src0[i + 3] * Src1[i + 3];
951  }
952  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
953  return Result;
954  else
955  return __ESIMD_NS::saturate<T1>(Result);
956 }
957 
968 template <typename T, typename U, int SZ,
969  class Sat = __ESIMD_NS::saturation_off_tag>
970 ESIMD_NODEBUG ESIMD_INLINE
971  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
972  std::is_floating_point<T>::value &&
973  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
974  std::is_floating_point<U>::value,
975  __ESIMD_NS::simd<T, SZ>>
976  dph(__ESIMD_NS::simd<T, SZ> src0, U src1, Sat sat = {}) {
977  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
978 
979  __ESIMD_NS::simd<float, SZ> Src1 = src1;
980  __ESIMD_NS::simd<float, SZ> Result;
981 #pragma unroll
982  for (int i = 0; i < SZ; i += 4) {
983  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
984  src0[i + 2] * Src1[i + 2] + 1.0 * Src1[i + 3];
985  }
986  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
987  return Result;
988  else
989  return __ESIMD_NS::saturate<T>(Result);
990 }
991 
1002 template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
1003 ESIMD_NODEBUG
1004  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
1005  std::is_floating_point<T>::value,
1006  __ESIMD_NS::simd<T, SZ>>
1007  line(__ESIMD_NS::simd<T, 4> src0, __ESIMD_NS::simd<T, SZ> src1,
1008  Sat sat = {}) {
1009  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
1010 
1011  __ESIMD_NS::simd<T, SZ> Src1 = src1;
1012  __ESIMD_NS::simd<T, SZ> Result;
1013 #pragma unroll
1014  for (int i = 0; i < SZ; i += 4) {
1015  Result.select<4, 1>(i) = src0[0] * src1[i] + src0[3];
1016  }
1017 
1018  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1019  return Result;
1020  else
1021  return __ESIMD_NS::saturate<T>(Result);
1022 }
1023 
1035 template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
1036 ESIMD_NODEBUG
1037  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
1038  std::is_floating_point<T>::value,
1039  __ESIMD_NS::simd<T, SZ>>
1040  line(float P, float Q, __ESIMD_NS::simd<T, SZ> src1, Sat sat = {}) {
1041  __ESIMD_NS::simd<T, 4> Src0 = P;
1042  Src0(3) = Q;
1043  return esimd::line<T>(Src0, src1, sat);
1044 }
1045 
1046 #endif
1047 
1054 template <typename T, int SZ>
1055 __ESIMD_API __ESIMD_NS::simd<T, SZ> frc(__ESIMD_NS::simd<T, SZ> src0) {
1056  __ESIMD_NS::simd<float, SZ> Src0 = src0;
1057  return __esimd_frc(Src0.data());
1058 }
1059 
1065 template <typename T> __ESIMD_API T frc(T src0) {
1066  __ESIMD_NS::simd<T, 1> Src0 = src0;
1067  __ESIMD_NS::simd<T, 1> Result = esimd::frc<T>(Src0);
1068  return Result[0];
1069 }
1070 
1071 // lzd
1072 template <typename RT, typename T0, int SZ,
1073  class Sat = __ESIMD_NS::saturation_off_tag>
1074 __ESIMD_API __ESIMD_NS::simd<RT, SZ> lzd(__ESIMD_NS::simd<T0, SZ> src0,
1075  Sat sat = {}) {
1076  // Saturation parameter ignored
1077  __ESIMD_NS::simd<__ESIMD_NS::uint, SZ> Src0 = src0;
1078  return __esimd_lzd<__ESIMD_NS::uint, SZ>(Src0.data());
1079 }
1080 
1081 template <typename RT, typename T0, class Sat = __ESIMD_NS::saturation_off_tag>
1082 ESIMD_NODEBUG
1083  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<RT>::value &&
1084  __ESIMD_DNS::is_esimd_scalar<T0>::value,
1085  std::remove_const_t<RT>>
1086  lzd(T0 src0, Sat sat = {}) {
1087  __ESIMD_NS::simd<T0, 1> Src0 = src0;
1088  __ESIMD_NS::simd<RT, 1> Result = esimd::lzd<RT>(Src0);
1089  return Result[0];
1090 }
1091 
1092 // lrp
1093 #if defined(ESIMD_GEN7_5) || defined(ESIMD_GEN8) || defined(ESIMD_GEN8_5) || \
1094  defined(ESIMD_GEN9) || defined(ESIMD_GEN9_5)
1095 
1096 template <int SZ, typename U, typename V,
1097  class Sat = __ESIMD_NS::saturation_off_tag>
1098 __ESIMD_API __ESIMD_NS::simd<float, SZ> lrp(__ESIMD_NS::simd<float, SZ> src0,
1099  U src1, V src2, Sat sat = {}) {
1100  static_assert(SZ >= 4 && (SZ & 0x3) == 0,
1101  "vector size must be a multiple of 4");
1102  __ESIMD_NS::simd<float, SZ> Src1 = src1;
1103  __ESIMD_NS::simd<float, SZ> Src2 = src2;
1104  __ESIMD_NS::simd<float, SZ> Result =
1105  __esimd_lrp<SZ>(src0.data(), Src1.data(), Src2.data());
1106 
1107  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1108  return Result;
1109  else
1110  return __ESIMD_NS::saturate<float>(Result);
1111 }
1112 
1113 #else
1114 
1115 // The old implementation is to generate vISA IRs for lrp.
1116 // Now We change to use direct mul/add, and hope to generate mad instructions
1117 // at the end, to still get the performance as good as HW solution.
1118 // The only input and return types for these APIs are floats.
1119 // In order to be able to use the old emu code, we keep the template argument
1120 // for the type, although the type "T" can only be float.
1121 // We use std::enable_if to force the float type only.
1122 // If the gen is not specified we warn the programmer that they are potentially
1123 // using less efficient implementation.
1124 template <typename T, int SZ, typename U, typename V,
1125  class Sat = __ESIMD_NS::saturation_off_tag>
1126 ESIMD_NODEBUG ESIMD_INLINE
1127  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
1128  std::is_floating_point<T>::value &&
1129  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
1130  std::is_floating_point<U>::value,
1131  __ESIMD_NS::simd<T, SZ>>
1132  lrp(__ESIMD_NS::simd<T, SZ> src0, U src1, V src2, Sat sat = {}) {
1133 
1134  __ESIMD_NS::simd<float, SZ> Src1 = src1;
1135  __ESIMD_NS::simd<float, SZ> Src2 = src2;
1136  __ESIMD_NS::simd<float, SZ> Result;
1137  Result = Src1 * src0 + Src2 * (1.0f - src0);
1138  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1139  return Result;
1140  else
1141  return __ESIMD_NS::saturate<T>(Result);
1142 }
1143 #endif
1144 
1146 
1149 
1151 template <typename T0, typename T1, int SZ>
1152 __ESIMD_API __ESIMD_NS::simd<T0, SZ> bf_reverse(__ESIMD_NS::simd<T1, SZ> src0) {
1153  __ESIMD_NS::simd<unsigned, SZ> Src0 = src0;
1154  return __esimd_bfrev<unsigned>(Src0.data());
1155 }
1156 
1158 template <typename T0, typename T1>
1159 ESIMD_NODEBUG
1160  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1161  __ESIMD_DNS::is_esimd_scalar<T1>::value,
1162  std::remove_const_t<T0>>
1163  bf_reverse(T1 src0) {
1164  __ESIMD_NS::simd<T1, 1> Src0 = src0;
1165  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_reverse<T0>(Src0);
1166  return Result[0];
1167 }
1168 
1170 template <typename T0, typename T1, int SZ, typename U, typename V, typename W>
1171 ESIMD_NODEBUG ESIMD_INLINE
1172  std::enable_if_t<std::is_integral<T1>::value, __ESIMD_NS::simd<T0, SZ>>
1173  bf_insert(U src0, V src1, W src2, __ESIMD_NS::simd<T1, SZ> src3) {
1174  typedef typename __ESIMD_DNS::dword_type<T1> DT1;
1175  static_assert(std::is_integral<DT1>::value && sizeof(DT1) == sizeof(int),
1176  "operand conversion failed");
1177  __ESIMD_NS::simd<DT1, SZ> Src0 = src0;
1178  __ESIMD_NS::simd<DT1, SZ> Src1 = src1;
1179  __ESIMD_NS::simd<DT1, SZ> Src2 = src2;
1180  __ESIMD_NS::simd<DT1, SZ> Src3 = src3;
1181 
1182  return __esimd_bfi<DT1>(Src0.data(), Src1.data(), Src2.data(), Src3.data());
1183 }
1184 
1186 template <typename T0, typename T1, typename T2, typename T3, typename T4>
1187 ESIMD_NODEBUG
1188  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1189  __ESIMD_DNS::is_esimd_scalar<T4>::value,
1190  std::remove_const_t<T0>>
1191  bf_insert(T1 src0, T2 src1, T3 src2, T4 src3) {
1192  __ESIMD_NS::simd<T4, 1> Src3 = src3;
1193  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_insert<T0>(src0, src1, src2, Src3);
1194  return Result[0];
1195 }
1196 
1198 template <typename T0, typename T1, int SZ, typename U, typename V>
1199 ESIMD_NODEBUG ESIMD_INLINE
1200  std::enable_if_t<std::is_integral<T1>::value, __ESIMD_NS::simd<T0, SZ>>
1201  bf_extract(U src0, V src1, __ESIMD_NS::simd<T1, SZ> src2) {
1202  typedef typename __ESIMD_DNS::dword_type<T1> DT1;
1203  static_assert(std::is_integral<DT1>::value && sizeof(DT1) == sizeof(int),
1204  "operand conversion failed");
1205  __ESIMD_NS::simd<DT1, SZ> Src0 = src0;
1206  __ESIMD_NS::simd<DT1, SZ> Src1 = src1;
1207  __ESIMD_NS::simd<DT1, SZ> Src2 = src2;
1208 
1209  return __esimd_sbfe<DT1>(Src0.data(), Src1.data(), Src2.data());
1210 }
1211 
1213 template <typename T0, typename T1, typename T2, typename T3>
1214 ESIMD_NODEBUG
1215  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1216  __ESIMD_DNS::is_esimd_scalar<T3>::value,
1217  std::remove_const_t<T0>>
1218  bf_extract(T1 src0, T2 src1, T3 src2) {
1219  __ESIMD_NS::simd<T3, 1> Src2 = src2;
1220  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_extract<T0>(src0, src1, Src2);
1221  return Result[0];
1222 }
1223 
1225 
1228 
1229 // sincos
1230 template <int SZ, typename U, class Sat = __ESIMD_NS::saturation_off_tag>
1231 __ESIMD_API __ESIMD_NS::simd<float, SZ>
1232 sincos(__ESIMD_NS::simd<float, SZ> &dstcos, U src0, Sat sat = {}) {
1233  dstcos = __ESIMD_NS::cos(src0, sat);
1234  return __ESIMD_NS::sin(src0, sat);
1235 }
1236 
1237 // atan
1238 
1240 namespace detail {
1241 constexpr double __ESIMD_CONST_PI = 3.1415926535897932384626433832795;
1242 } // namespace detail
1244 
1245 template <typename T, int SZ>
1246 ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<T, SZ>
1247 atan(__ESIMD_NS::simd<T, SZ> src0) {
1248  static_assert(std::is_floating_point<T>::value,
1249  "Floating point argument type is expected.");
1250  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
1251 
1252  __ESIMD_NS::simd<T, SZ> OneP((T)1.0);
1253  __ESIMD_NS::simd<T, SZ> OneN((T)-1.0);
1254  __ESIMD_NS::simd<T, SZ> sign;
1255  __ESIMD_NS::simd_mask<SZ> Gt1 = Src0 > T(1.0);
1256 
1257  sign.merge(OneN, OneP, src0 < 0);
1258 
1259  Src0.merge(__ESIMD_NS::inv(Src0), Gt1);
1260 
1261  __ESIMD_NS::simd<T, SZ> Src0P2 = Src0 * Src0;
1262  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
1263 
1264  __ESIMD_NS::simd<T, SZ> Result =
1265  (Src0P4 * T(0.185696) + ((Src0 * T(0.787997) + T(0.63693)) * Src0P2) +
1266  Src0) /
1267  (((((Src0 * -T(0.000121387) + T(0.00202308)) * Src0P2) +
1268  (Src0 * -T(0.0149145)) + T(0.182569)) *
1269  Src0P4) +
1270  ((Src0 * T(0.395889) + T(1.12158)) * Src0P2) + (Src0 * T(0.636918)) +
1271  T(1.0));
1272 
1273  Result.merge(Result - T(detail::__ESIMD_CONST_PI) / T(2.0), Gt1);
1274 
1275  return __ESIMD_NS::abs(Result) * sign;
1276 }
1277 
1278 template <typename T> __ESIMD_API T atan(T src0) {
1279  static_assert(std::is_floating_point<T>::value,
1280  "Floating point argument type is expected.");
1281  __ESIMD_NS::simd<T, 1> Src0 = src0;
1282  __ESIMD_NS::simd<T, 1> Result = esimd::atan(Src0);
1283  return Result[0];
1284 }
1285 
1286 // acos
1287 
1288 template <typename T, int SZ>
1289 ESIMD_NODEBUG ESIMD_INLINE
1290  std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
1291  acos(__ESIMD_NS::simd<T, SZ> src0) {
1292  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
1293 
1294  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
1295  __ESIMD_NS::simd_mask<SZ> TooBig = Src0 >= T(0.999998);
1296 
1297  // Replace oversized values to ensure no possibility of sqrt of
1298  // a negative value later
1299  Src0.merge(T(0.0), TooBig);
1300 
1301  __ESIMD_NS::simd<T, SZ> Src01m = T(1.0) - Src0;
1302 
1303  __ESIMD_NS::simd<T, SZ> Src0P2 = Src01m * Src01m;
1304  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
1305 
1306  __ESIMD_NS::simd<T, SZ> Result =
1307  (((Src01m * T(0.015098965761299077) - T(0.005516443930088506)) * Src0P4) +
1308  ((Src01m * T(0.047654245891495528) + T(0.163910606547823220)) * Src0P2) +
1309  Src01m * T(2.000291665285952400) - T(0.000007239283986332)) *
1310  __ESIMD_NS::rsqrt(Src01m * T(2.0));
1311 
1312  Result.merge(T(0.0), TooBig);
1313  Result.merge(T(detail::__ESIMD_CONST_PI) - Result, Neg);
1314  return Result;
1315 }
1316 
1317 template <typename T>
1318 __ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> acos(T src0) {
1319  __ESIMD_NS::simd<T, 1> Src0 = src0;
1320  __ESIMD_NS::simd<T, 1> Result = esimd::acos(Src0);
1321  return Result[0];
1322 }
1323 
1324 // asin
1325 
1326 template <typename T, int SZ>
1327 ESIMD_NODEBUG ESIMD_INLINE
1328  std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
1329  asin(__ESIMD_NS::simd<T, SZ> src0) {
1330  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
1331 
1332  __ESIMD_NS::simd<T, SZ> Result =
1333  T(detail::__ESIMD_CONST_PI / 2.0) - esimd::acos(__ESIMD_NS::abs(src0));
1334 
1335  Result.merge(-Result, Neg);
1336  return Result;
1337 }
1338 
1339 template <typename T>
1340 __ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> asin(T src0) {
1341  __ESIMD_NS::simd<T, 1> Src0 = src0;
1342  __ESIMD_NS::simd<T, 1> Result = esimd::asin(Src0);
1343  return Result[0];
1344 }
1346 
1349 
1350 /* atan2_fast - a fast atan2 implementation */
1351 /* vector input */
1352 template <int N>
1353 __ESIMD_NS::simd<float, N> atan2_fast(__ESIMD_NS::simd<float, N> y,
1354  __ESIMD_NS::simd<float, N> x);
1355 /* scalar input */
1356 template <typename T> float atan2_fast(T y, T x);
1357 
1358 /* atan2 - atan2 implementation */
1359 /* For Vector input */
1360 template <int N>
1361 __ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
1362  __ESIMD_NS::simd<float, N> x);
1363 /* scalar Input */
1364 template <typename T> float atan2(T y, T x);
1365 
1366 /* fmod: */
1367 /* vector input */
1368 template <int N>
1369 __ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
1370  __ESIMD_NS::simd<float, N> x);
1371 /* scalar Input */
1372 template <typename T> float fmod(T y, T x);
1373 
1374 /* sin_emu - EU emulation for sin(x) */
1375 /* For Vector input */
1376 template <int N>
1377 __ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x);
1378 /* scalar Input */
1379 template <typename T> float sin_emu(T x);
1380 
1381 /* cos_emu - EU emulation for cos(x) */
1382 /* For Vector input */
1383 template <int N>
1384 __ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x);
1385 
1386 /* scalar Input */
1387 template <typename T> float cos_emu(T x);
1388 
1389 /* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
1390 /* float input */
1391 float tanh_cody_waite(float x);
1392 /* vector input */
1393 template <int N>
1394 __ESIMD_NS::simd<float, N> tanh_cody_waite(__ESIMD_NS::simd<float, N> x);
1395 /* tanh - opencl like implementation for tanh(x) */
1396 /* float input */
1397 float tanh(float x);
1398 /* vector input */
1399 template <int N> __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x);
1400 
1401 /* ------------------------- Extended Math Routines
1402  * -------------------------------------------------*/
1403 
1404 // For vector input
1405 template <int N>
1406 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1407 atan2_fast(__ESIMD_NS::simd<float, N> y, __ESIMD_NS::simd<float, N> x) {
1408  /* smallest such that 1.0+CONST_DBL_EPSILON != 1.0 */
1409  constexpr float CONST_DBL_EPSILON = 0.00001f;
1410  __ESIMD_NS::simd<float, N> OneP(1.0f);
1411  __ESIMD_NS::simd<float, N> OneN(-1.0f);
1412  __ESIMD_NS::simd<float, N> sign;
1413  __ESIMD_NS::simd<float, N> atan2;
1414  __ESIMD_NS::simd<float, N> r;
1415  __ESIMD_NS::simd_mask<N> mask = x < 0;
1416  __ESIMD_NS::simd<float, N> abs_y = __ESIMD_NS::abs(y) + CONST_DBL_EPSILON;
1417 
1418  r.merge((x + abs_y) / (abs_y - x), (x - abs_y) / (x + abs_y), mask);
1419  atan2.merge(float(detail::__ESIMD_CONST_PI) * 0.75f,
1420  float(detail::__ESIMD_CONST_PI) * 0.25f, mask);
1421  atan2 += (0.1963f * r * r - 0.9817f) * r;
1422 
1423  sign.merge(OneN, OneP, y < 0);
1424 
1425  return atan2 * sign;
1426 }
1427 
1428 // For Scalar Input
1429 template <> ESIMD_INLINE float atan2_fast(float y, float x) {
1430  __ESIMD_NS::simd<float, 1> vy = y;
1431  __ESIMD_NS::simd<float, 1> vx = x;
1432  __ESIMD_NS::simd<float, 1> atan2 = esimd::atan2_fast(vy, vx);
1433  return atan2[0];
1434 }
1435 
1436 // atan2
1437 // For Vector input
1438 template <int N>
1439 ESIMD_INLINE __ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
1440  __ESIMD_NS::simd<float, N> x) {
1441  __ESIMD_NS::simd<float, N> v_distance;
1442  __ESIMD_NS::simd<float, N> atan2;
1443  __ESIMD_NS::simd_mask<N> mask;
1444 
1445  constexpr float CONST_DBL_EPSILON = 0.00001f;
1446 
1447  mask = (x < -CONST_DBL_EPSILON && y < CONST_DBL_EPSILON && y >= 0.f);
1448  atan2.merge(float(detail::__ESIMD_CONST_PI), 0.f, mask);
1449  mask = (x < -CONST_DBL_EPSILON && y > -CONST_DBL_EPSILON && y < 0);
1450  atan2.merge(float(-detail::__ESIMD_CONST_PI), mask);
1451  mask = (x < CONST_DBL_EPSILON && __ESIMD_NS::abs(y) > CONST_DBL_EPSILON);
1452  v_distance = __ESIMD_NS::sqrt(x * x + y * y);
1453  atan2.merge(2.0f * esimd::atan((v_distance - x) / y), mask);
1454 
1455  mask = (x > 0.f);
1456  atan2.merge(2.0f * esimd::atan(y / (v_distance + x)), mask);
1457 
1458  return atan2;
1459 }
1460 
1461 // For Scalar Input
1462 template <> ESIMD_INLINE float atan2(float y, float x) {
1463  __ESIMD_NS::simd<float, 1> vy = y;
1464  __ESIMD_NS::simd<float, 1> vx = x;
1465  __ESIMD_NS::simd<float, 1> atan2 = esimd::atan2(vy, vx);
1466  return atan2[0];
1467 }
1468 
1469 // fmod:
1470 // For Vector input
1471 template <int N>
1472 ESIMD_INLINE __ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
1473  __ESIMD_NS::simd<float, N> x) {
1474  __ESIMD_NS::simd<float, N> abs_x = __ESIMD_NS::abs(x);
1475  __ESIMD_NS::simd<float, N> abs_y = __ESIMD_NS::abs(y);
1476 
1477  auto fmod_sign_mask = (y.template bit_cast_view<int32_t>()) & 0x80000000;
1478 
1479  __ESIMD_NS::simd<float, N> reminder =
1480  abs_y - abs_x * __ESIMD_NS::trunc<float>(abs_y / abs_x);
1481 
1482  abs_x.merge(0.0f, reminder >= 0);
1483  __ESIMD_NS::simd<float, N> fmod = reminder + abs_x;
1484  __ESIMD_NS::simd<float, N> fmod_abs = __ESIMD_NS::abs(fmod);
1485 
1486  auto fmod_bits =
1487  (fmod_abs.template bit_cast_view<int32_t>()) | fmod_sign_mask;
1488  return fmod_bits.template bit_cast_view<float>();
1489 }
1490 
1491 // For Scalar Input
1492 template <> ESIMD_INLINE float fmod(float y, float x) {
1493  return fmod(__ESIMD_NS::simd<float, 1>(y), __ESIMD_NS::simd<float, 1>(x))[0];
1494 }
1495 
1496 // sin_emu - EU emulation for sin(x)
1497 // For Vector input
1498 template <int N>
1499 ESIMD_INLINE __ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x) {
1500  __ESIMD_NS::simd<float, N> x1;
1501  __ESIMD_NS::simd<float, N> x2;
1502  __ESIMD_NS::simd<float, N> t3;
1503 
1504  __ESIMD_NS::simd<float, N> sign;
1505  __ESIMD_NS::simd<float, N> fTrig;
1506  __ESIMD_NS::simd<float, N> TwoPI(float(detail::__ESIMD_CONST_PI) * 2.0f);
1507  __ESIMD_NS::simd<float, N> CmpI((float)detail::__ESIMD_CONST_PI);
1508  __ESIMD_NS::simd<float, N> OneP(1.0f);
1509  __ESIMD_NS::simd<float, N> OneN(-1.0f);
1510 
1511  x = esimd::fmod(x, TwoPI);
1512  x.merge(TwoPI + x, x < 0);
1513 
1514  x1.merge(CmpI - x, x - CmpI, (x <= float(detail::__ESIMD_CONST_PI)));
1515  x1.merge(x, (x <= float(detail::__ESIMD_CONST_PI) * 0.5f));
1516  x1.merge(TwoPI - x, (x > float(detail::__ESIMD_CONST_PI) * 1.5f));
1517 
1518  sign.merge(OneN, OneP, (x > float(detail::__ESIMD_CONST_PI)));
1519 
1520  x2 = x1 * x1;
1521  t3 = x2 * x1 * 0.1666667f;
1522 
1523  fTrig =
1524  x1 + t3 * (OneN + x2 * 0.05f *
1525  (OneP + x2 * 0.0238095f *
1526  (OneN + x2 * 0.0138889f *
1527  (OneP - x2 * 0.0090909f))));
1528  fTrig *= sign;
1529  return fTrig;
1530 }
1531 
1532 // scalar Input
1533 template <> ESIMD_INLINE float sin_emu(float x0) {
1534  return esimd::sin_emu(__ESIMD_NS::simd<float, 1>(x0))[0];
1535 }
1536 
1537 // cos_emu - EU emulation for sin(x)
1538 // For Vector input
1539 template <int N>
1540 ESIMD_INLINE __ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x) {
1541  return esimd::sin_emu(0.5f * float(detail::__ESIMD_CONST_PI) - x);
1542 }
1543 
1544 // scalar Input
1545 template <> ESIMD_INLINE float cos_emu(float x0) {
1546  return esimd::cos_emu(__ESIMD_NS::simd<float, 1>(x0))[0];
1547 }
1548 
1550 namespace detail {
1551 
1552 template <int N>
1553 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1554 tanh_cody_waite_impl(__ESIMD_NS::simd<float, N> x) {
1555  /*
1556  * 0 x_small x_medium x_large
1557  * | x | rational polynomial | 1 - 2/(1 + exp(2*x)) | 1
1558  *
1559  * rational polynomial for single precision = x + x * (g * (p[1] * g + p[0]) /
1560  * (g + q[0]) g = x^2 p0 = -0.82377 28127 E+00 p1 = -0.38310 10665 E-02 q0 =
1561  * 0.24713 19654 E+01 q1 = 1.00000 00000 E+00
1562  *
1563  */
1564 
1565  constexpr float p0 = -0.8237728127E+00f;
1566  constexpr float p1 = -0.3831010665E-02f;
1567  constexpr float q0 = 0.2471319654E+01f;
1568  constexpr float q1 = 1.0000000000E+00f;
1569  constexpr float xsmall = 4.22863966691620432990E-04f;
1570  constexpr float xmedium = 0.54930614433405484570f;
1571  constexpr float xlarge = 8.66433975699931636772f;
1572 
1573  using RT = __ESIMD_NS::simd<float, N>;
1574 
1575  RT absX = __ESIMD_NS::abs(x);
1576  RT g = absX * absX;
1577 
1578  RT sign;
1579  sign.merge(-1.f, 1.f, x < 0.f);
1580 
1581  auto isLarge = absX > xlarge;
1582  auto minor = absX <= xlarge;
1583  auto isGtMed = minor & (absX > xmedium);
1584  auto isGtSmall = (absX > xsmall) & (absX <= xmedium);
1585 
1586  RT res;
1587  res.merge(sign, x, isLarge);
1588  auto temp = __ESIMD_NS::exp(absX * 2.0f) + 1.f;
1589  temp = ((temp - 2.f) / temp) * sign;
1590  res.merge(temp, isGtMed);
1591  res.merge((absX + absX * g * (g * p1 + p0) / (g + q0)) * sign, isGtSmall);
1592 
1593  return res;
1594 }
1595 
1596 template <int N>
1597 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1598 tanh_impl(__ESIMD_NS::simd<float, N> x) {
1599  /*
1600  * 0 x_small x_large
1601  * | x | ( exp(x) - exp(-x) ) / ( exp(x) + exp(-x) ) | 1
1602  *
1603  */
1604 
1605  constexpr float xsmall = 0.000045f; // same as exp(-10.0f)
1606  constexpr float xlarge = 40.f;
1607 
1608  using RT = __ESIMD_NS::simd<float, N>;
1609 
1610  RT absX = __ESIMD_NS::abs(x);
1611 
1612  RT sign;
1613  sign.merge(-1.f, 1.f, x < 0.f);
1614 
1615  auto isLarge = (absX > xlarge);
1616  auto isLessE = (absX <= xlarge);
1617 
1618  RT res;
1619  res.merge(sign, x, isLarge);
1620 
1621  RT exp;
1622  exp = __ESIMD_NS::exp(absX * 2.f);
1623 
1624  res.merge(((exp - 1.f) / (exp + 1.f)) * sign, (absX > xsmall) & isLessE);
1625 
1626  return res;
1627 }
1628 } // namespace detail
1630 
1631 /* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
1632 /* float input */
1633 ESIMD_INLINE float tanh_cody_waite(float x) {
1634  return detail::tanh_cody_waite_impl(__ESIMD_NS::simd<float, 1>(x))[0];
1635 }
1636 /* vector input */
1637 template <int N>
1638 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1639 tanh_cody_waite(__ESIMD_NS::simd<float, N> x) {
1640  return detail::tanh_cody_waite_impl(x);
1641 }
1642 
1643 /* tanh - opencl like implementation for tanh(x) */
1644 /* float input */
1645 ESIMD_INLINE float tanh(float x) {
1646  return esimd::detail::tanh_impl(__ESIMD_NS::simd<float, 1>(x))[0];
1647 }
1648 /* vector input */
1649 template <int N>
1650 ESIMD_INLINE __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x) {
1651  return esimd::detail::tanh_impl(x);
1652 }
1653 
1654 template <typename T, int N>
1655 __ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> v1,
1656  __ESIMD_NS::simd<T, N> v2) {
1657  auto retv = __esimd_dp4<T, N>(v1.data(), v2.data());
1658  return retv;
1659 }
1660 
1662 
1666 
1679 template <argument_type src1_precision, argument_type src2_precision,
1680  typename T, int systolic_depth, int repeat_count, typename T0,
1681  typename T1, typename T2, int N, int N1, int N2,
1682  typename Sat = __ESIMD_NS::saturation_off_tag>
1683 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::native::dpas()")
1684 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1685  __ESIMD_NS::simd<T0, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1686  __ESIMD_NS::simd<T2, N2> src2,
1687  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1688  auto result =
1689  __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T0, T1, T2,
1690  src1_precision, src2_precision>(src0, src1, src2);
1691  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1692  return result;
1693  else
1694  return __ESIMD_NS::saturate<T>(result);
1695 }
1696 
1707 template <argument_type src1_precision, argument_type src2_precision,
1708  int systolic_depth, int repeat_count, typename T, typename T1,
1709  typename T2, int N, int N1, int N2,
1710  typename Sat = __ESIMD_NS::saturation_off_tag>
1711 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
1712 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1713  __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1714  __ESIMD_NS::simd<T2, N2> src2,
1715  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1716  return dpas<src1_precision, src2_precision, T, systolic_depth, repeat_count>(
1717  src0, src1, src2, sat);
1718 }
1719 
1728 template <argument_type src1_precision, argument_type src2_precision,
1729  int systolic_depth, int repeat_count, typename T, typename T1,
1730  typename T2, int N, int N1, int N2,
1731  typename Sat = __ESIMD_NS::saturation_off_tag>
1732 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
1733 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1734  __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
1735  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1736 
1737  __ESIMD_NS::simd<T, N> result =
1738  __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T1, T2,
1739  src1_precision, src2_precision>(src1, src2);
1740 
1741  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1742  return result;
1743  else
1744  return __ESIMD_NS::saturate<T>(result);
1745 }
1746 
1757 template <argument_type src1_precision, argument_type src2_precision,
1758  int systolic_depth, int repeat_count, typename T, typename T1,
1759  typename T2, int N, int N1, int N2,
1760  typename Sat = __ESIMD_NS::saturation_off_tag>
1761 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
1762 __ESIMD_API __ESIMD_NS::simd<T, N> dpasw(
1763  __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1764  __ESIMD_NS::simd<T2, N2> src2,
1765  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1766 
1767  __ESIMD_NS::simd<T, N> result =
1768  __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
1769  src1_precision, src2_precision>(src0, src1, src2);
1770  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1771  return result;
1772  else
1773  return __ESIMD_NS::saturate<T>(result);
1774 }
1775 
1784 template <argument_type src1_precision, argument_type src2_precision,
1785  int systolic_depth, int repeat_count, typename T, typename T1,
1786  typename T2, int N, int N1, int N2,
1787  typename Sat = __ESIMD_NS::saturation_off_tag>
1788 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
1789 __ESIMD_API __ESIMD_NS::simd<T, N> dpasw2(
1790  __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
1791  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1792 
1793  __ESIMD_NS::simd<T, N> result =
1794  __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
1795  src1_precision, src2_precision>(src1, src2);
1796 
1797  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1798  return result;
1799  else
1800  return __ESIMD_NS::saturate<T>(result);
1801 }
1803 
1806 
1811 enum class bfn_t : uint8_t { x = 0xAA, y = 0xCC, z = 0xF0 };
1812 
1813 static constexpr bfn_t operator~(bfn_t x) {
1814  uint8_t val = static_cast<uint8_t>(x);
1815  uint8_t res = ~val;
1816  return static_cast<bfn_t>(res);
1817 }
1818 
1819 static constexpr bfn_t operator|(bfn_t x, bfn_t y) {
1820  uint8_t arg0 = static_cast<uint8_t>(x);
1821  uint8_t arg1 = static_cast<uint8_t>(y);
1822  uint8_t res = arg0 | arg1;
1823  return static_cast<bfn_t>(res);
1824 }
1825 
1826 static constexpr bfn_t operator&(bfn_t x, bfn_t y) {
1827  uint8_t arg0 = static_cast<uint8_t>(x);
1828  uint8_t arg1 = static_cast<uint8_t>(y);
1829  uint8_t res = arg0 & arg1;
1830  return static_cast<bfn_t>(res);
1831 }
1832 
1833 static constexpr bfn_t operator^(bfn_t x, bfn_t y) {
1834  uint8_t arg0 = static_cast<uint8_t>(x);
1835  uint8_t arg1 = static_cast<uint8_t>(y);
1836  uint8_t res = arg0 ^ arg1;
1837  return static_cast<bfn_t>(res);
1838 }
1839 
1848 template <bfn_t FuncControl, typename T, int N>
1849 __ESIMD_API std::enable_if_t<std::is_integral_v<T>, __ESIMD_NS::simd<T, N>>
1850 bfn(__ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T, N> src1,
1851  __ESIMD_NS::simd<T, N> src2) {
1852  if constexpr ((sizeof(T) == 8) || ((sizeof(T) == 1) && (N % 4 == 0)) ||
1853  ((sizeof(T) == 2) && (N % 2 == 0))) {
1854  // Bitcast Nx8-byte vectors to 2xN vectors of 4-byte integers.
1855  // Bitcast Nx1-byte vectors to N/4 vectors of 4-byte integers.
1856  // Bitcast Nx2-byte vectors to N/2 vectors of 4-byte integers.
1857  auto Result = __ESIMD_ENS::bfn<FuncControl>(
1858  src0.template bit_cast_view<int32_t>().read(),
1859  src1.template bit_cast_view<int32_t>().read(),
1860  src2.template bit_cast_view<int32_t>().read());
1861  return Result.template bit_cast_view<T>();
1862  } else if constexpr (sizeof(T) == 2 || sizeof(T) == 4) {
1863  constexpr uint8_t FC = static_cast<uint8_t>(FuncControl);
1864  return __esimd_bfn<FC, T, N>(src0.data(), src1.data(), src2.data());
1865  } else if constexpr (N % 2 == 0) {
1866  // Bitcast Nx1-byte vectors (N is even) to N/2 vectors of 2-byte integers.
1867  auto Result = __ESIMD_ENS::bfn<FuncControl>(
1868  src0.template bit_cast_view<int16_t>().read(),
1869  src1.template bit_cast_view<int16_t>().read(),
1870  src2.template bit_cast_view<int16_t>().read());
1871  return Result.template bit_cast_view<T>();
1872  } else {
1873  // Odd number of 1-byte elements.
1874  __ESIMD_NS::simd<T, N + 1> Src0, Src1, Src2;
1875  Src0.template select<N, 1>() = src0;
1876  Src1.template select<N, 1>() = src1;
1877  Src2.template select<N, 1>() = src2;
1878  auto Result = __ESIMD_ENS::bfn<FuncControl>(Src0, Src1, Src2);
1879  return Result.template select<N, 1>();
1880  }
1881 }
1882 
1890 template <bfn_t FuncControl, typename T>
1891 ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t<
1892  __ESIMD_DNS::is_esimd_scalar<T>::value && std::is_integral_v<T>, T>
1893 bfn(T src0, T src1, T src2) {
1894  __ESIMD_NS::simd<T, 1> Src0 = src0;
1895  __ESIMD_NS::simd<T, 1> Src1 = src1;
1896  __ESIMD_NS::simd<T, 1> Src2 = src2;
1897  __ESIMD_NS::simd<T, 1> Result =
1898  esimd::bfn<FuncControl, T, 1>(Src0, Src1, Src2);
1899  return Result[0];
1900 }
1901 
1903 
1904 } // namespace ext::intel::experimental::esimd
1905 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1906 } // namespace sycl
sycl::_V1::ext::intel::experimental::esimd::bf_insert
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T4 >::value, std::remove_const_t< T0 > > bf_insert(T1 src0, T2 src1, T3 src2, T4 src3)
bf_insert
Definition: math.hpp:1191
math.hpp
math_intrin.hpp
sycl::_V1::ext::intel::experimental::esimd::dpasw2
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpasw2(sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
DPASW2.
Definition: math.hpp:1789
sycl::_V1::ext::intel::experimental::esimd::frc
__ESIMD_API T frc(T src0)
Performs truncate-to-minus-infinity fraction operation of src0.
Definition: math.hpp:1065
sycl::_V1::ext::intel::experimental::esimd::atan
__ESIMD_API T atan(T src0)
Definition: math.hpp:1278
sycl::_V1::ext::intel::experimental::esimd::bfn_t::z
@ z
sycl::_V1::ext::intel::experimental::esimd::rol
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T2, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), std::remove_const_t< T0 > > rol(T1 src0, T2 src1)
Rotate left operation with two scalar inputs.
Definition: math.hpp:332
sycl::_V1::ext::intel::experimental::esimd::dpasw
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpasw(sycl::ext::intel::esimd::simd< T, N > src0, sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
DPASW.
Definition: math.hpp:1762
T
sycl::_V1::ext::intel::experimental::esimd::operator&
static constexpr bfn_t operator&(bfn_t x, bfn_t y)
Definition: math.hpp:1826
sycl::_V1::ext::intel::experimental::esimd::dph
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T, SZ > > dph(sycl::ext::intel::esimd::simd< T, SZ > src0, U src1, Sat sat={})
Dot product on groups of 4 elements.
Definition: math.hpp:976
sycl::_V1::ext::intel::esimd::sin
__ESIMD_API T sin(T src, Sat sat={})
Scalar version.
Definition: math.hpp:401
sycl::_V1::ext::intel::esimd::cos
__ESIMD_API T cos(T src, Sat sat={})
Scalar version.
Definition: math.hpp:405
sycl::_V1::ext::intel::experimental::esimd::cos_emu
ESIMD_INLINE float cos_emu(float x0)
Definition: math.hpp:1545
sycl::_V1::ext::intel::esimd::detail::is_saturation_tag_v
constexpr bool is_saturation_tag_v
Definition: common.hpp:76
sycl::_V1::ext::intel::experimental::esimd::mod
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value, std::remove_const_t< T0 > > mod(T0 src0, T1 src1)
Modulo (scalar version)
Definition: math.hpp:631
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:11
defines_elementary.hpp
sycl::_V1::detail::int64_t
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:34
sycl::_V1::ext::intel::experimental::esimd::operator^
static constexpr bfn_t operator^(bfn_t x, bfn_t y)
Definition: math.hpp:1833
sycl::_V1::ext::intel::experimental::esimd::dpas
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > dpas(sycl::ext::intel::esimd::simd< T1, N1 > src1, sycl::ext::intel::esimd::simd< T2, N2 > src2, std::enable_if_t< __ESIMD_DNS::is_saturation_tag_v< Sat >, Sat > sat={})
DPAS.
Definition: math.hpp:1733
sycl::_V1::ext::intel::experimental::esimd::tanh_cody_waite
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N > tanh_cody_waite(sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1639
sycl::_V1::ext::intel::experimental::esimd::line
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value, sycl::ext::intel::esimd::simd< T, SZ > > line(float P, float Q, sycl::ext::intel::esimd::simd< T, SZ > src1, Sat sat={})
Linear equation.
Definition: math.hpp:1040
sycl::_V1::ext::intel::experimental::esimd::atan2_fast
ESIMD_INLINE float atan2_fast(float y, float x)
Definition: math.hpp:1429
sycl::_V1::ext::intel::math::inv
std::enable_if_t< std::is_same_v< Tp, sycl::half2 >, sycl::half2 > inv(Tp x)
Definition: math.hpp:144
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14
util.hpp
sycl::_V1::ext::intel::experimental::esimd::dp2
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T1 >::value &&std::is_floating_point< T1 >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > > dp2(sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
Dot product on groups of 4 elements.
Definition: math.hpp:874
sycl::_V1::ext::intel::experimental::esimd::addc
__ESIMD_API uint32_t addc(uint32_t &carry, uint32_t src0, uint32_t src1)
Definition: math.hpp:527
sycl::_V1::ext::intel::experimental::esimd::operator|
static constexpr bfn_t operator|(bfn_t x, bfn_t y)
Definition: math.hpp:1819
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:46
sycl::_V1::ext::intel::experimental::esimd::bfn_t::x
@ x
sycl::_V1::ext::intel::experimental::esimd::asin
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, T > asin(T src0)
Definition: math.hpp:1340
sycl::_V1::ext::intel::experimental::esimd::subb
__ESIMD_API uint32_t subb(uint32_t &borrow, uint32_t src0, uint32_t src1)
Definition: math.hpp:564
sycl::_V1::ext::intel::experimental::esimd::lzd
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< RT >::value &&__ESIMD_DNS::is_esimd_scalar< T0 >::value, std::remove_const_t< RT > > lzd(T0 src0, Sat sat={})
Definition: math.hpp:1086
sycl::_V1::ext::intel::experimental::esimd::dp4
sycl::ext::intel::esimd::simd< T, N > dp4(sycl::ext::intel::esimd::simd< T, N > v1, sycl::ext::intel::esimd::simd< T, N > v2)
Definition: math.hpp:1655
sycl::_V1::ext::intel::experimental::esimd::tanh
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N > tanh(sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1650
sycl::_V1::ext::intel::experimental::esimd::sincos
__ESIMD_API sycl::ext::intel::esimd::simd< float, SZ > sincos(sycl::ext::intel::esimd::simd< float, SZ > &dstcos, U src0, Sat sat={})
Definition: math.hpp:1232
types.hpp
sycl::_V1::ext::intel::experimental::esimd::bf_extract
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T3 >::value, std::remove_const_t< T0 > > bf_extract(T1 src0, T2 src1, T3 src2)
bf_extract
Definition: math.hpp:1218
sycl::_V1::ext::intel::experimental::esimd::ror
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&sycl::ext::intel::esimd::detail::is_type< T0, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T1, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >) &&sycl::ext::intel::esimd::detail::is_type< T2, int16_t, uint16_t, int32_t, uint32_t, int64_t, uint64_t >), std::remove_const_t< T0 > > ror(T1 src0, T2 src1)
Rotate right operation with two scalar inputs.
Definition: math.hpp:398
sycl::_V1::ext::intel::experimental::esimd::imul_impl
__ESIMD_API sycl::ext::intel::esimd::simd< T, N > imul_impl(sycl::ext::intel::esimd::simd< T, N > &rmd, sycl::ext::intel::esimd::simd< T0, N > src0, sycl::ext::intel::esimd::simd< T1, N > src1)
Computes the 64-bit result of two 32-bit element vectors src0 and src1 multiplication.
Definition: math.hpp:415
sycl::_V1::ext::intel::esimd::exp
ESIMD_NODEBUG ESIMD_INLINE T exp(T src0, Sat sat={})
Definition: math.hpp:497
common.hpp
sycl::_V1::ext::oneapi::experimental::simd
std::experimental::simd< T, simd_abi::native_fixed_size< T, N > > simd
Definition: invoke_simd.hpp:87
sycl::_V1::ext::intel::experimental::esimd::bfn
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T >::value &&std::is_integral_v< T >, T > bfn(T src0, T src1, T src2)
Performs binary function computation with three scalar operands.
Definition: math.hpp:1893
sycl::_V1::ext::intel::experimental::esimd::bfn_t::y
@ y
sycl::_V1::ext::intel::experimental::esimd::div
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< RT >::value &&__ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value, std::remove_const_t< RT > > div(sycl::ext::intel::esimd::simd< std::remove_const_t< RT >, 1 > &remainder, T0 src0, T1 src1)
Integral division (scalar version).
Definition: math.hpp:689
sycl::_V1::ext::intel::math::sqrt
std::enable_if_t< std::is_same_v< Tp, sycl::half2 >, sycl::half2 > sqrt(Tp x)
Definition: math.hpp:186
simd
Definition: simd.hpp:1027
sycl::_V1::ext::intel::experimental::esimd::fmod
ESIMD_INLINE float fmod(float y, float x)
Definition: math.hpp:1492
sycl::_V1::ext::intel::esimd::abs
__ESIMD_API std::enable_if_t< detail::is_esimd_scalar< T1 >::value, std::remove_const_t< T1 > > abs(T1 src0)
Get absolute value (scalar version).
Definition: math.hpp:166
sycl::_V1::ext::intel::experimental::esimd::atan2
ESIMD_INLINE float atan2(float y, float x)
Definition: math.hpp:1462
std
Definition: accessor.hpp:3914
sycl::_V1::ext::intel::experimental::esimd::quot
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value, std::remove_const_t< T0 > > quot(T0 src0, T1 src1)
Integral quotient (scalar version)
Definition: math.hpp:600
sycl::_V1::ext::intel::experimental::esimd::lrp
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T >::value &&std::is_floating_point< T >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T, SZ > > lrp(sycl::ext::intel::esimd::simd< T, SZ > src0, U src1, V src2, Sat sat={})
Definition: math.hpp:1132
sycl::_V1::ext::intel::experimental::esimd::imul
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_dword_type< T >::value &&__ESIMD_DNS::is_dword_type< T0 >::value &&__ESIMD_DNS::is_dword_type< T1 >::value, T > imul(T &rmd, T0 src0, T1 src1)
Computes the 64-bit multiply result of two scalar 32-bit integer values src0 and src1.
Definition: math.hpp:489
sycl::_V1::image_channel_order::r
@ r
sycl::_V1::ext::intel::experimental::esimd::operator~
static constexpr bfn_t operator~(bfn_t x)
Definition: math.hpp:1813
sycl::_V1::ext::intel::experimental::esimd::asr
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > > asr(T1 src0, T2 src1, Sat sat={})
Arithmetical Shift Right (scalar version)
Definition: math.hpp:216
sycl::_V1::ext::intel::experimental::esimd::shl
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > > shl(T1 src0, T2 src1, Sat sat={})
Shift left operation (scalar version)
Definition: math.hpp:104
sycl::_V1::ext::intel::experimental::esimd::shr
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > > shr(T1 src0, T2 src1, Sat sat={})
Shift right operation (scalar version)
Definition: math.hpp:265
sycl::_V1::ext::intel::experimental::esimd::acos
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, T > acos(T src0)
Definition: math.hpp:1318
sycl::_V1::ext::intel::experimental::esimd::bfn_t
bfn_t
This enum is used to encode all possible logical operations performed on the 3 input operands.
Definition: math.hpp:1811
sycl::_V1::ext::intel::math::rsqrt
std::enable_if_t< std::is_same_v< Tp, sycl::half2 >, sycl::half2 > rsqrt(Tp x)
Definition: math.hpp:207
sycl::_V1::ext::intel::experimental::esimd::dp3
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_fp_or_dword_type< T1 >::value &&std::is_floating_point< T1 >::value &&__ESIMD_DNS::is_fp_or_dword_type< U >::value &&std::is_floating_point< U >::value, sycl::ext::intel::esimd::simd< T0, SZ > > dp3(sycl::ext::intel::esimd::simd< T1, SZ > src0, U src1, Sat sat={})
Dot product on groups of 4 elements.
Definition: math.hpp:907
sycl::_V1::ext::intel::experimental::esimd::bf_reverse
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value, std::remove_const_t< T0 > > bf_reverse(T1 src0)
bf_reverse
Definition: math.hpp:1163
sycl::_V1::ext::intel::experimental::esimd::sin_emu
ESIMD_INLINE float sin_emu(float x0)
Definition: math.hpp:1533
sycl::_V1::ext::intel::experimental::esimd::lsr
__ESIMD_API std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T0 >::value &&__ESIMD_DNS::is_esimd_scalar< T1 >::value &&__ESIMD_DNS::is_esimd_scalar< T2 >::value &&std::is_integral< T0 >::value &&std::is_integral< T1 >::value &&std::is_integral< T2 >::value, std::remove_const_t< T0 > > lsr(T1 src0, T2 src1, Sat sat={})
Logical Shift Right (scalar version)
Definition: math.hpp:160