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 
409 // imul
410 #ifndef ESIMD_HAS_LONG_LONG
411 // use mulh instruction for high half
412 template <typename T0, typename T1, typename U, int SZ>
413 ESIMD_NODEBUG
414  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_dword_type<T0>::value &&
415  __ESIMD_DNS::is_dword_type<T1>::value &&
416  __ESIMD_DNS::is_dword_type<U>::value,
417  __ESIMD_NS::simd<T0, SZ>>
418  imul(__ESIMD_NS::simd<T0, SZ> &rmd, __ESIMD_NS::simd<T1, SZ> src0, U src1) {
419  using ComputationTy = __ESIMD_DNS::computation_type_t<decltype(src0), U>;
420  ComputationTy Src0 = src0;
421  ComputationTy Src1 = src1;
422  rmd = Src0 * Src1;
423  if constexpr (std::is_unsigned<T0>::value)
424  return __esimd_umulh(Src0.data(), Src1.data());
425  else
426  return __esimd_smulh(Src0.data(), Src1.data());
427 }
428 
429 #else
430 // imul bdw+ version: use qw=dw*dw multiply.
431 // We need to special case SZ==1 to avoid "error: when select size is 1, the
432 // stride must also be 1" on the selects.
433 template <typename T0, typename T1, typename U, int SZ>
434 __ESIMD_API
435  std::enable_if_t<__ESIMD_DNS::is_dword_type<T0>::value &&
436  __ESIMD_DNS::is_dword_type<T1>::value &&
437  __ESIMD_DNS::is_dword_type<U>::value && SZ == 1,
438  __ESIMD_NS::simd<T0, SZ>>
439  imul(__ESIMD_NS::simd<T0, SZ> &rmd, __ESIMD_NS::simd<T1, SZ> src0, U src1) {
440  using ComputationTy =
441  __ESIMD_DNS::computation_type_t<decltype(rmd), long long>;
442  ComputationTy Product = convert<long long>(src0);
443  Product *= src1;
444  rmd = Product.bit_cast_view<T0>().select<1, 1>[0];
445  return Product.bit_cast_view<T0>().select<1, 1>[1];
446 }
447 
448 template <typename T0, typename T1, typename U, int SZ>
449 __ESIMD_API
450  std::enable_if_t<__ESIMD_DNS::is_dword_type<T0>::value &&
451  __ESIMD_DNS::is_dword_type<T1>::value &&
452  __ESIMD_DNS::is_dword_type<U>::value && SZ != 1,
453  __ESIMD_NS::simd<T0, SZ>>
454  imul(__ESIMD_NS::simd<T0, SZ> &rmd, __ESIMD_NS::simd<T1, SZ> src0, U src1) {
455  using ComputationTy =
456  __ESIMD_DNS::computation_type_t<decltype(rmd), long long>;
457  ComputationTy Product = convert<long long>(src0);
458  Product *= src1;
459  rmd = Product.bit_cast_view<T0>().select<SZ, 2>(0);
460  return Product.bit_cast_view<T0>().select<SZ, 2>(1);
461 }
462 #endif
463 
464 // TODO: document
465 template <typename T0, typename T1, typename U, int SZ>
466 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<U>::value,
467  __ESIMD_NS::simd<T0, SZ>>
468 imul(__ESIMD_NS::simd<T0, SZ> &rmd, U src0, __ESIMD_NS::simd<T1, SZ> src1) {
469  return esimd::imul(rmd, src1, src0);
470 }
471 
472 // TODO: document
473 template <typename T0, typename T, typename U>
474 ESIMD_NODEBUG
475  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T>::value &&
476  __ESIMD_DNS::is_esimd_scalar<U>::value &&
477  __ESIMD_DNS::is_esimd_scalar<T0>::value,
478  T0>
479  imul(__ESIMD_NS::simd<T0, 1> &rmd, T src0, U src1) {
480  __ESIMD_NS::simd<T, 1> src_0 = src0;
481  __ESIMD_NS::simd<U, 1> src_1 = src1;
482  __ESIMD_NS::simd<T0, 1> res =
483  esimd::imul(rmd, src_0.select_all(), src_1.select_all());
484  return res[0];
485 }
486 
494 template <typename T, int SZ, typename U>
495 __ESIMD_API
496  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
497  __ESIMD_NS::simd<T, SZ>>
498  quot(__ESIMD_NS::simd<T, SZ> src0, U src1) {
499  return src0 / src1;
500 }
501 
508 template <typename T0, typename T1>
509 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
510  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
511  std::is_integral<T0>::value &&
512  std::is_integral<T1>::value,
513  std::remove_const_t<T0>>
514 quot(T0 src0, T1 src1) {
515  return src0 / src1;
516 }
517 
525 template <typename T, int SZ, typename U>
526 __ESIMD_API
527  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
528  __ESIMD_NS::simd<T, SZ>>
529  mod(__ESIMD_NS::simd<T, SZ> src0, U src1) {
530  return src0 % src1;
531 }
532 
539 template <typename T0, typename T1>
540 __ESIMD_API std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
541  __ESIMD_DNS::is_esimd_scalar<T1>::value &&
542  std::is_integral<T0>::value &&
543  std::is_integral<T1>::value,
544  std::remove_const_t<T0>>
545 mod(T0 src0, T1 src1) {
546  return src0 % src1;
547 }
548 
558 template <typename T, int SZ, typename U>
559 __ESIMD_API
560  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value,
561  __ESIMD_NS::simd<T, SZ>>
562  div(__ESIMD_NS::simd<T, SZ> &remainder, __ESIMD_NS::simd<T, SZ> src0,
563  U src1) {
564  remainder = src0 % src1;
565  return src0 / src1;
566 }
567 
577 template <typename T, int SZ, typename U>
578 __ESIMD_API
579  std::enable_if_t<std::is_integral<T>::value && std::is_integral<U>::value &&
580  __ESIMD_DNS::is_esimd_scalar<U>::value,
581  __ESIMD_NS::simd<T, SZ>>
582  div(__ESIMD_NS::simd<T, SZ> &remainder, U src0,
583  __ESIMD_NS::simd<T, SZ> src1) {
584  remainder = src0 % src1;
585  return src0 / src1;
586 }
587 
597 template <typename RT, typename T0, typename T1>
598 ESIMD_NODEBUG
599  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<RT>::value &&
600  __ESIMD_DNS::is_esimd_scalar<T0>::value &&
601  __ESIMD_DNS::is_esimd_scalar<T1>::value,
602  std::remove_const_t<RT>>
603  div(__ESIMD_NS::simd<std::remove_const_t<RT>, 1> &remainder, T0 src0,
604  T1 src1) {
605  remainder[0] = src0 % src1;
606  return src0 / src1;
607 }
608 
609 // Dot product builtins
610 #if defined(ESIMD_GEN7_5) || defined(ESIMD_GEN8) || defined(ESIMD_GEN8_5) || \
611  defined(ESIMD_GEN9) || defined(ESIMD_GEN9_5)
612 
623 template <typename T0, typename T1, int SZ, typename U,
624  class Sat = __ESIMD_NS::saturation_off_tag>
625 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dp2(__ESIMD_NS::simd<T1, SZ> src0, U src1,
626  Sat sat = {}) {
627  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
628  __ESIMD_NS::simd<float, SZ> Src0 = src0;
629  __ESIMD_NS::simd<float, SZ> Src1 = src1;
630  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp2(Src0.data(), Src1.data());
631  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
632  return Result;
633  else
634  return __ESIMD_NS::saturate<T0>(Result);
635 }
636 
647 template <typename T0, typename T1, int SZ, typename U,
648  class Sat = __ESIMD_NS::saturation_off_tag>
649 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dp3(__ESIMD_NS::simd<T1, SZ> src0, U src1,
650  Sat sat = {}) {
651  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
652  __ESIMD_NS::simd<float, SZ> Src0 = src0;
653  __ESIMD_NS::simd<float, SZ> Src1 = src1;
654  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp3(Src0.data(), Src1.data());
655  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
656  return Result;
657  else
658  return __ESIMD_NS::saturate<T0>(Result);
659 }
660 
671 template <typename T0, typename T1, int SZ, typename U,
672  class Sat = __ESIMD_NS::saturation_off_tag>
673 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dp4(__ESIMD_NS::simd<T1, SZ> src0, U src1,
674  Sat sat = {}) {
675  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
676  __ESIMD_NS::simd<float, SZ> Src0 = src0;
677  __ESIMD_NS::simd<float, SZ> Src1 = src1;
678  __ESIMD_NS::simd<float, SZ> Result = __esimd_dp4(Src0.data(), Src1.data());
679  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
680  return Result;
681  else
682  return __ESIMD_NS::saturate<T0>(Result);
683 }
684 
695 template <typename T0, typename T1, typename U, int SZ,
696  class Sat = __ESIMD_NS::saturation_off_tag>
697 __ESIMD_API __ESIMD_NS::simd<T0, SZ> dph(__ESIMD_NS::simd<T1, SZ> src0, U src1,
698  Sat sat = {}) {
699  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
700  __ESIMD_NS::simd<float, SZ> Src0 = src0;
701  __ESIMD_NS::simd<float, SZ> Src1 = src1;
702  __ESIMD_NS::simd<float, SZ> Result = __esimd_dph(Src0.data(), Src1.data());
703  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
704  return Result;
705  else
706  return __ESIMD_NS::saturate<T0>(Result);
707 }
708 
720 template <typename RT, typename T1, typename T2, int SZ,
721  class Sat = __ESIMD_NS::saturation_off_tag>
722 __ESIMD_API __ESIMD_NS::simd<RT, SZ> line(__ESIMD_NS::simd<T1, 4> src0,
723  __ESIMD_NS::simd<T2, SZ> src1,
724  Sat sat = {}) {
725  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
726 
727  __ESIMD_NS::simd<float, 4> Src0 = src0;
728  __ESIMD_NS::simd<float, SZ> Src1 = src1;
729  __ESIMD_NS::simd<float, SZ> Result = __esimd_line(Src0.data(), Src1.data());
730 
731  __ESIMD_NS::simd<RT, SZ> Result;
732  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
733  return Result;
734  else
735  return __ESIMD_NS::saturate<RT>(Result);
736 }
737 
749 template <typename RT, typename T, int SZ,
750  class Sat = __ESIMD_NS::saturation_off_tag>
751 __ESIMD_API __ESIMD_NS::simd<RT, SZ>
752 line(float P, float Q, __ESIMD_NS::simd<T, SZ> src1, Sat sat = {}) {
753  __ESIMD_NS::simd<float, 4> Src0 = P;
754  Src0(3) = Q;
755  return esimd::line<RT>(Src0, src1, sat);
756 }
757 
758 #else
759 // The old implementation is to generate vISA IRs for dp2/dp3/dp4/dph/line.
760 // Now We change to use direct mul/add, and hope to generate mad instructions
761 // at the end, to still get the performance as good as HW solution.
762 // We rely on "pragma unroll" to get better code.
763 // The only input and return types for these APIs are floats.
764 // In order to be able to use the old emu code, we keep the template argument
765 // for the type, although the type "T" can only be float.
766 // We use enable_if to force the float type only.
767 // If the gen is not specified we warn the programmer that they are potentially
768 // using a less efficient implementation if not on GEN10 or above.
769 
780 template <typename T0, typename T1, int SZ, typename U,
781  class Sat = __ESIMD_NS::saturation_off_tag>
782 ESIMD_NODEBUG ESIMD_INLINE
783  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
784  std::is_floating_point<T1>::value &&
785  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
786  std::is_floating_point<U>::value,
787  __ESIMD_NS::simd<T0, SZ>>
788  dp2(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
789  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
790 
791  __ESIMD_NS::simd<float, SZ> Src1 = src1;
792  __ESIMD_NS::simd<float, SZ> Result;
793 #pragma unroll
794  for (int i = 0; i < SZ; i += 4) {
795  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1];
796  }
797  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
798  return Result;
799  else
800  return __ESIMD_NS::saturate<T1>(Result);
801 }
802 
813 template <typename T0, typename T1, int SZ, typename U,
814  class Sat = __ESIMD_NS::saturation_off_tag>
815 ESIMD_NODEBUG ESIMD_INLINE
816  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
817  std::is_floating_point<T1>::value &&
818  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
819  std::is_floating_point<U>::value,
820  __ESIMD_NS::simd<T0, SZ>>
821  dp3(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
822  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
823 
824  __ESIMD_NS::simd<float, SZ> Src1 = src1;
825  __ESIMD_NS::simd<float, SZ> Result;
826 #pragma unroll
827  for (int i = 0; i < SZ; i += 4) {
828  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
829  src0[i + 2] * Src1[i + 2];
830  }
831  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
832  return Result;
833  else
834  return __ESIMD_NS::saturate<T1>(Result);
835 }
836 
847 template <typename T0, typename T1, int SZ, typename U,
848  class Sat = __ESIMD_NS::saturation_off_tag>
849 ESIMD_NODEBUG ESIMD_INLINE
850  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T1>::value &&
851  std::is_floating_point<T1>::value &&
852  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
853  std::is_floating_point<U>::value,
854  __ESIMD_NS::simd<T0, SZ>>
855  dp4(__ESIMD_NS::simd<T1, SZ> src0, U src1, Sat sat = {}) {
856  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
857 
858  __ESIMD_NS::simd<T1, SZ> Src1 = src1;
859  __ESIMD_NS::simd<float, SZ> Result;
860 #pragma unroll
861  for (int i = 0; i < SZ; i += 4) {
862  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
863  src0[i + 2] * Src1[i + 2] +
864  src0[i + 3] * Src1[i + 3];
865  }
866  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
867  return Result;
868  else
869  return __ESIMD_NS::saturate<T1>(Result);
870 }
871 
882 template <typename T, typename U, int SZ,
883  class Sat = __ESIMD_NS::saturation_off_tag>
884 ESIMD_NODEBUG ESIMD_INLINE
885  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
886  std::is_floating_point<T>::value &&
887  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
888  std::is_floating_point<U>::value,
889  __ESIMD_NS::simd<T, SZ>>
890  dph(__ESIMD_NS::simd<T, SZ> src0, U src1, Sat sat = {}) {
891  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
892 
893  __ESIMD_NS::simd<float, SZ> Src1 = src1;
894  __ESIMD_NS::simd<float, SZ> Result;
895 #pragma unroll
896  for (int i = 0; i < SZ; i += 4) {
897  Result.select<4, 1>(i) = src0[i] * Src1[i] + src0[i + 1] * Src1[i + 1] +
898  src0[i + 2] * Src1[i + 2] + 1.0 * Src1[i + 3];
899  }
900  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
901  return Result;
902  else
903  return __ESIMD_NS::saturate<T>(Result);
904 }
905 
916 template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
917 ESIMD_NODEBUG
918  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
919  std::is_floating_point<T>::value,
920  __ESIMD_NS::simd<T, SZ>>
921  line(__ESIMD_NS::simd<T, 4> src0, __ESIMD_NS::simd<T, SZ> src1,
922  Sat sat = {}) {
923  static_assert(SZ % 4 == 0, "result size is not a multiple of 4");
924 
925  __ESIMD_NS::simd<T, SZ> Src1 = src1;
926  __ESIMD_NS::simd<T, SZ> Result;
927 #pragma unroll
928  for (int i = 0; i < SZ; i += 4) {
929  Result.select<4, 1>(i) = src0[0] * src1[i] + src0[3];
930  }
931 
932  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
933  return Result;
934  else
935  return __ESIMD_NS::saturate<T>(Result);
936 }
937 
949 template <typename T, int SZ, class Sat = __ESIMD_NS::saturation_off_tag>
950 ESIMD_NODEBUG
951  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
952  std::is_floating_point<T>::value,
953  __ESIMD_NS::simd<T, SZ>>
954  line(float P, float Q, __ESIMD_NS::simd<T, SZ> src1, Sat sat = {}) {
955  __ESIMD_NS::simd<T, 4> Src0 = P;
956  Src0(3) = Q;
957  return esimd::line<T>(Src0, src1, sat);
958 }
959 
960 #endif
961 
968 template <typename T, int SZ>
969 __ESIMD_API __ESIMD_NS::simd<T, SZ> frc(__ESIMD_NS::simd<T, SZ> src0) {
970  __ESIMD_NS::simd<float, SZ> Src0 = src0;
971  return __esimd_frc(Src0.data());
972 }
973 
979 template <typename T> __ESIMD_API T frc(T src0) {
980  __ESIMD_NS::simd<T, 1> Src0 = src0;
981  __ESIMD_NS::simd<T, 1> Result = esimd::frc<T>(Src0);
982  return Result[0];
983 }
984 
985 // lzd
986 template <typename RT, typename T0, int SZ,
987  class Sat = __ESIMD_NS::saturation_off_tag>
988 __ESIMD_API __ESIMD_NS::simd<RT, SZ> lzd(__ESIMD_NS::simd<T0, SZ> src0,
989  Sat sat = {}) {
990  // Saturation parameter ignored
991  __ESIMD_NS::simd<uint, SZ> Src0 = src0;
992  return __esimd_lzd<uint>(Src0.data());
993 }
994 
995 template <typename RT, typename T0, class Sat = __ESIMD_NS::saturation_off_tag>
996 ESIMD_NODEBUG
997  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<RT>::value &&
998  __ESIMD_DNS::is_esimd_scalar<T0>::value,
999  std::remove_const_t<RT>>
1000  lzd(T0 src0, Sat sat = {}) {
1001  __ESIMD_NS::simd<T0, 1> Src0 = src0;
1002  __ESIMD_NS::simd<RT, 1> Result = esimd::lzd<RT>(Src0);
1003  return Result[0];
1004 }
1005 
1006 // lrp
1007 #if defined(ESIMD_GEN7_5) || defined(ESIMD_GEN8) || defined(ESIMD_GEN8_5) || \
1008  defined(ESIMD_GEN9) || defined(ESIMD_GEN9_5)
1009 
1010 template <int SZ, typename U, typename V,
1011  class Sat = __ESIMD_NS::saturation_off_tag>
1012 __ESIMD_API __ESIMD_NS::simd<float, SZ> lrp(__ESIMD_NS::simd<float, SZ> src0,
1013  U src1, V src2, Sat sat = {}) {
1014  static_assert(SZ >= 4 && (SZ & 0x3) == 0,
1015  "vector size must be a multiple of 4");
1016  __ESIMD_NS::simd<float, SZ> Src1 = src1;
1017  __ESIMD_NS::simd<float, SZ> Src2 = src2;
1018  __ESIMD_NS::simd<float, SZ> Result =
1019  __esimd_lrp<SZ>(src0.data(), Src1.data(), Src2.data());
1020 
1021  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1022  return Result;
1023  else
1024  return __ESIMD_NS::saturate<float>(Result);
1025 }
1026 
1027 #else
1028 
1029 // The old implementation is to generate vISA IRs for lrp.
1030 // Now We change to use direct mul/add, and hope to generate mad instructions
1031 // at the end, to still get the performance as good as HW solution.
1032 // The only input and return types for these APIs are floats.
1033 // In order to be able to use the old emu code, we keep the template argument
1034 // for the type, although the type "T" can only be float.
1035 // We use enable_if to force the float type only.
1036 // If the gen is not specified we warn the programmer that they are potentially
1037 // using less efficient implementation.
1038 template <typename T, int SZ, typename U, typename V,
1039  class Sat = __ESIMD_NS::saturation_off_tag>
1040 ESIMD_NODEBUG ESIMD_INLINE
1041  std::enable_if_t<__ESIMD_DNS::is_fp_or_dword_type<T>::value &&
1042  std::is_floating_point<T>::value &&
1043  __ESIMD_DNS::is_fp_or_dword_type<U>::value &&
1044  std::is_floating_point<U>::value,
1045  __ESIMD_NS::simd<T, SZ>>
1046  lrp(__ESIMD_NS::simd<T, SZ> src0, U src1, V src2, Sat sat = {}) {
1047 
1048  __ESIMD_NS::simd<float, SZ> Src1 = src1;
1049  __ESIMD_NS::simd<float, SZ> Src2 = src2;
1050  __ESIMD_NS::simd<float, SZ> Result;
1051  Result = Src1 * src0 + Src2 * (1.0f - src0);
1052  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1053  return Result;
1054  else
1055  return __ESIMD_NS::saturate<T>(Result);
1056 }
1057 #endif
1058 
1060 
1063 
1065 template <typename T0, typename T1, int SZ>
1066 __ESIMD_API __ESIMD_NS::simd<T0, SZ> bf_reverse(__ESIMD_NS::simd<T1, SZ> src0) {
1067  __ESIMD_NS::simd<unsigned, SZ> Src0 = src0;
1068  return __esimd_bfrev<unsigned>(Src0.data());
1069 }
1070 
1072 template <typename T0, typename T1>
1073 ESIMD_NODEBUG
1074  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1075  __ESIMD_DNS::is_esimd_scalar<T1>::value,
1076  std::remove_const_t<T0>>
1077  bf_reverse(T1 src0) {
1078  __ESIMD_NS::simd<T1, 1> Src0 = src0;
1079  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_reverse<T0>(Src0);
1080  return Result[0];
1081 }
1082 
1084 template <typename T0, typename T1, int SZ, typename U, typename V, typename W>
1085 ESIMD_NODEBUG ESIMD_INLINE
1086  std::enable_if_t<std::is_integral<T1>::value, __ESIMD_NS::simd<T0, SZ>>
1087  bf_insert(U src0, V src1, W src2, __ESIMD_NS::simd<T1, SZ> src3) {
1088  typedef typename __ESIMD_DNS::dword_type<T1> DT1;
1089  static_assert(std::is_integral<DT1>::value && sizeof(DT1) == sizeof(int),
1090  "operand conversion failed");
1091  __ESIMD_NS::simd<DT1, SZ> Src0 = src0;
1092  __ESIMD_NS::simd<DT1, SZ> Src1 = src1;
1093  __ESIMD_NS::simd<DT1, SZ> Src2 = src2;
1094  __ESIMD_NS::simd<DT1, SZ> Src3 = src3;
1095 
1096  return __esimd_bfi<DT1>(Src0.data(), Src1.data(), Src2.data(), Src3.data());
1097 }
1098 
1100 template <typename T0, typename T1, typename T2, typename T3, typename T4>
1101 ESIMD_NODEBUG
1102  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1103  __ESIMD_DNS::is_esimd_scalar<T4>::value,
1104  std::remove_const_t<T0>>
1105  bf_insert(T1 src0, T2 src1, T3 src2, T4 src3) {
1106  __ESIMD_NS::simd<T4, 1> Src3 = src3;
1107  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_insert<T0>(src0, src1, src2, Src3);
1108  return Result[0];
1109 }
1110 
1112 template <typename T0, typename T1, int SZ, typename U, typename V>
1113 ESIMD_NODEBUG ESIMD_INLINE
1114  std::enable_if_t<std::is_integral<T1>::value, __ESIMD_NS::simd<T0, SZ>>
1115  bf_extract(U src0, V src1, __ESIMD_NS::simd<T1, SZ> src2) {
1116  typedef typename __ESIMD_DNS::dword_type<T1> DT1;
1117  static_assert(std::is_integral<DT1>::value && sizeof(DT1) == sizeof(int),
1118  "operand conversion failed");
1119  __ESIMD_NS::simd<DT1, SZ> Src0 = src0;
1120  __ESIMD_NS::simd<DT1, SZ> Src1 = src1;
1121  __ESIMD_NS::simd<DT1, SZ> Src2 = src2;
1122 
1123  return __esimd_sbfe<DT1>(Src0.data(), Src1.data(), Src2.data());
1124 }
1125 
1127 template <typename T0, typename T1, typename T2, typename T3>
1128 ESIMD_NODEBUG
1129  ESIMD_INLINE std::enable_if_t<__ESIMD_DNS::is_esimd_scalar<T0>::value &&
1130  __ESIMD_DNS::is_esimd_scalar<T3>::value,
1131  std::remove_const_t<T0>>
1132  bf_extract(T1 src0, T2 src1, T3 src2) {
1133  __ESIMD_NS::simd<T3, 1> Src2 = src2;
1134  __ESIMD_NS::simd<T0, 1> Result = esimd::bf_extract<T0>(src0, src1, Src2);
1135  return Result[0];
1136 }
1137 
1139 
1142 
1143 // sincos
1144 template <int SZ, typename U, class Sat = __ESIMD_NS::saturation_off_tag>
1145 __ESIMD_API __ESIMD_NS::simd<float, SZ>
1146 sincos(__ESIMD_NS::simd<float, SZ> &dstcos, U src0, Sat sat = {}) {
1147  dstcos = __ESIMD_NS::cos(src0, sat);
1148  return __ESIMD_NS::sin(src0, sat);
1149 }
1150 
1151 // atan
1152 
1154 namespace detail {
1155 constexpr double HDR_CONST_PI = 3.1415926535897932384626433832795;
1156 } // namespace detail
1158 
1159 template <typename T, int SZ>
1160 ESIMD_NODEBUG ESIMD_INLINE
1161  std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
1162  atan(__ESIMD_NS::simd<T, SZ> src0) {
1163  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
1164 
1165  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
1166  __ESIMD_NS::simd_mask<SZ> Gt1 = Src0 > T(1.0);
1167 
1168  Src0.merge(__ESIMD_NS::inv(Src0), Gt1);
1169 
1170  __ESIMD_NS::simd<T, SZ> Src0P2 = Src0 * Src0;
1171  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
1172 
1173  __ESIMD_NS::simd<T, SZ> Result =
1174  (Src0P4 * T(0.185696) + ((Src0 * T(0.787997) + T(0.63693)) * Src0P2) +
1175  Src0) /
1176  (((((Src0 * -T(0.000121387) + T(0.00202308)) * Src0P2) +
1177  (Src0 * -T(0.0149145)) + T(0.182569)) *
1178  Src0P4) +
1179  ((Src0 * T(0.395889) + T(1.12158)) * Src0P2) + (Src0 * T(0.636918)) +
1180  T(1.0));
1181 
1182  Result.merge(Result - T(detail::HDR_CONST_PI / 2.0), Gt1);
1183  Result.merge(Result, Neg);
1184  return Result;
1185 }
1186 
1187 template <typename T>
1188 __ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> atan(T src0) {
1189  __ESIMD_NS::simd<T, 1> Src0 = src0;
1190  __ESIMD_NS::simd<T, 1> Result = esimd::atan(Src0);
1191  return Result[0];
1192 }
1193 
1194 // acos
1195 
1196 template <typename T, int SZ>
1197 ESIMD_NODEBUG ESIMD_INLINE
1198  std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
1199  acos(__ESIMD_NS::simd<T, SZ> src0) {
1200  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
1201 
1202  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
1203  __ESIMD_NS::simd_mask<SZ> TooBig = Src0 >= T(0.999998);
1204 
1205  // Replace oversized values to ensure no possibility of sqrt of
1206  // a negative value later
1207  Src0.merge(T(0.0), TooBig);
1208 
1209  __ESIMD_NS::simd<T, SZ> Src01m = T(1.0) - Src0;
1210 
1211  __ESIMD_NS::simd<T, SZ> Src0P2 = Src01m * Src01m;
1212  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
1213 
1214  __ESIMD_NS::simd<T, SZ> Result =
1215  (((Src01m * T(0.015098965761299077) - T(0.005516443930088506)) * Src0P4) +
1216  ((Src01m * T(0.047654245891495528) + T(0.163910606547823220)) * Src0P2) +
1217  Src01m * T(2.000291665285952400) - T(0.000007239283986332)) *
1218  __ESIMD_NS::rsqrt(Src01m * T(2.0));
1219 
1220  Result.merge(T(0.0), TooBig);
1221  Result.merge(T(detail::HDR_CONST_PI) - Result, Neg);
1222  return Result;
1223 }
1224 
1225 template <typename T>
1226 __ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> acos(T src0) {
1227  __ESIMD_NS::simd<T, 1> Src0 = src0;
1228  __ESIMD_NS::simd<T, 1> Result = esimd::acos(Src0);
1229  return Result[0];
1230 }
1231 
1232 // asin
1233 
1234 template <typename T, int SZ>
1235 ESIMD_NODEBUG ESIMD_INLINE
1236  std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
1237  asin(__ESIMD_NS::simd<T, SZ> src0) {
1238  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
1239 
1240  __ESIMD_NS::simd<T, SZ> Result =
1241  T(detail::HDR_CONST_PI / 2.0) - esimd::acos(__ESIMD_NS::abs(src0));
1242 
1243  Result.merge(-Result, Neg);
1244  return Result;
1245 }
1246 
1247 template <typename T>
1248 __ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> asin(T src0) {
1249  __ESIMD_NS::simd<T, 1> Src0 = src0;
1250  __ESIMD_NS::simd<T, 1> Result = esimd::asin(Src0);
1251  return Result[0];
1252 }
1254 
1257 
1258 /* atan2_fast - a fast atan2 implementation */
1259 /* vector input */
1260 template <int N>
1261 __ESIMD_NS::simd<float, N> atan2_fast(__ESIMD_NS::simd<float, N> y,
1262  __ESIMD_NS::simd<float, N> x);
1263 /* scalar input */
1264 template <typename T> float atan2_fast(T y, T x);
1265 
1266 /* atan2 - atan2 implementation */
1267 /* For Vector input */
1268 template <int N>
1269 __ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
1270  __ESIMD_NS::simd<float, N> x);
1271 /* scalar Input */
1272 template <typename T> float atan2(T y, T x);
1273 
1274 /* fmod: */
1275 /* vector input */
1276 template <int N>
1277 __ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
1278  __ESIMD_NS::simd<float, N> x);
1279 /* scalar Input */
1280 template <typename T> float fmod(T y, T x);
1281 
1282 /* sin_emu - EU emulation for sin(x) */
1283 /* For Vector input */
1284 template <int N>
1285 __ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x);
1286 /* scalar Input */
1287 template <typename T> float sin_emu(T x);
1288 
1289 /* cos_emu - EU emulation for cos(x) */
1290 /* For Vector input */
1291 template <int N>
1292 __ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x);
1293 
1294 /* scalar Input */
1295 template <typename T> float cos_emu(T x);
1296 
1297 /* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
1298 /* float input */
1299 float tanh_cody_waite(float x);
1300 /* vector input */
1301 template <int N>
1302 __ESIMD_NS::simd<float, N> tanh_cody_waite(__ESIMD_NS::simd<float, N> x);
1303 /* tanh - opencl like implementation for tanh(x) */
1304 /* float input */
1305 float tanh(float x);
1306 /* vector input */
1307 template <int N> __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x);
1308 
1309 /* ------------------------- Extended Math Routines
1310  * -------------------------------------------------*/
1311 
1313 
1314 namespace detail {
1315 static auto constexpr CONST_PI = 3.14159f;
1316 static auto constexpr CMPI = 3.14159265f;
1317 } // namespace detail
1318 
1320 
1321 // For vector input
1322 template <int N>
1323 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1324 atan2_fast(__ESIMD_NS::simd<float, N> y, __ESIMD_NS::simd<float, N> x) {
1325  __ESIMD_NS::simd<float, N> a0;
1326  __ESIMD_NS::simd<float, N> a1;
1327  __ESIMD_NS::simd<float, N> atan2;
1328 
1329  __ESIMD_NS::simd_mask<N> mask = (y >= 0.0f);
1330  a0.merge(detail::CONST_PI * 0.5f, detail::CONST_PI * 1.5f, mask);
1331  a1.merge(0, detail::CONST_PI * 2.0f, mask);
1332 
1333  a1.merge(detail::CONST_PI, x < 0.0f);
1334 
1335  __ESIMD_NS::simd<float, N> xy = x * y;
1336  __ESIMD_NS::simd<float, N> x2 = x * x;
1337  __ESIMD_NS::simd<float, N> y2 = y * y;
1338 
1339  /* smallest such that 1.0+CONST_DBL_EPSILON != 1.0 */
1340  constexpr auto CONST_DBL_EPSILON = 0.00001f;
1341 
1342  a0 -= (xy / (y2 + x2 * 0.28f + CONST_DBL_EPSILON));
1343  a1 += (xy / (x2 + y2 * 0.28f + CONST_DBL_EPSILON));
1344 
1345  atan2.merge(a1, a0, y2 <= x2);
1346  return atan2;
1347 }
1348 
1349 // For Scalar Input
1350 template <> ESIMD_INLINE float atan2_fast(float y, float x) {
1351  __ESIMD_NS::simd<float, 1> vy = y;
1352  __ESIMD_NS::simd<float, 1> vx = x;
1353  __ESIMD_NS::simd<float, 1> atan2 = esimd::atan2_fast(vy, vx);
1354  return atan2[0];
1355 }
1356 
1357 // atan2
1358 // For Vector input
1359 template <int N>
1360 ESIMD_INLINE __ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
1361  __ESIMD_NS::simd<float, N> x) {
1362  __ESIMD_NS::simd<float, N> v_distance;
1363  __ESIMD_NS::simd<float, N> v_y0;
1364  __ESIMD_NS::simd<float, N> atan2;
1365  __ESIMD_NS::simd_mask<N> mask;
1366 
1367  mask = (x < 0);
1368  v_y0.merge(detail::CONST_PI, 0, mask);
1369  v_distance = __ESIMD_NS::sqrt(x * x + y * y);
1370  mask = (__ESIMD_NS::abs<float>(y) < 0.000001f);
1371  atan2.merge(v_y0, (2 * esimd::atan((v_distance - x) / y)), mask);
1372  return atan2;
1373 }
1374 
1375 // For Scalar Input
1376 template <> ESIMD_INLINE float atan2(float y, float x) {
1377  float v_distance;
1378  float v_y0;
1379  __ESIMD_NS::simd<float, 1> atan2;
1380  __ESIMD_NS::simd_mask<1> mask;
1381 
1382  mask = (x < 0);
1383  v_y0 = mask[0] ? detail::CONST_PI : 0;
1384  v_distance = __ESIMD_NS::sqrt<float>(x * x + y * y);
1385  mask = (__ESIMD_NS::abs<float>(y) < 0.000001f);
1386  atan2.merge(v_y0, (2 * esimd::atan((v_distance - x) / y)), mask);
1387  return atan2[0];
1388 }
1389 
1390 // fmod:
1391 // For Vector input
1392 template <int N>
1393 ESIMD_INLINE __ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
1394  __ESIMD_NS::simd<float, N> x) {
1395  __ESIMD_NS::simd<float, N> abs_x = __ESIMD_NS::abs(x);
1396  __ESIMD_NS::simd<float, N> abs_y = __ESIMD_NS::abs(y);
1397  auto fmod_sign_mask = (y.template bit_cast_view<int32_t>()) & 0x80000000;
1398 
1399  __ESIMD_NS::simd<float, N> reminder =
1400  abs_y - abs_x * __ESIMD_NS::trunc<float>(abs_y / abs_x);
1401 
1402  abs_x.merge(0.0, reminder >= 0);
1403  __ESIMD_NS::simd<float, N> fmod = reminder + abs_x;
1404  __ESIMD_NS::simd<float, N> fmod_abs = __ESIMD_NS::abs(fmod);
1405 
1406  auto fmod_bits =
1407  (fmod_abs.template bit_cast_view<int32_t>()) | fmod_sign_mask;
1408  return fmod_bits.template bit_cast_view<float>();
1409 }
1410 
1411 // For Scalar Input
1412 template <> ESIMD_INLINE float fmod(float y, float x) {
1413  return fmod(__ESIMD_NS::simd<float, 1>(y), __ESIMD_NS::simd<float, 1>(x))[0];
1414 }
1415 
1416 // sin_emu - EU emulation for sin(x)
1417 // For Vector input
1418 template <int N>
1419 ESIMD_INLINE __ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x) {
1420  __ESIMD_NS::simd<float, N> x1;
1421  __ESIMD_NS::simd<float, N> x2;
1422  __ESIMD_NS::simd<float, N> t3;
1423 
1424  __ESIMD_NS::simd<float, N> sign;
1425  __ESIMD_NS::simd<float, N> fTrig;
1426  __ESIMD_NS::simd<float, N> TwoPI(6.2831853f);
1427  __ESIMD_NS::simd<float, N> CmpI(detail::CMPI);
1428  __ESIMD_NS::simd<float, N> OneP(1.f);
1429  __ESIMD_NS::simd<float, N> OneN(-1.f);
1430 
1431  x = esimd::fmod(x, TwoPI);
1432 
1433  x1.merge(CmpI - x, x - CmpI, (x <= detail::CMPI));
1434  x1.merge(x, (x <= detail::CMPI * 0.5f));
1435  x1.merge(CmpI * 2 - x, (x > detail::CMPI * 1.5f));
1436 
1437  sign.merge(OneN, OneP, (x > detail::CMPI));
1438 
1439  x2 = x1 * x1;
1440  t3 = x2 * x1 * 0.1666667f;
1441 
1442  fTrig =
1443  x1 + t3 * (OneN + x2 * 0.05f *
1444  (OneP + x2 * 0.0238095f *
1445  (OneN + x2 * 0.0138889f *
1446  (OneP - x2 * 0.0090909f))));
1447  fTrig *= sign;
1448  return fTrig;
1449 }
1450 
1451 // scalar Input
1452 template <typename T> ESIMD_INLINE float sin_emu(T x0) {
1453  __ESIMD_NS::simd<float, 1> x1;
1454  __ESIMD_NS::simd<float, 1> x2;
1455  __ESIMD_NS::simd<float, 1> t3;
1456 
1457  __ESIMD_NS::simd<float, 1> sign;
1458  __ESIMD_NS::simd<float, 1> fTrig;
1459  float TwoPI = detail::CMPI * 2.0f;
1460 
1461  __ESIMD_NS::simd<float, 1> x = esimd::fmod(x0, TwoPI);
1462 
1463  __ESIMD_NS::simd<float, 1> CmpI(detail::CMPI);
1464  __ESIMD_NS::simd<float, 1> OneP(1.f);
1465  __ESIMD_NS::simd<float, 1> OneN(-1.f);
1466 
1467  x1.merge(CmpI - x, x - CmpI, (x <= detail::CMPI));
1468  x1.merge(x, (x <= detail::CMPI * 0.5f));
1469  x1.merge(CmpI * 2.0f - x, (x > detail::CMPI * 1.5f));
1470 
1471  sign.merge(OneN, OneP, (x > detail::CMPI));
1472 
1473  x2 = x1 * x1;
1474  t3 = x2 * x1 * 0.1666667f;
1475 
1476  fTrig =
1477  x1 + t3 * (OneN + x2 * 0.05f *
1478  (OneP + x2 * 0.0238095f *
1479  (OneN + x2 * 0.0138889f *
1480  (OneP - x2 * 0.0090909f))));
1481  fTrig *= sign;
1482  return fTrig[0];
1483 }
1484 
1485 // cos_emu - EU emulation for sin(x)
1486 // For Vector input
1487 template <int N>
1488 ESIMD_INLINE __ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x) {
1489  __ESIMD_NS::simd<float, N> x1;
1490  __ESIMD_NS::simd<float, N> x2;
1491  __ESIMD_NS::simd<float, N> t2;
1492  __ESIMD_NS::simd<float, N> t3;
1493 
1494  __ESIMD_NS::simd<float, N> sign;
1495  __ESIMD_NS::simd<float, N> fTrig;
1496  __ESIMD_NS::simd<float, N> TwoPI(6.2831853f);
1497  __ESIMD_NS::simd<float, N> CmpI(detail::CMPI);
1498  __ESIMD_NS::simd<float, N> OneP(1.f);
1499  __ESIMD_NS::simd<float, N> OneN(-1.f);
1500 
1501  x = esimd::fmod(x, TwoPI);
1502 
1503  x1.merge(x - detail::CMPI * 0.5f, CmpI * 1.5f - x, (x <= detail::CMPI));
1504  x1.merge(CmpI * 0.5f - x, (x <= detail::CMPI * 0.5f));
1505  x1.merge(x - detail::CMPI * 1.5f, (x > detail::CMPI * 1.5f));
1506 
1507  sign.merge(1, -1, ((x < detail::CMPI * 0.5f) | (x >= detail::CMPI * 1.5f)));
1508 
1509  x2 = x1 * x1;
1510  t3 = x2 * x1 * 0.1666667f;
1511  fTrig =
1512  x1 + t3 * (OneN + x2 * 0.05f *
1513  (OneP + x2 * 0.0238095f *
1514  (OneN + x2 * 0.0138889f *
1515  (OneP - x2 * 0.0090909f))));
1516  fTrig *= sign;
1517  return fTrig;
1518 }
1519 
1520 // scalar Input
1521 template <typename T> ESIMD_INLINE float cos_emu(T x0) {
1522  __ESIMD_NS::simd<float, 1> x1;
1523  __ESIMD_NS::simd<float, 1> x2;
1524  __ESIMD_NS::simd<float, 1> t3;
1525 
1526  __ESIMD_NS::simd<float, 1> sign;
1527  __ESIMD_NS::simd<float, 1> fTrig;
1528  float TwoPI = detail::CMPI * 2.0f;
1529 
1530  __ESIMD_NS::simd<float, 1> x = esimd::fmod(x0, TwoPI);
1531 
1532  __ESIMD_NS::simd<float, 1> CmpI(detail::CMPI);
1533  __ESIMD_NS::simd<float, 1> OneP(1.f);
1534  __ESIMD_NS::simd<float, 1> OneN(-1.f);
1535 
1536  x1.merge(x - detail::CMPI * 0.5f, CmpI * 1.5f - x, (x <= detail::CMPI));
1537  x1.merge(CmpI * 0.5f - x, (x <= detail::CMPI * 0.5f));
1538  x1.merge(x - detail::CMPI * 1.5f, (x > detail::CMPI * 1.5f));
1539 
1540  sign.merge(OneP, OneN,
1541  ((x < detail::CMPI * 0.5f) | (x >= detail::CMPI * 1.5f)));
1542 
1543  x2 = x1 * x1;
1544  t3 = x2 * x1 * 0.1666667f;
1545  fTrig =
1546  x1 + t3 * (OneN + x2 * 0.05f *
1547  (OneP + x2 * 0.0238095f *
1548  (OneN + x2 * 0.0138889f *
1549  (OneP - x2 * 0.0090909f))));
1550  fTrig *= sign;
1551  return fTrig[0];
1552 }
1553 
1555 namespace detail {
1556 
1557 template <int N>
1558 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1559 tanh_cody_waite_impl(__ESIMD_NS::simd<float, N> x) {
1560  /*
1561  * 0 x_small x_medium x_large
1562  * | x | rational polynomial | 1 - 2/(1 + exp(2*x)) | 1
1563  *
1564  * rational polynomial for single precision = x + x * (g * (p[1] * g + p[0]) /
1565  * (g + q[0]) g = x^2 p0 = -0.82377 28127 E+00 p1 = -0.38310 10665 E-02 q0 =
1566  * 0.24713 19654 E+01 q1 = 1.00000 00000 E+00
1567  *
1568  */
1569 
1570  constexpr float p0 = -0.8237728127E+00f;
1571  constexpr float p1 = -0.3831010665E-02f;
1572  constexpr float q0 = 0.2471319654E+01f;
1573  constexpr float q1 = 1.0000000000E+00f;
1574  constexpr float xsmall = 4.22863966691620432990E-04f;
1575  constexpr float xmedium = 0.54930614433405484570f;
1576  constexpr float xlarge = 8.66433975699931636772f;
1577  constexpr float log2E = 1.442695f; // same as esimd::log(e)
1578 
1579  using RT = __ESIMD_NS::simd<float, N>;
1580 
1581  RT absX = __ESIMD_NS::abs(x);
1582  RT g = absX * absX;
1583 
1584  RT sign;
1585  sign.merge(-1.f, 1.f, x < 0.f);
1586 
1587  auto isLarge = absX > xlarge;
1588  auto minor = absX <= xlarge;
1589  auto isGtMed = minor & (absX > xmedium);
1590  auto isGtSmall = (absX > xsmall) & (absX <= xmedium);
1591 
1592  RT res;
1593  res.merge(sign, x, isLarge);
1594  auto temp = __ESIMD_NS::exp(absX * 2.0f * log2E) + 1.f;
1595  temp = ((temp - 2.f) / temp) * sign;
1596  res.merge(temp, isGtMed);
1597  res.merge((absX + absX * g * (g * p1 + p0) / (g + q0)) * sign, isGtSmall);
1598 
1599  return res;
1600 }
1601 
1602 template <int N>
1603 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1604 tanh_impl(__ESIMD_NS::simd<float, N> x) {
1605  /*
1606  * 0 x_small x_large
1607  * | x | ( exp(x) - exp(-x) ) / ( exp(x) + exp(-x) ) | 1
1608  *
1609  */
1610 
1611  constexpr float xsmall = 0.000045f; // same as exp(-10.0f)
1612  constexpr float xlarge = 88.f;
1613  constexpr float log2E = 1.442695f; // same as esimd::log(e)
1614 
1615  using RT = __ESIMD_NS::simd<float, N>;
1616 
1617  RT absX = __ESIMD_NS::abs(x);
1618 
1619  RT sign;
1620  sign.merge(-1.f, 1.f, x < 0.f);
1621 
1622  auto isLarge = (absX > xlarge);
1623  auto isLessE = (absX <= xlarge);
1624 
1625  RT res;
1626  res.merge(sign, x, isLarge);
1627 
1628  RT exp;
1629  exp = __ESIMD_NS::exp(absX * 2.f * log2E);
1630 
1631  res.merge(((exp - 1.f) / (exp + 1.f)) * sign, (absX > xsmall) & isLessE);
1632 
1633  return res;
1634 }
1635 } // namespace detail
1637 
1638 /* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
1639 /* float input */
1640 ESIMD_INLINE float tanh_cody_waite(float x) {
1641  return detail::tanh_cody_waite_impl(__ESIMD_NS::simd<float, 1>(x))[0];
1642 }
1643 /* vector input */
1644 template <int N>
1645 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1646 tanh_cody_waite(__ESIMD_NS::simd<float, N> x) {
1647  return detail::tanh_cody_waite_impl(x);
1648 }
1649 
1650 /* tanh - opencl like implementation for tanh(x) */
1651 /* float input */
1652 ESIMD_INLINE float tanh(float x) {
1653  return esimd::detail::tanh_impl(__ESIMD_NS::simd<float, 1>(x))[0];
1654 }
1655 /* vector input */
1656 template <int N>
1657 ESIMD_INLINE __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x) {
1658  return esimd::detail::tanh_impl(x);
1659 }
1660 
1661 template <typename T, int N>
1662 __ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> v1,
1663  __ESIMD_NS::simd<T, N> v2) {
1664  auto retv = __esimd_dp4<T, N>(v1.data(), v2.data());
1665  return retv;
1666 }
1667 
1669 
1673 
1686 template <argument_type src1_precision, argument_type src2_precision,
1687  typename T, int systolic_depth, int repeat_count, typename T0,
1688  typename T1, typename T2, int N, int N1, int N2,
1689  typename Sat = __ESIMD_NS::saturation_off_tag>
1690 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::native::dpas()")
1691 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1692  __ESIMD_NS::simd<T0, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1693  __ESIMD_NS::simd<T2, N2> src2,
1694  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1695  auto result =
1696  __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T0, T1, T2,
1697  src1_precision, src2_precision>(src0, src1, src2);
1698  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1699  return result;
1700  else
1701  return __ESIMD_NS::saturate<T>(result);
1702 }
1703 
1714 template <argument_type src1_precision, argument_type src2_precision,
1715  int systolic_depth, int repeat_count, typename T, typename T1,
1716  typename T2, int N, int N1, int N2,
1717  typename Sat = __ESIMD_NS::saturation_off_tag>
1718 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
1719 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1720  __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1721  __ESIMD_NS::simd<T2, N2> src2,
1722  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1723  return dpas<src1_precision, src2_precision, T, systolic_depth, repeat_count>(
1724  src0, src1, src2, sat);
1725 }
1726 
1735 template <argument_type src1_precision, argument_type src2_precision,
1736  int systolic_depth, int repeat_count, typename T, typename T1,
1737  typename T2, int N, int N1, int N2,
1738  typename Sat = __ESIMD_NS::saturation_off_tag>
1739 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
1740 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1741  __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
1742  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1743 
1744  __ESIMD_NS::simd<T, N> result =
1745  __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T1, T2,
1746  src1_precision, src2_precision>(src1, src2);
1747 
1748  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1749  return result;
1750  else
1751  return __ESIMD_NS::saturate<T>(result);
1752 }
1753 
1764 template <argument_type src1_precision, argument_type src2_precision,
1765  int systolic_depth, int repeat_count, typename T, typename T1,
1766  typename T2, int N, int N1, int N2,
1767  typename Sat = __ESIMD_NS::saturation_off_tag>
1768 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
1769 __ESIMD_API __ESIMD_NS::simd<T, N> dpasw(
1770  __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1771  __ESIMD_NS::simd<T2, N2> src2,
1772  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1773 
1774  __ESIMD_NS::simd<T, N> result =
1775  __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
1776  src1_precision, src2_precision>(src0, src1, src2);
1777  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1778  return result;
1779  else
1780  return __ESIMD_NS::saturate<T>(result);
1781 }
1782 
1791 template <argument_type src1_precision, argument_type src2_precision,
1792  int systolic_depth, int repeat_count, typename T, typename T1,
1793  typename T2, int N, int N1, int N2,
1794  typename Sat = __ESIMD_NS::saturation_off_tag>
1795 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
1796 __ESIMD_API __ESIMD_NS::simd<T, N> dpasw2(
1797  __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
1798  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1799  constexpr bool is_4xhf =
1800  std::is_same_v<T, __ESIMD_DNS::__raw_t<sycl::half>> &&
1801  src1_precision == src2_precision && src1_precision == argument_type::FP16;
1802 
1803  __ESIMD_NS::simd<T, N> result =
1804  __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
1805  src1_precision, src2_precision>(src1, src2);
1806 
1807  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1808  return result;
1809  else
1810  return __ESIMD_NS::saturate<T>(result);
1811 }
1813 
1814 } // namespace ext::intel::experimental::esimd
1815 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1816 } // 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:1105
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:1796
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:979
sycl::_V1::ext::intel::experimental::esimd::cos_emu
ESIMD_INLINE float cos_emu(T x0)
Definition: math.hpp:1521
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:1769
T
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:890
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::esimd::detail::is_saturation_tag_v
constexpr bool is_saturation_tag_v
Definition: common.hpp:76
sycl::_V1::ext::intel::esimd::inv
__ESIMD_API T inv(T src, Sat sat={})
Scalar version.
Definition: math.hpp:375
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:545
__SYCL_INLINE_VER_NAMESPACE
#define __SYCL_INLINE_VER_NAMESPACE(X)
Definition: defines_elementary.hpp:13
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::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:1740
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:1646
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:954
sycl::_V1::ext::intel::experimental::esimd::atan2_fast
ESIMD_INLINE float atan2_fast(float y, float x)
Definition: math.hpp:1350
sycl
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:13
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:788
sycl::_V1::ext::intel::experimental::esimd::imul
ESIMD_NODEBUG ESIMD_INLINE std::enable_if_t< __ESIMD_DNS::is_esimd_scalar< T >::value &&__ESIMD_DNS::is_esimd_scalar< U >::value &&__ESIMD_DNS::is_esimd_scalar< T0 >::value, T0 > imul(sycl::ext::intel::esimd::simd< T0, 1 > &rmd, T src0, U src1)
Definition: math.hpp:479
sycl::_V1::ext::intel::experimental::esimd::atan
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, T > atan(T src0)
Definition: math.hpp:1188
__SYCL_DEPRECATED
#define __SYCL_DEPRECATED(message)
Definition: defines_elementary.hpp:45
sycl::_V1::detail::enable_if_t
typename std::enable_if< B, T >::type enable_if_t
Definition: stl_type_traits.hpp:24
sycl::_V1::sign
detail::enable_if_t< detail::is_genfloat< T >::value, T > sign(T x) __NOEXC
Definition: builtins.hpp:622
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:1248
sycl::_V1::ext::intel::esimd::sqrt
__ESIMD_API T sqrt(T src, Sat sat={})
Scalar version.
Definition: math.hpp:389
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:1000
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:1662
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:1657
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:1146
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:1132
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::remainder
detail::enable_if_t< detail::is_genfloat< T >::value, T > remainder(T x, T y) __NOEXC
Definition: builtins.hpp:399
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:90
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:603
simd
Definition: simd.hpp:1027
sycl::_V1::ext::intel::experimental::esimd::fmod
ESIMD_INLINE float fmod(float y, float x)
Definition: math.hpp:1412
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:1376
std
Definition: accessor.hpp:3071
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:514
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:1046
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:1226
sycl::_V1::ext::intel::esimd::rsqrt
__ESIMD_API T rsqrt(T src, Sat sat={})
Scalar version.
Definition: math.hpp:397
sycl::_V1::ext::intel::experimental::esimd::sin_emu
ESIMD_INLINE float sin_emu(T x0)
Definition: math.hpp:1452
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:821
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:1077
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