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 
19 
20 namespace sycl {
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<__ESIMD_NS::uint, SZ> Src0 = src0;
992  return __esimd_lzd<__ESIMD_NS::uint, SZ>(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 __ESIMD_CONST_PI = 3.1415926535897932384626433832795;
1156 } // namespace detail
1158 
1159 template <typename T, int SZ>
1160 ESIMD_NODEBUG ESIMD_INLINE __ESIMD_NS::simd<T, SZ>
1161 atan(__ESIMD_NS::simd<T, SZ> src0) {
1162  static_assert(std::is_floating_point<T>::value,
1163  "Floating point argument type is expected.");
1164  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
1165 
1166  __ESIMD_NS::simd<T, SZ> OneP((T)1.0);
1167  __ESIMD_NS::simd<T, SZ> OneN((T)-1.0);
1168  __ESIMD_NS::simd<T, SZ> sign;
1169  __ESIMD_NS::simd_mask<SZ> Gt1 = Src0 > T(1.0);
1170 
1171  sign.merge(OneN, OneP, src0 < 0);
1172 
1173  Src0.merge(__ESIMD_NS::inv(Src0), Gt1);
1174 
1175  __ESIMD_NS::simd<T, SZ> Src0P2 = Src0 * Src0;
1176  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
1177 
1178  __ESIMD_NS::simd<T, SZ> Result =
1179  (Src0P4 * T(0.185696) + ((Src0 * T(0.787997) + T(0.63693)) * Src0P2) +
1180  Src0) /
1181  (((((Src0 * -T(0.000121387) + T(0.00202308)) * Src0P2) +
1182  (Src0 * -T(0.0149145)) + T(0.182569)) *
1183  Src0P4) +
1184  ((Src0 * T(0.395889) + T(1.12158)) * Src0P2) + (Src0 * T(0.636918)) +
1185  T(1.0));
1186 
1187  Result.merge(Result - T(detail::__ESIMD_CONST_PI) / T(2.0), Gt1);
1188 
1189  return __ESIMD_NS::abs(Result) * sign;
1190 }
1191 
1192 template <typename T> __ESIMD_API T atan(T src0) {
1193  static_assert(std::is_floating_point<T>::value,
1194  "Floating point argument type is expected.");
1195  __ESIMD_NS::simd<T, 1> Src0 = src0;
1196  __ESIMD_NS::simd<T, 1> Result = esimd::atan(Src0);
1197  return Result[0];
1198 }
1199 
1200 // acos
1201 
1202 template <typename T, int SZ>
1203 ESIMD_NODEBUG ESIMD_INLINE
1204  std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
1205  acos(__ESIMD_NS::simd<T, SZ> src0) {
1206  __ESIMD_NS::simd<T, SZ> Src0 = __ESIMD_NS::abs(src0);
1207 
1208  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
1209  __ESIMD_NS::simd_mask<SZ> TooBig = Src0 >= T(0.999998);
1210 
1211  // Replace oversized values to ensure no possibility of sqrt of
1212  // a negative value later
1213  Src0.merge(T(0.0), TooBig);
1214 
1215  __ESIMD_NS::simd<T, SZ> Src01m = T(1.0) - Src0;
1216 
1217  __ESIMD_NS::simd<T, SZ> Src0P2 = Src01m * Src01m;
1218  __ESIMD_NS::simd<T, SZ> Src0P4 = Src0P2 * Src0P2;
1219 
1220  __ESIMD_NS::simd<T, SZ> Result =
1221  (((Src01m * T(0.015098965761299077) - T(0.005516443930088506)) * Src0P4) +
1222  ((Src01m * T(0.047654245891495528) + T(0.163910606547823220)) * Src0P2) +
1223  Src01m * T(2.000291665285952400) - T(0.000007239283986332)) *
1224  __ESIMD_NS::rsqrt(Src01m * T(2.0));
1225 
1226  Result.merge(T(0.0), TooBig);
1227  Result.merge(T(detail::__ESIMD_CONST_PI) - Result, Neg);
1228  return Result;
1229 }
1230 
1231 template <typename T>
1232 __ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> acos(T src0) {
1233  __ESIMD_NS::simd<T, 1> Src0 = src0;
1234  __ESIMD_NS::simd<T, 1> Result = esimd::acos(Src0);
1235  return Result[0];
1236 }
1237 
1238 // asin
1239 
1240 template <typename T, int SZ>
1241 ESIMD_NODEBUG ESIMD_INLINE
1242  std::enable_if_t<std::is_floating_point<T>::value, __ESIMD_NS::simd<T, SZ>>
1243  asin(__ESIMD_NS::simd<T, SZ> src0) {
1244  __ESIMD_NS::simd_mask<SZ> Neg = src0 < T(0.0);
1245 
1246  __ESIMD_NS::simd<T, SZ> Result =
1247  T(detail::__ESIMD_CONST_PI / 2.0) - esimd::acos(__ESIMD_NS::abs(src0));
1248 
1249  Result.merge(-Result, Neg);
1250  return Result;
1251 }
1252 
1253 template <typename T>
1254 __ESIMD_API std::enable_if_t<std::is_floating_point<T>::value, T> asin(T src0) {
1255  __ESIMD_NS::simd<T, 1> Src0 = src0;
1256  __ESIMD_NS::simd<T, 1> Result = esimd::asin(Src0);
1257  return Result[0];
1258 }
1260 
1263 
1264 /* atan2_fast - a fast atan2 implementation */
1265 /* vector input */
1266 template <int N>
1267 __ESIMD_NS::simd<float, N> atan2_fast(__ESIMD_NS::simd<float, N> y,
1268  __ESIMD_NS::simd<float, N> x);
1269 /* scalar input */
1270 template <typename T> float atan2_fast(T y, T x);
1271 
1272 /* atan2 - atan2 implementation */
1273 /* For Vector input */
1274 template <int N>
1275 __ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
1276  __ESIMD_NS::simd<float, N> x);
1277 /* scalar Input */
1278 template <typename T> float atan2(T y, T x);
1279 
1280 /* fmod: */
1281 /* vector input */
1282 template <int N>
1283 __ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
1284  __ESIMD_NS::simd<float, N> x);
1285 /* scalar Input */
1286 template <typename T> float fmod(T y, T x);
1287 
1288 /* sin_emu - EU emulation for sin(x) */
1289 /* For Vector input */
1290 template <int N>
1291 __ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x);
1292 /* scalar Input */
1293 template <typename T> float sin_emu(T x);
1294 
1295 /* cos_emu - EU emulation for cos(x) */
1296 /* For Vector input */
1297 template <int N>
1298 __ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x);
1299 
1300 /* scalar Input */
1301 template <typename T> float cos_emu(T x);
1302 
1303 /* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
1304 /* float input */
1305 float tanh_cody_waite(float x);
1306 /* vector input */
1307 template <int N>
1308 __ESIMD_NS::simd<float, N> tanh_cody_waite(__ESIMD_NS::simd<float, N> x);
1309 /* tanh - opencl like implementation for tanh(x) */
1310 /* float input */
1311 float tanh(float x);
1312 /* vector input */
1313 template <int N> __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x);
1314 
1315 /* ------------------------- Extended Math Routines
1316  * -------------------------------------------------*/
1317 
1318 // For vector input
1319 template <int N>
1320 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1321 atan2_fast(__ESIMD_NS::simd<float, N> y, __ESIMD_NS::simd<float, N> x) {
1322  /* smallest such that 1.0+CONST_DBL_EPSILON != 1.0 */
1323  constexpr float CONST_DBL_EPSILON = 0.00001f;
1324  __ESIMD_NS::simd<float, N> OneP(1.0f);
1325  __ESIMD_NS::simd<float, N> OneN(-1.0f);
1326  __ESIMD_NS::simd<float, N> sign;
1327  __ESIMD_NS::simd<float, N> atan2;
1328  __ESIMD_NS::simd<float, N> r;
1329  __ESIMD_NS::simd_mask<N> mask = x < 0;
1330  __ESIMD_NS::simd<float, N> abs_y = __ESIMD_NS::abs(y) + CONST_DBL_EPSILON;
1331 
1332  r.merge((x + abs_y) / (abs_y - x), (x - abs_y) / (x + abs_y), mask);
1333  atan2.merge(float(detail::__ESIMD_CONST_PI) * 0.75f,
1334  float(detail::__ESIMD_CONST_PI) * 0.25f, mask);
1335  atan2 += (0.1963f * r * r - 0.9817f) * r;
1336 
1337  sign.merge(OneN, OneP, y < 0);
1338 
1339  return atan2 * sign;
1340 }
1341 
1342 // For Scalar Input
1343 template <> ESIMD_INLINE float atan2_fast(float y, float x) {
1344  __ESIMD_NS::simd<float, 1> vy = y;
1345  __ESIMD_NS::simd<float, 1> vx = x;
1346  __ESIMD_NS::simd<float, 1> atan2 = esimd::atan2_fast(vy, vx);
1347  return atan2[0];
1348 }
1349 
1350 // atan2
1351 // For Vector input
1352 template <int N>
1353 ESIMD_INLINE __ESIMD_NS::simd<float, N> atan2(__ESIMD_NS::simd<float, N> y,
1354  __ESIMD_NS::simd<float, N> x) {
1355  __ESIMD_NS::simd<float, N> v_distance;
1356  __ESIMD_NS::simd<float, N> atan2;
1357  __ESIMD_NS::simd_mask<N> mask;
1358 
1359  constexpr float CONST_DBL_EPSILON = 0.00001f;
1360 
1361  mask = (x < -CONST_DBL_EPSILON && y < CONST_DBL_EPSILON && y >= 0.f);
1362  atan2.merge(float(detail::__ESIMD_CONST_PI), 0.f, mask);
1363  mask = (x < -CONST_DBL_EPSILON && y > -CONST_DBL_EPSILON && y < 0);
1364  atan2.merge(float(-detail::__ESIMD_CONST_PI), mask);
1365  mask = (x < CONST_DBL_EPSILON && __ESIMD_NS::abs(y) > CONST_DBL_EPSILON);
1366  v_distance = __ESIMD_NS::sqrt(x * x + y * y);
1367  atan2.merge(2.0f * esimd::atan((v_distance - x) / y), mask);
1368 
1369  mask = (x > 0.f);
1370  atan2.merge(2.0f * esimd::atan(y / (v_distance + x)), mask);
1371 
1372  return atan2;
1373 }
1374 
1375 // For Scalar Input
1376 template <> ESIMD_INLINE float atan2(float y, float x) {
1377  __ESIMD_NS::simd<float, 1> vy = y;
1378  __ESIMD_NS::simd<float, 1> vx = x;
1379  __ESIMD_NS::simd<float, 1> atan2 = esimd::atan2(vy, vx);
1380  return atan2[0];
1381 }
1382 
1383 // fmod:
1384 // For Vector input
1385 template <int N>
1386 ESIMD_INLINE __ESIMD_NS::simd<float, N> fmod(__ESIMD_NS::simd<float, N> y,
1387  __ESIMD_NS::simd<float, N> x) {
1388  __ESIMD_NS::simd<float, N> abs_x = __ESIMD_NS::abs(x);
1389  __ESIMD_NS::simd<float, N> abs_y = __ESIMD_NS::abs(y);
1390 
1391  auto fmod_sign_mask = (y.template bit_cast_view<int32_t>()) & 0x80000000;
1392 
1393  __ESIMD_NS::simd<float, N> reminder =
1394  abs_y - abs_x * __ESIMD_NS::trunc<float>(abs_y / abs_x);
1395 
1396  abs_x.merge(0.0f, reminder >= 0);
1397  __ESIMD_NS::simd<float, N> fmod = reminder + abs_x;
1398  __ESIMD_NS::simd<float, N> fmod_abs = __ESIMD_NS::abs(fmod);
1399 
1400  auto fmod_bits =
1401  (fmod_abs.template bit_cast_view<int32_t>()) | fmod_sign_mask;
1402  return fmod_bits.template bit_cast_view<float>();
1403 }
1404 
1405 // For Scalar Input
1406 template <> ESIMD_INLINE float fmod(float y, float x) {
1407  return fmod(__ESIMD_NS::simd<float, 1>(y), __ESIMD_NS::simd<float, 1>(x))[0];
1408 }
1409 
1410 // sin_emu - EU emulation for sin(x)
1411 // For Vector input
1412 template <int N>
1413 ESIMD_INLINE __ESIMD_NS::simd<float, N> sin_emu(__ESIMD_NS::simd<float, N> x) {
1414  __ESIMD_NS::simd<float, N> x1;
1415  __ESIMD_NS::simd<float, N> x2;
1416  __ESIMD_NS::simd<float, N> t3;
1417 
1418  __ESIMD_NS::simd<float, N> sign;
1419  __ESIMD_NS::simd<float, N> fTrig;
1420  __ESIMD_NS::simd<float, N> TwoPI(float(detail::__ESIMD_CONST_PI) * 2.0f);
1421  __ESIMD_NS::simd<float, N> CmpI((float)detail::__ESIMD_CONST_PI);
1422  __ESIMD_NS::simd<float, N> OneP(1.0f);
1423  __ESIMD_NS::simd<float, N> OneN(-1.0f);
1424 
1425  x = esimd::fmod(x, TwoPI);
1426  x.merge(TwoPI + x, x < 0);
1427 
1428  x1.merge(CmpI - x, x - CmpI, (x <= float(detail::__ESIMD_CONST_PI)));
1429  x1.merge(x, (x <= float(detail::__ESIMD_CONST_PI) * 0.5f));
1430  x1.merge(TwoPI - x, (x > float(detail::__ESIMD_CONST_PI) * 1.5f));
1431 
1432  sign.merge(OneN, OneP, (x > float(detail::__ESIMD_CONST_PI)));
1433 
1434  x2 = x1 * x1;
1435  t3 = x2 * x1 * 0.1666667f;
1436 
1437  fTrig =
1438  x1 + t3 * (OneN + x2 * 0.05f *
1439  (OneP + x2 * 0.0238095f *
1440  (OneN + x2 * 0.0138889f *
1441  (OneP - x2 * 0.0090909f))));
1442  fTrig *= sign;
1443  return fTrig;
1444 }
1445 
1446 // scalar Input
1447 template <> ESIMD_INLINE float sin_emu(float x0) {
1448  return esimd::sin_emu(__ESIMD_NS::simd<float, 1>(x0))[0];
1449 }
1450 
1451 // cos_emu - EU emulation for sin(x)
1452 // For Vector input
1453 template <int N>
1454 ESIMD_INLINE __ESIMD_NS::simd<float, N> cos_emu(__ESIMD_NS::simd<float, N> x) {
1455  return esimd::sin_emu(0.5f * float(detail::__ESIMD_CONST_PI) - x);
1456 }
1457 
1458 // scalar Input
1459 template <> ESIMD_INLINE float cos_emu(float x0) {
1460  return esimd::cos_emu(__ESIMD_NS::simd<float, 1>(x0))[0];
1461 }
1462 
1464 namespace detail {
1465 
1466 template <int N>
1467 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1468 tanh_cody_waite_impl(__ESIMD_NS::simd<float, N> x) {
1469  /*
1470  * 0 x_small x_medium x_large
1471  * | x | rational polynomial | 1 - 2/(1 + exp(2*x)) | 1
1472  *
1473  * rational polynomial for single precision = x + x * (g * (p[1] * g + p[0]) /
1474  * (g + q[0]) g = x^2 p0 = -0.82377 28127 E+00 p1 = -0.38310 10665 E-02 q0 =
1475  * 0.24713 19654 E+01 q1 = 1.00000 00000 E+00
1476  *
1477  */
1478 
1479  constexpr float p0 = -0.8237728127E+00f;
1480  constexpr float p1 = -0.3831010665E-02f;
1481  constexpr float q0 = 0.2471319654E+01f;
1482  constexpr float q1 = 1.0000000000E+00f;
1483  constexpr float xsmall = 4.22863966691620432990E-04f;
1484  constexpr float xmedium = 0.54930614433405484570f;
1485  constexpr float xlarge = 8.66433975699931636772f;
1486 
1487  using RT = __ESIMD_NS::simd<float, N>;
1488 
1489  RT absX = __ESIMD_NS::abs(x);
1490  RT g = absX * absX;
1491 
1492  RT sign;
1493  sign.merge(-1.f, 1.f, x < 0.f);
1494 
1495  auto isLarge = absX > xlarge;
1496  auto minor = absX <= xlarge;
1497  auto isGtMed = minor & (absX > xmedium);
1498  auto isGtSmall = (absX > xsmall) & (absX <= xmedium);
1499 
1500  RT res;
1501  res.merge(sign, x, isLarge);
1502  auto temp = __ESIMD_NS::exp(absX * 2.0f) + 1.f;
1503  temp = ((temp - 2.f) / temp) * sign;
1504  res.merge(temp, isGtMed);
1505  res.merge((absX + absX * g * (g * p1 + p0) / (g + q0)) * sign, isGtSmall);
1506 
1507  return res;
1508 }
1509 
1510 template <int N>
1511 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1512 tanh_impl(__ESIMD_NS::simd<float, N> x) {
1513  /*
1514  * 0 x_small x_large
1515  * | x | ( exp(x) - exp(-x) ) / ( exp(x) + exp(-x) ) | 1
1516  *
1517  */
1518 
1519  constexpr float xsmall = 0.000045f; // same as exp(-10.0f)
1520  constexpr float xlarge = 40.f;
1521 
1522  using RT = __ESIMD_NS::simd<float, N>;
1523 
1524  RT absX = __ESIMD_NS::abs(x);
1525 
1526  RT sign;
1527  sign.merge(-1.f, 1.f, x < 0.f);
1528 
1529  auto isLarge = (absX > xlarge);
1530  auto isLessE = (absX <= xlarge);
1531 
1532  RT res;
1533  res.merge(sign, x, isLarge);
1534 
1535  RT exp;
1536  exp = __ESIMD_NS::exp(absX * 2.f);
1537 
1538  res.merge(((exp - 1.f) / (exp + 1.f)) * sign, (absX > xsmall) & isLessE);
1539 
1540  return res;
1541 }
1542 } // namespace detail
1544 
1545 /* tanh_cody_waite - Cody-Waite implementation for tanh(x) */
1546 /* float input */
1547 ESIMD_INLINE float tanh_cody_waite(float x) {
1548  return detail::tanh_cody_waite_impl(__ESIMD_NS::simd<float, 1>(x))[0];
1549 }
1550 /* vector input */
1551 template <int N>
1552 ESIMD_INLINE __ESIMD_NS::simd<float, N>
1553 tanh_cody_waite(__ESIMD_NS::simd<float, N> x) {
1554  return detail::tanh_cody_waite_impl(x);
1555 }
1556 
1557 /* tanh - opencl like implementation for tanh(x) */
1558 /* float input */
1559 ESIMD_INLINE float tanh(float x) {
1560  return esimd::detail::tanh_impl(__ESIMD_NS::simd<float, 1>(x))[0];
1561 }
1562 /* vector input */
1563 template <int N>
1564 ESIMD_INLINE __ESIMD_NS::simd<float, N> tanh(__ESIMD_NS::simd<float, N> x) {
1565  return esimd::detail::tanh_impl(x);
1566 }
1567 
1568 template <typename T, int N>
1569 __ESIMD_NS::simd<T, N> dp4(__ESIMD_NS::simd<T, N> v1,
1570  __ESIMD_NS::simd<T, N> v2) {
1571  auto retv = __esimd_dp4<T, N>(v1.data(), v2.data());
1572  return retv;
1573 }
1574 
1576 
1580 
1593 template <argument_type src1_precision, argument_type src2_precision,
1594  typename T, int systolic_depth, int repeat_count, typename T0,
1595  typename T1, typename T2, int N, int N1, int N2,
1596  typename Sat = __ESIMD_NS::saturation_off_tag>
1597 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::native::dpas()")
1598 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1599  __ESIMD_NS::simd<T0, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1600  __ESIMD_NS::simd<T2, N2> src2,
1601  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1602  auto result =
1603  __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T0, T1, T2,
1604  src1_precision, src2_precision>(src0, src1, src2);
1605  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1606  return result;
1607  else
1608  return __ESIMD_NS::saturate<T>(result);
1609 }
1610 
1621 template <argument_type src1_precision, argument_type src2_precision,
1622  int systolic_depth, int repeat_count, typename T, typename T1,
1623  typename T2, int N, int N1, int N2,
1624  typename Sat = __ESIMD_NS::saturation_off_tag>
1625 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
1626 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1627  __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1628  __ESIMD_NS::simd<T2, N2> src2,
1629  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1630  return dpas<src1_precision, src2_precision, T, systolic_depth, repeat_count>(
1631  src0, src1, src2, sat);
1632 }
1633 
1642 template <argument_type src1_precision, argument_type src2_precision,
1643  int systolic_depth, int repeat_count, typename T, typename T1,
1644  typename T2, int N, int N1, int N2,
1645  typename Sat = __ESIMD_NS::saturation_off_tag>
1646 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpas()")
1647 __ESIMD_API __ESIMD_NS::simd<T, N> dpas(
1648  __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
1649  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1650 
1651  __ESIMD_NS::simd<T, N> result =
1652  __ESIMD_NS::xmx::dpas<systolic_depth, repeat_count, T, T1, T2,
1653  src1_precision, src2_precision>(src1, src2);
1654 
1655  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1656  return result;
1657  else
1658  return __ESIMD_NS::saturate<T>(result);
1659 }
1660 
1671 template <argument_type src1_precision, argument_type src2_precision,
1672  int systolic_depth, int repeat_count, typename T, typename T1,
1673  typename T2, int N, int N1, int N2,
1674  typename Sat = __ESIMD_NS::saturation_off_tag>
1675 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
1676 __ESIMD_API __ESIMD_NS::simd<T, N> dpasw(
1677  __ESIMD_NS::simd<T, N> src0, __ESIMD_NS::simd<T1, N1> src1,
1678  __ESIMD_NS::simd<T2, N2> src2,
1679  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1680 
1681  __ESIMD_NS::simd<T, N> result =
1682  __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
1683  src1_precision, src2_precision>(src0, src1, src2);
1684  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1685  return result;
1686  else
1687  return __ESIMD_NS::saturate<T>(result);
1688 }
1689 
1698 template <argument_type src1_precision, argument_type src2_precision,
1699  int systolic_depth, int repeat_count, typename T, typename T1,
1700  typename T2, int N, int N1, int N2,
1701  typename Sat = __ESIMD_NS::saturation_off_tag>
1702 __SYCL_DEPRECATED("use sycl::ext::intel::esimd::xmx::dpasw()")
1703 __ESIMD_API __ESIMD_NS::simd<T, N> dpasw2(
1704  __ESIMD_NS::simd<T1, N1> src1, __ESIMD_NS::simd<T2, N2> src2,
1705  std::enable_if_t<__ESIMD_DNS::is_saturation_tag_v<Sat>, Sat> sat = {}) {
1706 
1707  __ESIMD_NS::simd<T, N> result =
1708  __ESIMD_NS::xmx::dpasw<systolic_depth, repeat_count, T, T1, T2,
1709  src1_precision, src2_precision>(src1, src2);
1710 
1711  if constexpr (std::is_same_v<Sat, __ESIMD_NS::saturation_off_tag>)
1712  return result;
1713  else
1714  return __ESIMD_NS::saturate<T>(result);
1715 }
1717 
1718 } // namespace ext::intel::experimental::esimd
1719 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
1720 } // namespace sycl
Definition: simd.hpp:1384
#define __SYCL_INLINE_VER_NAMESPACE(X)
#define __SYCL_DEPRECATED(message)
__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
__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
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
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
__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
__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
__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
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
__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
__ESIMD_API T cos(T src, Sat sat={})
Scalar version.
Definition: math.hpp:405
__ESIMD_API T inv(T src, Sat sat={})
Scalar version.
Definition: math.hpp:375
__ESIMD_API T sin(T src, Sat sat={})
Scalar version.
Definition: math.hpp:401
__ESIMD_API T frc(T src0)
Performs truncate-to-minus-infinity fraction operation of src0.
Definition: math.hpp:979
ESIMD_INLINE float sin_emu(float x0)
Definition: math.hpp:1447
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:1569
ESIMD_INLINE float fmod(float y, float x)
Definition: math.hpp:1406
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N > tanh(sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1564
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
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
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
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
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
ESIMD_INLINE float atan2(float y, float x)
Definition: math.hpp:1376
__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
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, T > asin(T src0)
Definition: math.hpp:1254
__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
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
ESIMD_INLINE float atan2_fast(float y, float x)
Definition: math.hpp:1343
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
ESIMD_INLINE float cos_emu(float x0)
Definition: math.hpp:1459
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
__ESIMD_API std::enable_if_t< std::is_floating_point< T >::value, T > acos(T src0)
Definition: math.hpp:1232
ESIMD_NODEBUG ESIMD_INLINE T exp(T src0, Sat sat={})
Definition: math.hpp:497
__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
__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
ESIMD_INLINE sycl::ext::intel::esimd::simd< float, N > tanh_cody_waite(sycl::ext::intel::esimd::simd< float, N > x)
Definition: math.hpp:1553
__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:1676
__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:1703
__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:1647
conditional< sizeof(long)==8, long, long long >::type int64_t
Definition: kernel_desc.hpp:34
typename std::enable_if< B, T >::type enable_if_t
sycl::half2 rsqrt(sycl::half2 x)
Definition: math.hpp:177
sycl::half2 sqrt(sycl::half2 x)
Definition: math.hpp:157
std::experimental::simd< T, simd_abi::native_fixed_size< T, N > > simd
Definition: invoke_simd.hpp:88
---— Error handling, matching OpenCL plugin semantics.
Definition: access.hpp:14