$darkmode
Eigen  5.0.1-dev
MathFunctions.h
1 // This file is part of Eigen, a lightweight C++ template library
2 // for linear algebra.
3 //
4 // Copyright (C) 2006-2010 Benoit Jacob <jacob.benoit.1@gmail.com>
5 // Copyright (c) 2021, NVIDIA CORPORATION. All rights reserved.
6 //
7 // This Source Code Form is subject to the terms of the Mozilla
8 // Public License v. 2.0. If a copy of the MPL was not distributed
9 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
10 
11 #ifndef EIGEN_MATHFUNCTIONS_H
12 #define EIGEN_MATHFUNCTIONS_H
13 
14 // TODO this should better be moved to NumTraits
15 // Source: WolframAlpha
16 #define EIGEN_PI 3.141592653589793238462643383279502884197169399375105820974944592307816406L
17 #define EIGEN_LOG2E 1.442695040888963407359924681001892137426645954152985934135449406931109219L
18 #define EIGEN_LN2 0.693147180559945309417232121458176568075500134360255254120680009493393621L
19 
20 // IWYU pragma: private
21 #include "./InternalHeaderCheck.h"
22 
23 namespace Eigen {
24 
25 namespace internal {
26 
48 template <typename T, typename dummy = void>
49 struct global_math_functions_filtering_base {
50  typedef T type;
51 };
52 
53 template <typename T>
54 struct always_void {
55  typedef void type;
56 };
57 
58 template <typename T>
59 struct global_math_functions_filtering_base<
60  T, typename always_void<typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl>::type> {
61  typedef typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl type;
62 };
63 
64 #define EIGEN_MATHFUNC_IMPL(func, scalar) \
65  Eigen::internal::func##_impl<typename Eigen::internal::global_math_functions_filtering_base<scalar>::type>
66 #define EIGEN_MATHFUNC_RETVAL(func, scalar) \
67  typename Eigen::internal::func##_retval< \
68  typename Eigen::internal::global_math_functions_filtering_base<scalar>::type>::type
69 
70 /****************************************************************************
71  * Implementation of real *
72  ****************************************************************************/
73 
74 template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
75 struct real_default_impl {
76  typedef typename NumTraits<Scalar>::Real RealScalar;
77  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x; }
78 };
79 
80 template <typename Scalar>
81 struct real_default_impl<Scalar, true> {
82  typedef typename NumTraits<Scalar>::Real RealScalar;
83  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
84  using std::real;
85  return real(x);
86  }
87 };
88 
89 template <typename Scalar>
90 struct real_impl : real_default_impl<Scalar> {};
91 
92 #if defined(EIGEN_GPU_COMPILE_PHASE)
93 template <typename T>
94 struct real_impl<std::complex<T>> {
95  typedef T RealScalar;
96  EIGEN_DEVICE_FUNC static inline T run(const std::complex<T>& x) { return x.real(); }
97 };
98 #endif
99 
100 template <typename Scalar>
101 struct real_retval {
102  typedef typename NumTraits<Scalar>::Real type;
103 };
104 
105 /****************************************************************************
106  * Implementation of imag *
107  ****************************************************************************/
108 
109 template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
110 struct imag_default_impl {
111  typedef typename NumTraits<Scalar>::Real RealScalar;
112  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar&) { return RealScalar(0); }
113 };
114 
115 template <typename Scalar>
116 struct imag_default_impl<Scalar, true> {
117  typedef typename NumTraits<Scalar>::Real RealScalar;
118  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
119  using std::imag;
120  return imag(x);
121  }
122 };
123 
124 template <typename Scalar>
125 struct imag_impl : imag_default_impl<Scalar> {};
126 
127 #if defined(EIGEN_GPU_COMPILE_PHASE)
128 template <typename T>
129 struct imag_impl<std::complex<T>> {
130  typedef T RealScalar;
131  EIGEN_DEVICE_FUNC static inline T run(const std::complex<T>& x) { return x.imag(); }
132 };
133 #endif
134 
135 template <typename Scalar>
136 struct imag_retval {
137  typedef typename NumTraits<Scalar>::Real type;
138 };
139 
140 /****************************************************************************
141  * Implementation of real_ref *
142  ****************************************************************************/
143 
144 template <typename Scalar>
145 struct real_ref_impl {
146  typedef typename NumTraits<Scalar>::Real RealScalar;
147  EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast<RealScalar*>(&x)[0]; }
148  EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) {
149  return reinterpret_cast<const RealScalar*>(&x)[0];
150  }
151 };
152 
153 template <typename Scalar>
154 struct real_ref_retval {
155  typedef typename NumTraits<Scalar>::Real& type;
156 };
157 
158 /****************************************************************************
159  * Implementation of imag_ref *
160  ****************************************************************************/
161 
162 template <typename Scalar, bool IsComplex>
163 struct imag_ref_default_impl {
164  typedef typename NumTraits<Scalar>::Real RealScalar;
165  EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast<RealScalar*>(&x)[1]; }
166  EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) {
167  return reinterpret_cast<const RealScalar*>(&x)[1];
168  }
169 };
170 
171 template <typename Scalar>
172 struct imag_ref_default_impl<Scalar, false> {
173  EIGEN_DEVICE_FUNC constexpr static Scalar run(Scalar&) { return Scalar(0); }
174  EIGEN_DEVICE_FUNC constexpr static const Scalar run(const Scalar&) { return Scalar(0); }
175 };
176 
177 template <typename Scalar>
178 struct imag_ref_impl : imag_ref_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
179 
180 template <typename Scalar>
181 struct imag_ref_retval {
182  typedef typename NumTraits<Scalar>::Real& type;
183 };
184 
185 } // namespace internal
186 
187 namespace numext {
188 
189 template <typename Scalar>
190 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x) {
191  return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x);
192 }
193 
194 template <typename Scalar>
195 EIGEN_DEVICE_FUNC inline internal::add_const_on_value_type_t<EIGEN_MATHFUNC_RETVAL(real_ref, Scalar)> real_ref(
196  const Scalar& x) {
197  return internal::real_ref_impl<Scalar>::run(x);
198 }
199 
200 template <typename Scalar>
201 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x) {
202  return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x);
203 }
204 
205 template <typename Scalar>
206 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x) {
207  return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x);
208 }
209 
210 template <typename Scalar>
211 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar select(const Scalar& mask, const Scalar& a, const Scalar& b) {
212  return numext::is_exactly_zero(mask) ? b : a;
213 }
214 
215 } // namespace numext
216 
217 namespace internal {
218 
219 /****************************************************************************
220  * Implementation of conj *
221  ****************************************************************************/
222 
223 template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
224 struct conj_default_impl {
225  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { return x; }
226 };
227 
228 template <typename Scalar>
229 struct conj_default_impl<Scalar, true> {
230  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
231  using std::conj;
232  return conj(x);
233  }
234 };
235 
236 template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
237 struct conj_impl : conj_default_impl<Scalar, IsComplex> {};
238 
239 template <typename Scalar>
240 struct conj_retval {
241  typedef Scalar type;
242 };
243 
244 /****************************************************************************
245  * Implementation of abs2 *
246  ****************************************************************************/
247 
248 template <typename Scalar, bool IsComplex>
249 struct abs2_impl_default {
250  typedef typename NumTraits<Scalar>::Real RealScalar;
251  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x * x; }
252 };
253 
254 template <typename Scalar>
255 struct abs2_impl_default<Scalar, true> // IsComplex
256 {
257  typedef typename NumTraits<Scalar>::Real RealScalar;
258  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
259  return numext::real(x) * numext::real(x) + numext::imag(x) * numext::imag(x);
260  }
261 };
262 
263 template <typename Scalar>
264 struct abs2_impl {
265  typedef typename NumTraits<Scalar>::Real RealScalar;
266  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
267  return abs2_impl_default<Scalar, NumTraits<Scalar>::IsComplex>::run(x);
268  }
269 };
270 
271 template <typename Scalar>
272 struct abs2_retval {
273  typedef typename NumTraits<Scalar>::Real type;
274 };
275 
276 /****************************************************************************
277  * Implementation of sqrt/rsqrt *
278  ****************************************************************************/
279 
280 template <typename Scalar>
281 struct sqrt_impl {
282  EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE Scalar run(const Scalar& x) {
283  EIGEN_USING_STD(sqrt);
284  return sqrt(x);
285  }
286 };
287 
288 // Complex sqrt defined in MathFunctionsImpl.h.
289 template <typename ComplexT>
290 EIGEN_DEVICE_FUNC ComplexT complex_sqrt(const ComplexT& a_x);
291 
292 // Custom implementation is faster than `std::sqrt`, works on
293 // GPU, and correctly handles special cases (unlike MSVC).
294 template <typename T>
295 struct sqrt_impl<std::complex<T>> {
296  EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x) { return complex_sqrt(x); }
297 };
298 
299 template <typename Scalar>
300 struct sqrt_retval {
301  typedef Scalar type;
302 };
303 
304 // Default implementation relies on numext::sqrt, at bottom of file.
305 template <typename T>
306 struct rsqrt_impl;
307 
308 // Complex rsqrt defined in MathFunctionsImpl.h.
309 template <typename ComplexT>
310 EIGEN_DEVICE_FUNC ComplexT complex_rsqrt(const ComplexT& a_x);
311 
312 template <typename T>
313 struct rsqrt_impl<std::complex<T>> {
314  EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x) {
315  return complex_rsqrt(x);
316  }
317 };
318 
319 template <typename Scalar>
320 struct rsqrt_retval {
321  typedef Scalar type;
322 };
323 
324 /****************************************************************************
325  * Implementation of norm1 *
326  ****************************************************************************/
327 
328 template <typename Scalar, bool IsComplex>
329 struct norm1_default_impl;
330 
331 template <typename Scalar>
332 struct norm1_default_impl<Scalar, true> {
333  typedef typename NumTraits<Scalar>::Real RealScalar;
334  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
335  EIGEN_USING_STD(abs);
336  return abs(numext::real(x)) + abs(numext::imag(x));
337  }
338 };
339 
340 template <typename Scalar>
341 struct norm1_default_impl<Scalar, false> {
342  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
343  EIGEN_USING_STD(abs);
344  return abs(x);
345  }
346 };
347 
348 template <typename Scalar>
349 struct norm1_impl : norm1_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
350 
351 template <typename Scalar>
352 struct norm1_retval {
353  typedef typename NumTraits<Scalar>::Real type;
354 };
355 
356 /****************************************************************************
357  * Implementation of hypot *
358  ****************************************************************************/
359 
360 template <typename Scalar>
361 struct hypot_impl;
362 
363 template <typename Scalar>
364 struct hypot_retval {
365  typedef typename NumTraits<Scalar>::Real type;
366 };
367 
368 /****************************************************************************
369  * Implementation of cast *
370  ****************************************************************************/
371 
372 template <typename OldType, typename NewType, typename EnableIf = void>
373 struct cast_impl {
374  EIGEN_DEVICE_FUNC static inline NewType run(const OldType& x) { return static_cast<NewType>(x); }
375 };
376 
377 template <typename OldType>
378 struct cast_impl<OldType, bool> {
379  EIGEN_DEVICE_FUNC static inline bool run(const OldType& x) { return x != OldType(0); }
380 };
381 
382 // Casting from S -> Complex<T> leads to an implicit conversion from S to T,
383 // generating warnings on clang. Here we explicitly cast the real component.
384 template <typename OldType, typename NewType>
385 struct cast_impl<OldType, NewType,
386  typename std::enable_if_t<!NumTraits<OldType>::IsComplex && NumTraits<NewType>::IsComplex>> {
387  EIGEN_DEVICE_FUNC static inline NewType run(const OldType& x) {
388  typedef typename NumTraits<NewType>::Real NewReal;
389  return static_cast<NewType>(static_cast<NewReal>(x));
390  }
391 };
392 
393 // here, for once, we're plainly returning NewType: we don't want cast to do weird things.
394 
395 template <typename OldType, typename NewType>
396 EIGEN_DEVICE_FUNC inline NewType cast(const OldType& x) {
397  return cast_impl<OldType, NewType>::run(x);
398 }
399 
400 /****************************************************************************
401  * Implementation of arg *
402  ****************************************************************************/
403 
404 // Visual Studio 2017 has a bug where arg(float) returns 0 for negative inputs.
405 // This seems to be fixed in VS 2019.
406 #if (!EIGEN_COMP_MSVC || EIGEN_COMP_MSVC >= 1920)
407 // std::arg is only defined for types of std::complex, or integer types or float/double/long double
408 template <typename Scalar, bool HasStdImpl = NumTraits<Scalar>::IsComplex || is_integral<Scalar>::value ||
409  is_same<Scalar, float>::value || is_same<Scalar, double>::value ||
410  is_same<Scalar, long double>::value>
411 struct arg_default_impl;
412 
413 template <typename Scalar>
414 struct arg_default_impl<Scalar, true> {
415  typedef typename NumTraits<Scalar>::Real RealScalar;
416  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
417  // There is no official ::arg on device in CUDA/HIP, so we always need to use std::arg.
418  using std::arg;
419  return static_cast<RealScalar>(arg(x));
420  }
421 };
422 
423 // Must be non-complex floating-point type (e.g. half/bfloat16).
424 template <typename Scalar>
425 struct arg_default_impl<Scalar, false> {
426  typedef typename NumTraits<Scalar>::Real RealScalar;
427  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
428  return (x < Scalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
429  }
430 };
431 #else
432 template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
433 struct arg_default_impl {
434  typedef typename NumTraits<Scalar>::Real RealScalar;
435  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
436  return (x < RealScalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
437  }
438 };
439 
440 template <typename Scalar>
441 struct arg_default_impl<Scalar, true> {
442  typedef typename NumTraits<Scalar>::Real RealScalar;
443  EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
444  EIGEN_USING_STD(arg);
445  return arg(x);
446  }
447 };
448 #endif
449 template <typename Scalar>
450 struct arg_impl : arg_default_impl<Scalar> {};
451 
452 template <typename Scalar>
453 struct arg_retval {
454  typedef typename NumTraits<Scalar>::Real type;
455 };
456 
457 /****************************************************************************
458  * Implementation of expm1 *
459  ****************************************************************************/
460 
461 // This implementation is based on GSL Math's expm1.
462 namespace std_fallback {
463 // fallback expm1 implementation in case there is no expm1(Scalar) function in namespace of Scalar,
464 // or that there is no suitable std::expm1 function available. Implementation
465 // attributed to Kahan. See: http://www.plunk.org/~hatch/rightway.php.
466 template <typename Scalar>
467 EIGEN_DEVICE_FUNC inline Scalar expm1(const Scalar& x) {
468  EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
469  typedef typename NumTraits<Scalar>::Real RealScalar;
470 
471  EIGEN_USING_STD(exp);
472  Scalar u = exp(x);
473  if (numext::equal_strict(u, Scalar(1))) {
474  return x;
475  }
476  Scalar um1 = u - RealScalar(1);
477  if (numext::equal_strict(um1, Scalar(-1))) {
478  return RealScalar(-1);
479  }
480 
481  EIGEN_USING_STD(log);
482  Scalar logu = log(u);
483  return numext::equal_strict(u, logu) ? u : (u - RealScalar(1)) * x / logu;
484 }
485 } // namespace std_fallback
486 
487 template <typename Scalar>
488 struct expm1_impl {
489  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
490  EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
491  EIGEN_USING_STD(expm1);
492  return expm1(x);
493  }
494 };
495 
496 template <typename Scalar>
497 struct expm1_retval {
498  typedef Scalar type;
499 };
500 
501 /****************************************************************************
502  * Implementation of log *
503  ****************************************************************************/
504 
505 // Complex log defined in MathFunctionsImpl.h.
506 template <typename ComplexT>
507 EIGEN_DEVICE_FUNC ComplexT complex_log(const ComplexT& z);
508 
509 template <typename Scalar>
510 struct log_impl {
511  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
512  EIGEN_USING_STD(log);
513  return static_cast<Scalar>(log(x));
514  }
515 };
516 
517 template <typename Scalar>
518 struct log_impl<std::complex<Scalar>> {
519  EIGEN_DEVICE_FUNC static inline std::complex<Scalar> run(const std::complex<Scalar>& z) { return complex_log(z); }
520 };
521 
522 /****************************************************************************
523  * Implementation of log1p *
524  ****************************************************************************/
525 
526 namespace std_fallback {
527 // fallback log1p implementation in case there is no log1p(Scalar) function in namespace of Scalar,
528 // or that there is no suitable std::log1p function available
529 template <typename Scalar>
530 EIGEN_DEVICE_FUNC inline Scalar log1p(const Scalar& x) {
531  EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
532  typedef typename NumTraits<Scalar>::Real RealScalar;
533  EIGEN_USING_STD(log);
534  Scalar x1p = RealScalar(1) + x;
535  Scalar log_1p = log_impl<Scalar>::run(x1p);
536  const bool is_small = numext::equal_strict(x1p, Scalar(1));
537  const bool is_inf = numext::equal_strict(x1p, log_1p);
538  return (is_small || is_inf) ? x : x * (log_1p / (x1p - RealScalar(1)));
539 }
540 } // namespace std_fallback
541 
542 template <typename Scalar>
543 struct log1p_impl {
544  EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
545 
546  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
547  EIGEN_USING_STD(log1p);
548  return log1p(x);
549  }
550 };
551 
552 // Specialization for complex types that are not supported by std::log1p.
553 template <typename RealScalar>
554 struct log1p_impl<std::complex<RealScalar>> {
555  EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
556 
557  EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(const std::complex<RealScalar>& x) {
558  return std_fallback::log1p(x);
559  }
560 };
561 
562 template <typename Scalar>
563 struct log1p_retval {
564  typedef Scalar type;
565 };
566 
567 /****************************************************************************
568  * Implementation of pow *
569  ****************************************************************************/
570 
571 template <typename ScalarX, typename ScalarY,
572  bool IsInteger = NumTraits<ScalarX>::IsInteger && NumTraits<ScalarY>::IsInteger>
573 struct pow_impl {
574  // typedef Scalar retval;
575  typedef typename ScalarBinaryOpTraits<ScalarX, ScalarY, internal::scalar_pow_op<ScalarX, ScalarY>>::ReturnType
576  result_type;
577  static EIGEN_DEVICE_FUNC inline result_type run(const ScalarX& x, const ScalarY& y) {
578  EIGEN_USING_STD(pow);
579  return pow(x, y);
580  }
581 };
582 
583 template <typename ScalarX, typename ScalarY>
584 struct pow_impl<ScalarX, ScalarY, true> {
585  typedef ScalarX result_type;
586  static EIGEN_DEVICE_FUNC inline ScalarX run(ScalarX x, ScalarY y) {
587  ScalarX res(1);
588  eigen_assert(!NumTraits<ScalarY>::IsSigned || y >= 0);
589  if (y & 1) res *= x;
590  y >>= 1;
591  while (y) {
592  x *= x;
593  if (y & 1) res *= x;
594  y >>= 1;
595  }
596  return res;
597  }
598 };
599 
600 enum { meta_floor_log2_terminate, meta_floor_log2_move_up, meta_floor_log2_move_down, meta_floor_log2_bogus };
601 
602 template <unsigned int n, int lower, int upper>
603 struct meta_floor_log2_selector {
604  enum {
605  middle = (lower + upper) / 2,
606  value = (upper <= lower + 1) ? int(meta_floor_log2_terminate)
607  : (n < (1 << middle)) ? int(meta_floor_log2_move_down)
608  : (n == 0) ? int(meta_floor_log2_bogus)
609  : int(meta_floor_log2_move_up)
610  };
611 };
612 
613 template <unsigned int n, int lower = 0, int upper = sizeof(unsigned int) * CHAR_BIT - 1,
614  int selector = meta_floor_log2_selector<n, lower, upper>::value>
615 struct meta_floor_log2 {};
616 
617 template <unsigned int n, int lower, int upper>
618 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_down> {
619  enum { value = meta_floor_log2<n, lower, meta_floor_log2_selector<n, lower, upper>::middle>::value };
620 };
621 
622 template <unsigned int n, int lower, int upper>
623 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_up> {
624  enum { value = meta_floor_log2<n, meta_floor_log2_selector<n, lower, upper>::middle, upper>::value };
625 };
626 
627 template <unsigned int n, int lower, int upper>
628 struct meta_floor_log2<n, lower, upper, meta_floor_log2_terminate> {
629  enum { value = (n >= ((unsigned int)(1) << (lower + 1))) ? lower + 1 : lower };
630 };
631 
632 template <unsigned int n, int lower, int upper>
633 struct meta_floor_log2<n, lower, upper, meta_floor_log2_bogus> {
634  // no value, error at compile time
635 };
636 
637 template <typename BitsType, typename EnableIf = void>
638 struct count_bits_impl {
639  static_assert(std::is_integral<BitsType>::value && std::is_unsigned<BitsType>::value,
640  "BitsType must be an unsigned integer");
641  static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
642  int n = CHAR_BIT * sizeof(BitsType);
643  int shift = n / 2;
644  while (bits > 0 && shift > 0) {
645  BitsType y = bits >> shift;
646  if (y > 0) {
647  n -= shift;
648  bits = y;
649  }
650  shift /= 2;
651  }
652  if (shift == 0) {
653  --n;
654  }
655  return n;
656  }
657 
658  static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
659  int n = CHAR_BIT * sizeof(BitsType);
660  int shift = n / 2;
661  while (bits > 0 && shift > 0) {
662  BitsType y = bits << shift;
663  if (y > 0) {
664  n -= shift;
665  bits = y;
666  }
667  shift /= 2;
668  }
669  if (shift == 0) {
670  --n;
671  }
672  return n;
673  }
674 };
675 
676 // Count leading zeros.
677 template <typename BitsType>
678 EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
679  return count_bits_impl<BitsType>::clz(bits);
680 }
681 
682 // Count trailing zeros.
683 template <typename BitsType>
684 EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
685  return count_bits_impl<BitsType>::ctz(bits);
686 }
687 
688 #if EIGEN_COMP_GNUC || EIGEN_COMP_CLANG
689 
690 template <typename BitsType>
691 struct count_bits_impl<
692  BitsType, std::enable_if_t<std::is_integral<BitsType>::value && sizeof(BitsType) <= sizeof(unsigned int)>> {
693  static constexpr int kNumBits = static_cast<int>(sizeof(BitsType) * CHAR_BIT);
694  static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
695  static constexpr int kLeadingBitsOffset = (sizeof(unsigned int) - sizeof(BitsType)) * CHAR_BIT;
696  return bits == 0 ? kNumBits : __builtin_clz(static_cast<unsigned int>(bits)) - kLeadingBitsOffset;
697  }
698 
699  static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
700  return bits == 0 ? kNumBits : __builtin_ctz(static_cast<unsigned int>(bits));
701  }
702 };
703 
704 template <typename BitsType>
705 struct count_bits_impl<BitsType,
706  std::enable_if_t<std::is_integral<BitsType>::value && sizeof(unsigned int) < sizeof(BitsType) &&
707  sizeof(BitsType) <= sizeof(unsigned long)>> {
708  static constexpr int kNumBits = static_cast<int>(sizeof(BitsType) * CHAR_BIT);
709  static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
710  static constexpr int kLeadingBitsOffset = (sizeof(unsigned long) - sizeof(BitsType)) * CHAR_BIT;
711  return bits == 0 ? kNumBits : __builtin_clzl(static_cast<unsigned long>(bits)) - kLeadingBitsOffset;
712  }
713 
714  static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
715  return bits == 0 ? kNumBits : __builtin_ctzl(static_cast<unsigned long>(bits));
716  }
717 };
718 
719 template <typename BitsType>
720 struct count_bits_impl<BitsType,
721  std::enable_if_t<std::is_integral<BitsType>::value && sizeof(unsigned long) < sizeof(BitsType) &&
722  sizeof(BitsType) <= sizeof(unsigned long long)>> {
723  static constexpr int kNumBits = static_cast<int>(sizeof(BitsType) * CHAR_BIT);
724  static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
725  static constexpr int kLeadingBitsOffset = (sizeof(unsigned long long) - sizeof(BitsType)) * CHAR_BIT;
726  return bits == 0 ? kNumBits : __builtin_clzll(static_cast<unsigned long long>(bits)) - kLeadingBitsOffset;
727  }
728 
729  static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
730  return bits == 0 ? kNumBits : __builtin_ctzll(static_cast<unsigned long long>(bits));
731  }
732 };
733 
734 #elif EIGEN_COMP_MSVC
735 
736 template <typename BitsType>
737 struct count_bits_impl<
738  BitsType, std::enable_if_t<std::is_integral<BitsType>::value && sizeof(BitsType) <= sizeof(unsigned long)>> {
739  static constexpr int kNumBits = static_cast<int>(sizeof(BitsType) * CHAR_BIT);
740  static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
741  unsigned long out;
742  _BitScanReverse(&out, static_cast<unsigned long>(bits));
743  return bits == 0 ? kNumBits : (kNumBits - 1) - static_cast<int>(out);
744  }
745 
746  static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
747  unsigned long out;
748  _BitScanForward(&out, static_cast<unsigned long>(bits));
749  return bits == 0 ? kNumBits : static_cast<int>(out);
750  }
751 };
752 
753 #ifdef _WIN64
754 
755 template <typename BitsType>
756 struct count_bits_impl<BitsType,
757  std::enable_if_t<std::is_integral<BitsType>::value && sizeof(unsigned long) < sizeof(BitsType) &&
758  sizeof(BitsType) <= sizeof(__int64)>> {
759  static constexpr int kNumBits = static_cast<int>(sizeof(BitsType) * CHAR_BIT);
760  static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
761  unsigned long out;
762  _BitScanReverse64(&out, static_cast<unsigned __int64>(bits));
763  return bits == 0 ? kNumBits : (kNumBits - 1) - static_cast<int>(out);
764  }
765 
766  static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
767  unsigned long out;
768  _BitScanForward64(&out, static_cast<unsigned __int64>(bits));
769  return bits == 0 ? kNumBits : static_cast<int>(out);
770  }
771 };
772 
773 #endif // _WIN64
774 
775 #endif // EIGEN_COMP_GNUC || EIGEN_COMP_CLANG
776 
777 template <typename BitsType>
778 struct log_2_impl {
779  static constexpr int kTotalBits = sizeof(BitsType) * CHAR_BIT;
780  static EIGEN_DEVICE_FUNC inline int run_ceil(const BitsType& x) {
781  const int n = kTotalBits - clz(x);
782  bool power_of_two = (x & (x - 1)) == 0;
783  return x == 0 ? 0 : power_of_two ? (n - 1) : n;
784  }
785  static EIGEN_DEVICE_FUNC inline int run_floor(const BitsType& x) {
786  const int n = kTotalBits - clz(x);
787  return x == 0 ? 0 : n - 1;
788  }
789 };
790 
791 template <typename BitsType>
792 int log2_ceil(const BitsType& x) {
793  return log_2_impl<BitsType>::run_ceil(x);
794 }
795 
796 template <typename BitsType>
797 int log2_floor(const BitsType& x) {
798  return log_2_impl<BitsType>::run_floor(x);
799 }
800 
801 // Implementation of is* functions
802 
803 template <typename T>
804 EIGEN_DEVICE_FUNC std::enable_if_t<!(std::numeric_limits<T>::has_infinity || std::numeric_limits<T>::has_quiet_NaN ||
805  std::numeric_limits<T>::has_signaling_NaN),
806  bool>
807 isfinite_impl(const T&) {
808  return true;
809 }
810 
811 template <typename T>
812 EIGEN_DEVICE_FUNC std::enable_if_t<(std::numeric_limits<T>::has_infinity || std::numeric_limits<T>::has_quiet_NaN ||
813  std::numeric_limits<T>::has_signaling_NaN) &&
814  (!NumTraits<T>::IsComplex),
815  bool>
816 isfinite_impl(const T& x) {
817  EIGEN_USING_STD(isfinite);
818  return isfinite EIGEN_NOT_A_MACRO(x);
819 }
820 
821 template <typename T>
822 EIGEN_DEVICE_FUNC std::enable_if_t<!std::numeric_limits<T>::has_infinity, bool> isinf_impl(const T&) {
823  return false;
824 }
825 
826 template <typename T>
827 EIGEN_DEVICE_FUNC std::enable_if_t<(std::numeric_limits<T>::has_infinity && !NumTraits<T>::IsComplex), bool> isinf_impl(
828  const T& x) {
829  EIGEN_USING_STD(isinf);
830  return isinf EIGEN_NOT_A_MACRO(x);
831 }
832 
833 template <typename T>
834 EIGEN_DEVICE_FUNC
835 std::enable_if_t<!(std::numeric_limits<T>::has_quiet_NaN || std::numeric_limits<T>::has_signaling_NaN), bool>
836 isnan_impl(const T&) {
837  return false;
838 }
839 
840 template <typename T>
841 EIGEN_DEVICE_FUNC std::enable_if_t<
842  (std::numeric_limits<T>::has_quiet_NaN || std::numeric_limits<T>::has_signaling_NaN) && (!NumTraits<T>::IsComplex),
843  bool>
844 isnan_impl(const T& x) {
845  EIGEN_USING_STD(isnan);
846  return isnan EIGEN_NOT_A_MACRO(x);
847 }
848 
849 // The following overload are defined at the end of this file
850 template <typename T>
851 EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x);
852 template <typename T>
853 EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x);
854 template <typename T>
855 EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x);
856 template <typename T>
857 EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS T ptanh_float(const T& a_x);
858 
859 /****************************************************************************
860  * Implementation of sign *
861  ****************************************************************************/
862 template <typename Scalar, bool IsComplex = (NumTraits<Scalar>::IsComplex != 0),
863  bool IsInteger = (NumTraits<Scalar>::IsInteger != 0)>
864 struct sign_impl {
865  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) { return Scalar((a > Scalar(0)) - (a < Scalar(0))); }
866 };
867 
868 template <typename Scalar>
869 struct sign_impl<Scalar, false, false> {
870  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) {
871  return (isnan_impl<Scalar>)(a) ? a : Scalar((a > Scalar(0)) - (a < Scalar(0)));
872  }
873 };
874 
875 template <typename Scalar, bool IsInteger>
876 struct sign_impl<Scalar, true, IsInteger> {
877  EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) {
878  using real_type = typename NumTraits<Scalar>::Real;
879  EIGEN_USING_STD(abs);
880  real_type aa = abs(a);
881  if (aa == real_type(0)) return Scalar(0);
882  aa = real_type(1) / aa;
883  return Scalar(numext::real(a) * aa, numext::imag(a) * aa);
884  }
885 };
886 
887 // The sign function for bool is the identity.
888 template <>
889 struct sign_impl<bool, false, true> {
890  EIGEN_DEVICE_FUNC static inline bool run(const bool& a) { return a; }
891 };
892 
893 template <typename Scalar>
894 struct sign_retval {
895  typedef Scalar type;
896 };
897 
898 // suppress "unary minus operator applied to unsigned type, result still unsigned" warnings on MSVC
899 // note: `0 - a` is distinct from `-a` when Scalar is a floating point type and `a` is zero
900 
901 template <typename Scalar, bool IsInteger = NumTraits<Scalar>::IsInteger>
902 struct negate_impl {
903  static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar run(const Scalar& a) { return -a; }
904 };
905 
906 template <typename Scalar>
907 struct negate_impl<Scalar, true> {
908  EIGEN_STATIC_ASSERT((!is_same<Scalar, bool>::value), NEGATE IS NOT DEFINED FOR BOOLEAN TYPES)
909  static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar run(const Scalar& a) { return Scalar(0) - a; }
910 };
911 
912 template <typename Scalar>
913 struct negate_retval {
914  typedef Scalar type;
915 };
916 
917 template <typename Scalar, bool IsInteger = NumTraits<typename unpacket_traits<Scalar>::type>::IsInteger>
918 struct nearest_integer_impl {
919  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_floor(const Scalar& x) {
920  EIGEN_USING_STD(floor) return floor(x);
921  }
922  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_ceil(const Scalar& x) {
923  EIGEN_USING_STD(ceil) return ceil(x);
924  }
925  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_rint(const Scalar& x) {
926  EIGEN_USING_STD(rint) return rint(x);
927  }
928  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_round(const Scalar& x) {
929  EIGEN_USING_STD(round) return round(x);
930  }
931  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_trunc(const Scalar& x) {
932  EIGEN_USING_STD(trunc) return trunc(x);
933  }
934 };
935 template <typename Scalar>
936 struct nearest_integer_impl<Scalar, true> {
937  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_floor(const Scalar& x) { return x; }
938  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_ceil(const Scalar& x) { return x; }
939  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_rint(const Scalar& x) { return x; }
940  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_round(const Scalar& x) { return x; }
941  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_trunc(const Scalar& x) { return x; }
942 };
943 
944 // Extra namespace to prevent leaking std::fma into Eigen::internal.
945 namespace has_fma_detail {
946 
947 template <typename T, typename EnableIf = void>
948 struct has_fma_impl : public std::false_type {};
949 
950 using std::fma;
951 
952 template <typename T>
953 struct has_fma_impl<
954  T, std::enable_if_t<std::is_same<T, decltype(fma(std::declval<T>(), std::declval<T>(), std::declval<T>()))>::value>>
955  : public std::true_type {};
956 
957 } // namespace has_fma_detail
958 
959 template <typename T>
960 struct has_fma : public has_fma_detail::has_fma_impl<T> {};
961 
962 // Default implementation.
963 template <typename T, typename Enable = void>
964 struct fma_impl {
965  static_assert(has_fma<T>::value, "No function fma(...) for type. Please provide an implementation.");
966 };
967 
968 // STD or ADL version if it exists.
969 template <typename T>
970 struct fma_impl<T, std::enable_if_t<has_fma<T>::value>> {
971  static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T run(const T& a, const T& b, const T& c) {
972  using std::fma;
973  return fma(a, b, c);
974  }
975 };
976 
977 #if defined(EIGEN_GPUCC)
978 template <>
979 struct has_fma<float> : public true_type {};
980 
981 template <>
982 struct fma_impl<float, void> {
983  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float run(const float& a, const float& b, const float& c) {
984  return ::fmaf(a, b, c);
985  }
986 };
987 
988 template <>
989 struct has_fma<double> : public true_type {};
990 
991 template <>
992 struct fma_impl<double, void> {
993  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double run(const double& a, const double& b, const double& c) {
994  return ::fma(a, b, c);
995  }
996 };
997 #endif
998 
999 // Basic multiply-add.
1000 template <typename Scalar, typename EnableIf = void>
1001 struct madd_impl {
1002  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run(const Scalar& x, const Scalar& y, const Scalar& z) {
1003  return x * y + z;
1004  }
1005 };
1006 
1007 #if EIGEN_SCALAR_MADD_USE_FMA
1008 template <typename Scalar>
1009 struct madd_impl<Scalar, std::enable_if_t<has_fma<Scalar>::value>> {
1010  static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run(const Scalar& x, const Scalar& y, const Scalar& z) {
1011  return fma_impl<Scalar>::run(x, y, z);
1012  }
1013 };
1014 #endif
1015 
1016 } // end namespace internal
1017 
1018 /****************************************************************************
1019  * Generic math functions *
1020  ****************************************************************************/
1021 
1022 namespace numext {
1023 
1024 #if (!defined(EIGEN_GPUCC) || defined(EIGEN_CONSTEXPR_ARE_DEVICE_FUNC))
1025 template <typename T>
1026 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y) {
1027  EIGEN_USING_STD(min)
1028  return min EIGEN_NOT_A_MACRO(x, y);
1029 }
1030 
1031 template <typename T>
1032 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y) {
1033  EIGEN_USING_STD(max)
1034  return max EIGEN_NOT_A_MACRO(x, y);
1035 }
1036 #else
1037 template <typename T>
1038 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y) {
1039  return y < x ? y : x;
1040 }
1041 template <>
1042 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float mini(const float& x, const float& y) {
1043  return fminf(x, y);
1044 }
1045 template <>
1046 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double mini(const double& x, const double& y) {
1047  return fmin(x, y);
1048 }
1049 
1050 #ifndef EIGEN_GPU_COMPILE_PHASE
1051 template <>
1052 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double mini(const long double& x, const long double& y) {
1053 #if defined(EIGEN_HIPCC)
1054  // no "fminl" on HIP yet
1055  return (x < y) ? x : y;
1056 #else
1057  return fminl(x, y);
1058 #endif
1059 }
1060 #endif
1061 
1062 template <typename T>
1063 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y) {
1064  return x < y ? y : x;
1065 }
1066 template <>
1067 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float maxi(const float& x, const float& y) {
1068  return fmaxf(x, y);
1069 }
1070 template <>
1071 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double maxi(const double& x, const double& y) {
1072  return fmax(x, y);
1073 }
1074 #ifndef EIGEN_GPU_COMPILE_PHASE
1075 template <>
1076 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double maxi(const long double& x, const long double& y) {
1077 #if defined(EIGEN_HIPCC)
1078  // no "fmaxl" on HIP yet
1079  return (x > y) ? x : y;
1080 #else
1081  return fmaxl(x, y);
1082 #endif
1083 }
1084 #endif
1085 #endif
1086 
1087 #if defined(SYCL_DEVICE_ONLY)
1088 
1089 #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1090  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \
1091  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \
1092  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \
1093  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1094 #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1095  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \
1096  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \
1097  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \
1098  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1099 #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1100  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \
1101  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1102  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \
1103  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1104 #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1105  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \
1106  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1107  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \
1108  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1109 #define SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(NAME, FUNC) \
1110  SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1111  SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC)
1112 #define SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(NAME, FUNC) \
1113  SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1114  SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC)
1115 #define SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(NAME, FUNC) \
1116  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1117  SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_double)
1118 #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(NAME, FUNC) \
1119  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1120  SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_double)
1121 #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(NAME, FUNC, RET_TYPE) \
1122  SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_float) \
1123  SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_double)
1124 
1125 #define SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1126  template <> \
1127  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE& x) { \
1128  return cl::sycl::FUNC(x); \
1129  }
1130 
1131 #define SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, TYPE) SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, TYPE, TYPE)
1132 
1133 #define SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE1, ARG_TYPE2) \
1134  template <> \
1135  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE1& x, const ARG_TYPE2& y) { \
1136  return cl::sycl::FUNC(x, y); \
1137  }
1138 
1139 #define SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1140  SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE, ARG_TYPE)
1141 
1142 #define SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, TYPE) SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, TYPE, TYPE)
1143 
1144 SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(mini, min)
1145 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(mini, fmin)
1146 SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(maxi, max)
1147 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(maxi, fmax)
1148 
1149 #endif
1150 
1151 template <typename Scalar>
1152 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(arg, Scalar) arg(const Scalar& x) {
1153  return EIGEN_MATHFUNC_IMPL(arg, Scalar)::run(x);
1154 }
1155 
1156 template <typename Scalar>
1157 EIGEN_DEVICE_FUNC inline internal::add_const_on_value_type_t<EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar)> imag_ref(
1158  const Scalar& x) {
1159  return internal::imag_ref_impl<Scalar>::run(x);
1160 }
1161 
1162 template <typename Scalar>
1163 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x) {
1164  return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(x);
1165 }
1166 
1167 template <typename Scalar>
1168 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x) {
1169  return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x);
1170 }
1171 
1172 template <typename Scalar>
1173 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(sign, Scalar) sign(const Scalar& x) {
1174  return EIGEN_MATHFUNC_IMPL(sign, Scalar)::run(x);
1175 }
1176 
1177 template <typename Scalar>
1178 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(negate, Scalar) negate(const Scalar& x) {
1179  return EIGEN_MATHFUNC_IMPL(negate, Scalar)::run(x);
1180 }
1181 
1182 template <typename Scalar>
1183 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x) {
1184  return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x);
1185 }
1186 
1187 EIGEN_DEVICE_FUNC inline bool abs2(bool x) { return x; }
1188 
1189 template <typename T>
1190 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T absdiff(const T& x, const T& y) {
1191  return x > y ? x - y : y - x;
1192 }
1193 template <>
1194 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float absdiff(const float& x, const float& y) {
1195  return fabsf(x - y);
1196 }
1197 template <>
1198 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double absdiff(const double& x, const double& y) {
1199  return fabs(x - y);
1200 }
1201 
1202 // HIP and CUDA do not support long double.
1203 #ifndef EIGEN_GPU_COMPILE_PHASE
1204 template <>
1205 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double absdiff(const long double& x, const long double& y) {
1206  return fabsl(x - y);
1207 }
1208 #endif
1209 
1210 template <typename Scalar>
1211 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x) {
1212  return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x);
1213 }
1214 
1215 template <typename Scalar>
1216 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(hypot, Scalar) hypot(const Scalar& x, const Scalar& y) {
1217  return EIGEN_MATHFUNC_IMPL(hypot, Scalar)::run(x, y);
1218 }
1219 
1220 #if defined(SYCL_DEVICE_ONLY)
1221 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(hypot, hypot)
1222 #endif
1223 
1224 template <typename Scalar>
1225 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(log1p, Scalar) log1p(const Scalar& x) {
1226  return EIGEN_MATHFUNC_IMPL(log1p, Scalar)::run(x);
1227 }
1228 
1229 #if defined(SYCL_DEVICE_ONLY)
1230 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log1p, log1p)
1231 #endif
1232 
1233 #if defined(EIGEN_GPUCC)
1234 template <>
1235 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float log1p(const float& x) {
1236  return ::log1pf(x);
1237 }
1238 
1239 template <>
1240 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double log1p(const double& x) {
1241  return ::log1p(x);
1242 }
1243 #endif
1244 
1245 template <typename ScalarX, typename ScalarY>
1246 EIGEN_DEVICE_FUNC inline typename internal::pow_impl<ScalarX, ScalarY>::result_type pow(const ScalarX& x,
1247  const ScalarY& y) {
1248  return internal::pow_impl<ScalarX, ScalarY>::run(x, y);
1249 }
1250 
1251 #if defined(SYCL_DEVICE_ONLY)
1252 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(pow, pow)
1253 #endif
1254 
1255 template <typename T>
1256 EIGEN_DEVICE_FUNC bool(isnan)(const T& x) {
1257  return internal::isnan_impl(x);
1258 }
1259 template <typename T>
1260 EIGEN_DEVICE_FUNC bool(isinf)(const T& x) {
1261  return internal::isinf_impl(x);
1262 }
1263 template <typename T>
1264 EIGEN_DEVICE_FUNC bool(isfinite)(const T& x) {
1265  return internal::isfinite_impl(x);
1266 }
1267 
1268 #if defined(SYCL_DEVICE_ONLY)
1269 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isnan, isnan, bool)
1270 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isinf, isinf, bool)
1271 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isfinite, isfinite, bool)
1272 #endif
1273 
1274 template <typename Scalar>
1275 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar rint(const Scalar& x) {
1276  return internal::nearest_integer_impl<Scalar>::run_rint(x);
1277 }
1278 
1279 template <typename Scalar>
1280 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar round(const Scalar& x) {
1281  return internal::nearest_integer_impl<Scalar>::run_round(x);
1282 }
1283 
1284 template <typename Scalar>
1285 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(floor)(const Scalar& x) {
1286  return internal::nearest_integer_impl<Scalar>::run_floor(x);
1287 }
1288 
1289 template <typename Scalar>
1290 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(ceil)(const Scalar& x) {
1291  return internal::nearest_integer_impl<Scalar>::run_ceil(x);
1292 }
1293 
1294 template <typename Scalar>
1295 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(trunc)(const Scalar& x) {
1296  return internal::nearest_integer_impl<Scalar>::run_trunc(x);
1297 }
1298 
1299 #if defined(SYCL_DEVICE_ONLY)
1300 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(round, round)
1301 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(floor, floor)
1302 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(ceil, ceil)
1303 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(trunc, trunc)
1304 #endif
1305 
1306 #if defined(EIGEN_GPUCC)
1307 template <>
1308 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float floor(const float& x) {
1309  return ::floorf(x);
1310 }
1311 template <>
1312 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double floor(const double& x) {
1313  return ::floor(x);
1314 }
1315 template <>
1316 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float ceil(const float& x) {
1317  return ::ceilf(x);
1318 }
1319 template <>
1320 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double ceil(const double& x) {
1321  return ::ceil(x);
1322 }
1323 template <>
1324 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float trunc(const float& x) {
1325  return ::truncf(x);
1326 }
1327 template <>
1328 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double trunc(const double& x) {
1329  return ::trunc(x);
1330 }
1331 #endif
1332 
1333 // Integer division with rounding up.
1334 // T is assumed to be an integer type with a>=0, and b>0
1335 template <typename T>
1336 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE constexpr T div_ceil(T a, T b) {
1337  using UnsignedT = typename internal::make_unsigned<T>::type;
1338  EIGEN_STATIC_ASSERT((NumTraits<T>::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES)
1339  // Note: explicitly declaring a and b as non-negative values allows the compiler to use better optimizations
1340  const UnsignedT ua = UnsignedT(a);
1341  const UnsignedT ub = UnsignedT(b);
1342  // Note: This form is used because it cannot overflow.
1343  return ua == 0 ? 0 : (ua - 1) / ub + 1;
1344 }
1345 
1346 // Integer round down to nearest power of b
1347 // T is assumed to be an integer type with a>=0, and b>0
1348 template <typename T, typename U>
1349 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE constexpr T round_down(T a, U b) {
1350  using UnsignedT = typename internal::make_unsigned<T>::type;
1351  using UnsignedU = typename internal::make_unsigned<U>::type;
1352  EIGEN_STATIC_ASSERT((NumTraits<T>::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES)
1353  EIGEN_STATIC_ASSERT((NumTraits<U>::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES)
1354  // Note: explicitly declaring a and b as non-negative values allows the compiler to use better optimizations
1355  const UnsignedT ua = UnsignedT(a);
1356  const UnsignedU ub = UnsignedU(b);
1357  return ub * (ua / ub);
1358 }
1359 
1362 constexpr int log2(int x) {
1363  unsigned int v(x);
1364  constexpr int table[32] = {0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30,
1365  8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31};
1366  v |= v >> 1;
1367  v |= v >> 2;
1368  v |= v >> 4;
1369  v |= v >> 8;
1370  v |= v >> 16;
1371  return table[(v * 0x07C4ACDDU) >> 27];
1372 }
1373 
1383 template <typename Scalar>
1384 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE EIGEN_MATHFUNC_RETVAL(sqrt, Scalar) sqrt(const Scalar& x) {
1385  return EIGEN_MATHFUNC_IMPL(sqrt, Scalar)::run(x);
1386 }
1387 
1388 // Boolean specialization, avoids implicit float to bool conversion (-Wimplicit-conversion-floating-point-to-bool).
1389 template <>
1390 EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS EIGEN_DEVICE_FUNC bool sqrt<bool>(const bool& x) {
1391  return x;
1392 }
1393 
1394 #if defined(SYCL_DEVICE_ONLY)
1395 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sqrt, sqrt)
1396 #endif
1397 
1399 template <typename T>
1400 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<!NumTraits<T>::IsComplex, T> cbrt(const T& x) {
1401  EIGEN_USING_STD(cbrt);
1402  return static_cast<T>(cbrt(x));
1403 }
1404 
1405 template <typename T>
1406 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<NumTraits<T>::IsComplex, T> cbrt(const T& x) {
1407  EIGEN_USING_STD(pow);
1408  return pow(x, typename NumTraits<T>::Real(1.0 / 3.0));
1409 }
1410 
1412 template <typename T>
1413 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T rsqrt(const T& x) {
1414  return internal::rsqrt_impl<T>::run(x);
1415 }
1416 
1417 template <typename T>
1418 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T log(const T& x) {
1419  return internal::log_impl<T>::run(x);
1420 }
1421 
1422 #if defined(SYCL_DEVICE_ONLY)
1423 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log, log)
1424 #endif
1425 
1426 #if defined(EIGEN_GPUCC)
1427 template <>
1428 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float log(const float& x) {
1429  return ::logf(x);
1430 }
1431 
1432 template <>
1433 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double log(const double& x) {
1434  return ::log(x);
1435 }
1436 #endif
1437 
1438 template <typename T>
1439 EIGEN_DEVICE_FUNC
1440 EIGEN_ALWAYS_INLINE std::enable_if_t<NumTraits<T>::IsSigned || NumTraits<T>::IsComplex, typename NumTraits<T>::Real>
1441 abs(const T& x) {
1442  EIGEN_USING_STD(abs);
1443  return abs(x);
1444 }
1445 
1446 template <typename T>
1447 EIGEN_DEVICE_FUNC
1448 EIGEN_ALWAYS_INLINE std::enable_if_t<!(NumTraits<T>::IsSigned || NumTraits<T>::IsComplex), typename NumTraits<T>::Real>
1449 abs(const T& x) {
1450  return x;
1451 }
1452 
1453 #if defined(SYCL_DEVICE_ONLY)
1454 SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(abs, abs)
1455 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(abs, fabs)
1456 #endif
1457 
1458 #if defined(EIGEN_GPUCC)
1459 template <>
1460 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float abs(const float& x) {
1461  return ::fabsf(x);
1462 }
1463 
1464 template <>
1465 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double abs(const double& x) {
1466  return ::fabs(x);
1467 }
1468 
1469 template <>
1470 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float abs(const std::complex<float>& x) {
1471  return ::hypotf(x.real(), x.imag());
1472 }
1473 
1474 template <>
1475 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double abs(const std::complex<double>& x) {
1476  return ::hypot(x.real(), x.imag());
1477 }
1478 #endif
1479 
1480 template <typename Scalar, bool IsInteger = NumTraits<Scalar>::IsInteger, bool IsSigned = NumTraits<Scalar>::IsSigned>
1481 struct signbit_impl;
1482 template <typename Scalar>
1483 struct signbit_impl<Scalar, false, true> {
1484  static constexpr size_t Size = sizeof(Scalar);
1485  static constexpr size_t Shift = (CHAR_BIT * Size) - 1;
1486  using intSize_t = typename get_integer_by_size<Size>::signed_type;
1487  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static Scalar run(const Scalar& x) {
1488  intSize_t a = bit_cast<intSize_t, Scalar>(x);
1489  a = a >> Shift;
1490  Scalar result = bit_cast<Scalar, intSize_t>(a);
1491  return result;
1492  }
1493 };
1494 template <typename Scalar>
1495 struct signbit_impl<Scalar, true, true> {
1496  static constexpr size_t Size = sizeof(Scalar);
1497  static constexpr size_t Shift = (CHAR_BIT * Size) - 1;
1498  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar run(const Scalar& x) { return x >> Shift; }
1499 };
1500 template <typename Scalar>
1501 struct signbit_impl<Scalar, true, false> {
1502  EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar run(const Scalar&) { return Scalar(0); }
1503 };
1504 template <typename Scalar>
1505 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar signbit(const Scalar& x) {
1506  return signbit_impl<Scalar>::run(x);
1507 }
1508 
1509 template <typename T>
1510 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T exp(const T& x) {
1511  EIGEN_USING_STD(exp);
1512  return exp(x);
1513 }
1514 
1515 // MSVC screws up some edge-cases for std::exp(complex).
1516 #ifdef EIGEN_COMP_MSVC
1517 template <typename RealScalar>
1518 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<RealScalar> exp(const std::complex<RealScalar>& x) {
1519  EIGEN_USING_STD(exp);
1520  // If z is (x,±∞) (for any finite x), the result is (NaN,NaN) and FE_INVALID is raised.
1521  // If z is (x,NaN) (for any finite x), the result is (NaN,NaN) and FE_INVALID may be raised.
1522  if ((isfinite)(real_ref(x)) && !(isfinite)(imag_ref(x))) {
1523  return std::complex<RealScalar>(NumTraits<RealScalar>::quiet_NaN(), NumTraits<RealScalar>::quiet_NaN());
1524  }
1525  // If z is (+∞,±∞), the result is (±∞,NaN) and FE_INVALID is raised (the sign of the real part is unspecified)
1526  // If z is (+∞,NaN), the result is (±∞,NaN) (the sign of the real part is unspecified)
1527  if ((real_ref(x) == NumTraits<RealScalar>::infinity() && !(isfinite)(imag_ref(x)))) {
1528  return std::complex<RealScalar>(NumTraits<RealScalar>::infinity(), NumTraits<RealScalar>::quiet_NaN());
1529  }
1530  return exp(x);
1531 }
1532 #endif
1533 
1534 #if defined(SYCL_DEVICE_ONLY)
1535 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp, exp)
1536 #endif
1537 
1538 #if defined(EIGEN_GPUCC)
1539 template <>
1540 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float exp(const float& x) {
1541  return ::expf(x);
1542 }
1543 
1544 template <>
1545 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double exp(const double& x) {
1546  return ::exp(x);
1547 }
1548 
1549 template <>
1550 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<float> exp(const std::complex<float>& x) {
1551  float com = ::expf(x.real());
1552  float res_real = com * ::cosf(x.imag());
1553  float res_imag = com * ::sinf(x.imag());
1554  return std::complex<float>(res_real, res_imag);
1555 }
1556 
1557 template <>
1558 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<double> exp(const std::complex<double>& x) {
1559  double com = ::exp(x.real());
1560  double res_real = com * ::cos(x.imag());
1561  double res_imag = com * ::sin(x.imag());
1562  return std::complex<double>(res_real, res_imag);
1563 }
1564 #endif
1565 
1566 template <typename T>
1567 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T exp2(const T& x) {
1568  EIGEN_USING_STD(exp2);
1569  return exp2(x);
1570 }
1571 
1572 // MSVC screws up some edge-cases for std::exp2(complex).
1573 #ifdef EIGEN_COMP_MSVC
1574 template <typename RealScalar>
1575 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<RealScalar> exp2(const std::complex<RealScalar>& x) {
1576  EIGEN_USING_STD(exp);
1577  // If z is (x,±∞) (for any finite x), the result is (NaN,NaN) and FE_INVALID is raised.
1578  // If z is (x,NaN) (for any finite x), the result is (NaN,NaN) and FE_INVALID may be raised.
1579  if ((isfinite)(real_ref(x)) && !(isfinite)(imag_ref(x))) {
1580  return std::complex<RealScalar>(NumTraits<RealScalar>::quiet_NaN(), NumTraits<RealScalar>::quiet_NaN());
1581  }
1582  // If z is (+∞,±∞), the result is (±∞,NaN) and FE_INVALID is raised (the sign of the real part is unspecified)
1583  // If z is (+∞,NaN), the result is (±∞,NaN) (the sign of the real part is unspecified)
1584  if ((real_ref(x) == NumTraits<RealScalar>::infinity() && !(isfinite)(imag_ref(x)))) {
1585  return std::complex<RealScalar>(NumTraits<RealScalar>::infinity(), NumTraits<RealScalar>::quiet_NaN());
1586  }
1587  return exp2(x);
1588 }
1589 #endif
1590 
1591 #if defined(SYCL_DEVICE_ONLY)
1592 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp2, exp2)
1593 #endif
1594 
1595 #if defined(EIGEN_GPUCC)
1596 template <>
1597 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float exp2(const float& x) {
1598  return ::exp2f(x);
1599 }
1600 
1601 template <>
1602 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double exp2(const double& x) {
1603  return ::exp2(x);
1604 }
1605 
1606 template <>
1607 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<float> exp2(const std::complex<float>& x) {
1608  float com = ::exp2f(x.real());
1609  float res_real = com * ::cosf(static_cast<float>(EIGEN_LN2) * x.imag());
1610  float res_imag = com * ::sinf(static_cast<float>(EIGEN_LN2) * x.imag());
1611  return std::complex<float>(res_real, res_imag);
1612 }
1613 
1614 template <>
1615 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<double> exp2(const std::complex<double>& x) {
1616  double com = ::exp2(x.real());
1617  double res_real = com * ::cos(static_cast<double>(EIGEN_LN2) * x.imag());
1618  double res_imag = com * ::sin(static_cast<double>(EIGEN_LN2) * x.imag());
1619  return std::complex<double>(res_real, res_imag);
1620 }
1621 #endif
1622 
1623 template <typename Scalar>
1624 EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(expm1, Scalar) expm1(const Scalar& x) {
1625  return EIGEN_MATHFUNC_IMPL(expm1, Scalar)::run(x);
1626 }
1627 
1628 #if defined(SYCL_DEVICE_ONLY)
1629 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(expm1, expm1)
1630 #endif
1631 
1632 #if defined(EIGEN_GPUCC)
1633 template <>
1634 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float expm1(const float& x) {
1635  return ::expm1f(x);
1636 }
1637 
1638 template <>
1639 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double expm1(const double& x) {
1640  return ::expm1(x);
1641 }
1642 #endif
1643 
1644 template <typename T>
1645 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T cos(const T& x) {
1646  EIGEN_USING_STD(cos);
1647  return cos(x);
1648 }
1649 
1650 #if defined(SYCL_DEVICE_ONLY)
1651 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cos, cos)
1652 #endif
1653 
1654 #if defined(EIGEN_GPUCC)
1655 template <>
1656 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float cos(const float& x) {
1657  return ::cosf(x);
1658 }
1659 
1660 template <>
1661 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double cos(const double& x) {
1662  return ::cos(x);
1663 }
1664 #endif
1665 
1666 template <typename T>
1667 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T sin(const T& x) {
1668  EIGEN_USING_STD(sin);
1669  return sin(x);
1670 }
1671 
1672 #if defined(SYCL_DEVICE_ONLY)
1673 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sin, sin)
1674 #endif
1675 
1676 #if defined(EIGEN_GPUCC)
1677 template <>
1678 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float sin(const float& x) {
1679  return ::sinf(x);
1680 }
1681 
1682 template <>
1683 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double sin(const double& x) {
1684  return ::sin(x);
1685 }
1686 #endif
1687 
1688 template <typename T>
1689 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T tan(const T& x) {
1690  EIGEN_USING_STD(tan);
1691  return tan(x);
1692 }
1693 
1694 #if defined(SYCL_DEVICE_ONLY)
1695 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tan, tan)
1696 #endif
1697 
1698 #if defined(EIGEN_GPUCC)
1699 template <>
1700 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tan(const float& x) {
1701  return ::tanf(x);
1702 }
1703 
1704 template <>
1705 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double tan(const double& x) {
1706  return ::tan(x);
1707 }
1708 #endif
1709 
1710 template <typename T>
1711 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acos(const T& x) {
1712  EIGEN_USING_STD(acos);
1713  return acos(x);
1714 }
1715 
1716 template <typename T>
1717 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acosh(const T& x) {
1718  EIGEN_USING_STD(acosh);
1719  return static_cast<T>(acosh(x));
1720 }
1721 
1722 #if defined(SYCL_DEVICE_ONLY)
1723 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acos, acos)
1724 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acosh, acosh)
1725 #endif
1726 
1727 #if defined(EIGEN_GPUCC)
1728 template <>
1729 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float acos(const float& x) {
1730  return ::acosf(x);
1731 }
1732 
1733 template <>
1734 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double acos(const double& x) {
1735  return ::acos(x);
1736 }
1737 #endif
1738 
1739 template <typename T>
1740 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asin(const T& x) {
1741  EIGEN_USING_STD(asin);
1742  return asin(x);
1743 }
1744 
1745 template <typename T>
1746 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asinh(const T& x) {
1747  EIGEN_USING_STD(asinh);
1748  return static_cast<T>(asinh(x));
1749 }
1750 
1751 #if defined(SYCL_DEVICE_ONLY)
1752 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asin, asin)
1753 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asinh, asinh)
1754 #endif
1755 
1756 #if defined(EIGEN_GPUCC)
1757 template <>
1758 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float asin(const float& x) {
1759  return ::asinf(x);
1760 }
1761 
1762 template <>
1763 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double asin(const double& x) {
1764  return ::asin(x);
1765 }
1766 #endif
1767 
1768 template <typename T>
1769 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atan(const T& x) {
1770  EIGEN_USING_STD(atan);
1771  return static_cast<T>(atan(x));
1772 }
1773 
1774 template <typename T, std::enable_if_t<!NumTraits<T>::IsComplex, int> = 0>
1775 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atan2(const T& y, const T& x) {
1776  EIGEN_USING_STD(atan2);
1777  return static_cast<T>(atan2(y, x));
1778 }
1779 
1780 template <typename T>
1781 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atanh(const T& x) {
1782  EIGEN_USING_STD(atanh);
1783  return static_cast<T>(atanh(x));
1784 }
1785 
1786 #if defined(SYCL_DEVICE_ONLY)
1787 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atan, atan)
1788 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atanh, atanh)
1789 #endif
1790 
1791 #if defined(EIGEN_GPUCC)
1792 template <>
1793 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float atan(const float& x) {
1794  return ::atanf(x);
1795 }
1796 
1797 template <>
1798 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double atan(const double& x) {
1799  return ::atan(x);
1800 }
1801 #endif
1802 
1803 template <typename T>
1804 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T cosh(const T& x) {
1805  EIGEN_USING_STD(cosh);
1806  return static_cast<T>(cosh(x));
1807 }
1808 
1809 #if defined(SYCL_DEVICE_ONLY)
1810 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cosh, cosh)
1811 #endif
1812 
1813 #if defined(EIGEN_GPUCC)
1814 template <>
1815 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float cosh(const float& x) {
1816  return ::coshf(x);
1817 }
1818 
1819 template <>
1820 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double cosh(const double& x) {
1821  return ::cosh(x);
1822 }
1823 #endif
1824 
1825 template <typename T>
1826 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T sinh(const T& x) {
1827  EIGEN_USING_STD(sinh);
1828  return static_cast<T>(sinh(x));
1829 }
1830 
1831 #if defined(SYCL_DEVICE_ONLY)
1832 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sinh, sinh)
1833 #endif
1834 
1835 #if defined(EIGEN_GPUCC)
1836 template <>
1837 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float sinh(const float& x) {
1838  return ::sinhf(x);
1839 }
1840 
1841 template <>
1842 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double sinh(const double& x) {
1843  return ::sinh(x);
1844 }
1845 #endif
1846 
1847 template <typename T>
1848 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T tanh(const T& x) {
1849  EIGEN_USING_STD(tanh);
1850  return tanh(x);
1851 }
1852 
1853 #if (!defined(EIGEN_GPUCC)) && EIGEN_FAST_MATH && !defined(SYCL_DEVICE_ONLY)
1854 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tanh(float x) { return internal::ptanh_float(x); }
1855 #endif
1856 
1857 #if defined(SYCL_DEVICE_ONLY)
1858 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tanh, tanh)
1859 #endif
1860 
1861 #if defined(EIGEN_GPUCC)
1862 template <>
1863 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tanh(const float& x) {
1864  return ::tanhf(x);
1865 }
1866 
1867 template <>
1868 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double tanh(const double& x) {
1869  return ::tanh(x);
1870 }
1871 #endif
1872 
1873 template <typename T>
1874 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T fmod(const T& a, const T& b) {
1875  EIGEN_USING_STD(fmod);
1876  return fmod(a, b);
1877 }
1878 
1879 #if defined(SYCL_DEVICE_ONLY)
1880 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(fmod, fmod)
1881 #endif
1882 
1883 #if defined(EIGEN_GPUCC)
1884 template <>
1885 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float fmod(const float& a, const float& b) {
1886  return ::fmodf(a, b);
1887 }
1888 
1889 template <>
1890 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double fmod(const double& a, const double& b) {
1891  return ::fmod(a, b);
1892 }
1893 #endif
1894 
1895 #if defined(SYCL_DEVICE_ONLY)
1896 #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY
1897 #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY
1898 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY
1899 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1900 #undef SYCL_SPECIALIZE_INTEGER_TYPES_BINARY
1901 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1902 #undef SYCL_SPECIALIZE_FLOATING_TYPES_BINARY
1903 #undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY
1904 #undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE
1905 #undef SYCL_SPECIALIZE_GEN_UNARY_FUNC
1906 #undef SYCL_SPECIALIZE_UNARY_FUNC
1907 #undef SYCL_SPECIALIZE_GEN1_BINARY_FUNC
1908 #undef SYCL_SPECIALIZE_GEN2_BINARY_FUNC
1909 #undef SYCL_SPECIALIZE_BINARY_FUNC
1910 #endif
1911 
1912 template <typename Scalar, typename Enable = std::enable_if_t<std::is_integral<Scalar>::value>>
1913 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar logical_shift_left(const Scalar& a, int n) {
1914  return a << n;
1915 }
1916 
1917 template <typename Scalar, typename Enable = std::enable_if_t<std::is_integral<Scalar>::value>>
1918 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar logical_shift_right(const Scalar& a, int n) {
1919  using UnsignedScalar = typename numext::get_integer_by_size<sizeof(Scalar)>::unsigned_type;
1920  return bit_cast<Scalar, UnsignedScalar>(bit_cast<UnsignedScalar, Scalar>(a) >> n);
1921 }
1922 
1923 template <typename Scalar, typename Enable = std::enable_if_t<std::is_integral<Scalar>::value>>
1924 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar arithmetic_shift_right(const Scalar& a, int n) {
1925  using SignedScalar = typename numext::get_integer_by_size<sizeof(Scalar)>::signed_type;
1926  return bit_cast<Scalar, SignedScalar>(bit_cast<SignedScalar, Scalar>(a) >> n);
1927 }
1928 
1929 template <typename Scalar>
1930 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar fma(const Scalar& x, const Scalar& y, const Scalar& z) {
1931  return internal::fma_impl<Scalar>::run(x, y, z);
1932 }
1933 
1934 // Multiply-add.
1935 template <typename Scalar>
1936 EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar madd(const Scalar& x, const Scalar& y, const Scalar& z) {
1937  return internal::madd_impl<Scalar>::run(x, y, z);
1938 }
1939 
1940 } // end namespace numext
1941 
1942 namespace internal {
1943 
1944 template <typename T>
1945 EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x) {
1946  return (numext::isfinite)(numext::real(x)) && (numext::isfinite)(numext::imag(x));
1947 }
1948 
1949 template <typename T>
1950 EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x) {
1951  return (numext::isnan)(numext::real(x)) || (numext::isnan)(numext::imag(x));
1952 }
1953 
1954 template <typename T>
1955 EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x) {
1956  return ((numext::isinf)(numext::real(x)) || (numext::isinf)(numext::imag(x))) && (!(numext::isnan)(x));
1957 }
1958 
1959 /****************************************************************************
1960  * Implementation of fuzzy comparisons *
1961  ****************************************************************************/
1962 
1963 template <typename Scalar, bool IsComplex, bool IsInteger>
1964 struct scalar_fuzzy_default_impl {};
1965 
1966 template <typename Scalar>
1967 struct scalar_fuzzy_default_impl<Scalar, false, false> {
1968  typedef typename NumTraits<Scalar>::Real RealScalar;
1969  template <typename OtherScalar>
1970  EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
1971  const RealScalar& prec) {
1972  return numext::abs(x) <= numext::abs(y) * prec;
1973  }
1974  EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) {
1975  return numext::abs(x - y) <= numext::mini(numext::abs(x), numext::abs(y)) * prec;
1976  }
1977  EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar& prec) {
1978  return x <= y || isApprox(x, y, prec);
1979  }
1980 };
1981 
1982 template <typename Scalar>
1983 struct scalar_fuzzy_default_impl<Scalar, false, true> {
1984  typedef typename NumTraits<Scalar>::Real RealScalar;
1985  template <typename OtherScalar>
1986  EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const Scalar&, const RealScalar&) {
1987  return x == Scalar(0);
1988  }
1989  EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar&) { return x == y; }
1990  EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar&) {
1991  return x <= y;
1992  }
1993 };
1994 
1995 template <typename Scalar>
1996 struct scalar_fuzzy_default_impl<Scalar, true, false> {
1997  typedef typename NumTraits<Scalar>::Real RealScalar;
1998  template <typename OtherScalar>
1999  EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
2000  const RealScalar& prec) {
2001  return numext::abs2(x) <= numext::abs2(y) * prec * prec;
2002  }
2003  EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) {
2004  return numext::abs2(x - y) <= numext::mini(numext::abs2(x), numext::abs2(y)) * prec * prec;
2005  }
2006 };
2007 
2008 template <typename Scalar>
2009 struct scalar_fuzzy_impl
2010  : scalar_fuzzy_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {};
2011 
2012 template <typename Scalar, typename OtherScalar>
2013 EIGEN_DEVICE_FUNC inline bool isMuchSmallerThan(
2014  const Scalar& x, const OtherScalar& y,
2015  const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
2016  return scalar_fuzzy_impl<Scalar>::template isMuchSmallerThan<OtherScalar>(x, y, precision);
2017 }
2018 
2019 template <typename Scalar>
2020 EIGEN_DEVICE_FUNC inline bool isApprox(
2021  const Scalar& x, const Scalar& y,
2022  const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
2023  return scalar_fuzzy_impl<Scalar>::isApprox(x, y, precision);
2024 }
2025 
2026 template <typename Scalar>
2027 EIGEN_DEVICE_FUNC inline bool isApproxOrLessThan(
2028  const Scalar& x, const Scalar& y,
2029  const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
2030  return scalar_fuzzy_impl<Scalar>::isApproxOrLessThan(x, y, precision);
2031 }
2032 
2033 /******************************************
2034 *** The special case of the bool type ***
2035 ******************************************/
2036 
2037 template <>
2038 struct scalar_fuzzy_impl<bool> {
2039  typedef bool RealScalar;
2040 
2041  template <typename OtherScalar>
2042  EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const bool& x, const bool&, const bool&) {
2043  return !x;
2044  }
2045 
2046  EIGEN_DEVICE_FUNC static inline bool isApprox(bool x, bool y, bool) { return x == y; }
2047 
2048  EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const bool& x, const bool& y, const bool&) {
2049  return (!x) || y;
2050  }
2051 };
2052 
2053 } // end namespace internal
2054 
2055 // Default implementations that rely on other numext implementations
2056 namespace internal {
2057 
2058 // Specialization for complex types that are not supported by std::expm1.
2059 template <typename RealScalar>
2060 struct expm1_impl<std::complex<RealScalar>> {
2061  EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
2062 
2063  EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(const std::complex<RealScalar>& x) {
2064  RealScalar xr = x.real();
2065  RealScalar xi = x.imag();
2066  // expm1(z) = exp(z) - 1
2067  // = exp(x + i * y) - 1
2068  // = exp(x) * (cos(y) + i * sin(y)) - 1
2069  // = exp(x) * cos(y) - 1 + i * exp(x) * sin(y)
2070  // Imag(expm1(z)) = exp(x) * sin(y)
2071  // Real(expm1(z)) = exp(x) * cos(y) - 1
2072  // = exp(x) * cos(y) - 1.
2073  // = expm1(x) + exp(x) * (cos(y) - 1)
2074  // = expm1(x) + exp(x) * (2 * sin(y / 2) ** 2)
2075  RealScalar erm1 = numext::expm1<RealScalar>(xr);
2076  RealScalar er = erm1 + RealScalar(1.);
2077  RealScalar sin2 = numext::sin(xi / RealScalar(2.));
2078  sin2 = sin2 * sin2;
2079  RealScalar s = numext::sin(xi);
2080  RealScalar real_part = erm1 - RealScalar(2.) * er * sin2;
2081  return std::complex<RealScalar>(real_part, er * s);
2082  }
2083 };
2084 
2085 template <typename T>
2086 struct rsqrt_impl {
2087  EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE T run(const T& x) { return T(1) / numext::sqrt(x); }
2088 };
2089 
2090 #if defined(EIGEN_GPU_COMPILE_PHASE)
2091 template <typename T>
2092 struct conj_impl<std::complex<T>, true> {
2093  EIGEN_DEVICE_FUNC static inline std::complex<T> run(const std::complex<T>& x) {
2094  return std::complex<T>(numext::real(x), -numext::imag(x));
2095  }
2096 };
2097 #endif
2098 
2099 } // end namespace internal
2100 
2101 } // end namespace Eigen
2102 
2103 #endif // EIGEN_MATHFUNCTIONS_H
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_tanh_op< typename Derived::Scalar >, const Derived > tanh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_sinh_op< typename Derived::Scalar >, const Derived > sinh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_isfinite_op< typename Derived::Scalar >, const Derived > isfinite(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_sqrt_op< typename Derived::Scalar >, const Derived > sqrt(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_conjugate_op< typename Derived::Scalar >, const Derived > conj(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_arg_op< typename Derived::Scalar >, const Derived > arg(const Eigen::ArrayBase< Derived > &x)
Namespace containing all symbols from the Eigen library.
Definition: B01_Experimental.dox:1
Definition: BFloat16.h:231
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_asin_op< typename Derived::Scalar >, const Derived > asin(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_acos_op< typename Derived::Scalar >, const Derived > acos(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_exp2_op< typename Derived::Scalar >, const Derived > exp2(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_cos_op< typename Derived::Scalar >, const Derived > cos(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_imag_op< typename Derived::Scalar >, const Derived > imag(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log1p_op< typename Derived::Scalar >, const Derived > log1p(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_real_op< typename Derived::Scalar >, const Derived > real(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_rsqrt_op< typename Derived::Scalar >, const Derived > rsqrt(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_abs_op< typename Derived::Scalar >, const Derived > abs(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_cosh_op< typename Derived::Scalar >, const Derived > cosh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log_op< typename Derived::Scalar >, const Derived > log(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_tan_op< typename Derived::Scalar >, const Derived > tan(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_expm1_op< typename Derived::Scalar >, const Derived > expm1(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_atanh_op< typename Derived::Scalar >, const Derived > atanh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_atan_op< typename Derived::Scalar >, const Derived > atan(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_sin_op< typename Derived::Scalar >, const Derived > sin(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_exp_op< typename Derived::Scalar >, const Derived > exp(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_cbrt_op< typename Derived::Scalar >, const Derived > cbrt(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_acosh_op< typename Derived::Scalar >, const Derived > acosh(const Eigen::ArrayBase< Derived > &x)
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_asinh_op< typename Derived::Scalar >, const Derived > asinh(const Eigen::ArrayBase< Derived > &x)