Eigen  5.0.0-dev
Loading...
Searching...
No Matches
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
23namespace Eigen {
24
25namespace internal {
26
47
48template <typename T, typename dummy = void>
49struct global_math_functions_filtering_base {
50 typedef T type;
51};
52
53template <typename T>
54struct always_void {
55 typedef void type;
56};
57
58template <typename T>
59struct 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
74template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
75struct 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
80template <typename Scalar>
81struct 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
89template <typename Scalar>
90struct real_impl : real_default_impl<Scalar> {};
91
92#if defined(EIGEN_GPU_COMPILE_PHASE)
93template <typename T>
94struct 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
100template <typename Scalar>
101struct real_retval {
102 typedef typename NumTraits<Scalar>::Real type;
103};
104
105/****************************************************************************
106 * Implementation of imag *
107 ****************************************************************************/
108
109template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
110struct 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
115template <typename Scalar>
116struct 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
124template <typename Scalar>
125struct imag_impl : imag_default_impl<Scalar> {};
126
127#if defined(EIGEN_GPU_COMPILE_PHASE)
128template <typename T>
129struct 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
135template <typename Scalar>
136struct imag_retval {
137 typedef typename NumTraits<Scalar>::Real type;
138};
139
140/****************************************************************************
141 * Implementation of real_ref *
142 ****************************************************************************/
143
144template <typename Scalar>
145struct 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
153template <typename Scalar>
154struct real_ref_retval {
155 typedef typename NumTraits<Scalar>::Real& type;
156};
157
158/****************************************************************************
159 * Implementation of imag_ref *
160 ****************************************************************************/
161
162template <typename Scalar, bool IsComplex>
163struct 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
171template <typename Scalar>
172struct 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
177template <typename Scalar>
178struct imag_ref_impl : imag_ref_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
179
180template <typename Scalar>
181struct imag_ref_retval {
182 typedef typename NumTraits<Scalar>::Real& type;
183};
184
185} // namespace internal
186
187namespace numext {
188
189template <typename Scalar>
190EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x) {
191 return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x);
192}
193
194template <typename Scalar>
195EIGEN_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
200template <typename Scalar>
201EIGEN_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
205template <typename Scalar>
206EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x) {
207 return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x);
208}
209
210template <typename Scalar>
211EIGEN_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
217namespace internal {
218
219/****************************************************************************
220 * Implementation of conj *
221 ****************************************************************************/
222
223template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
224struct conj_default_impl {
225 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { return x; }
226};
227
228template <typename Scalar>
229struct 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
236template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
237struct conj_impl : conj_default_impl<Scalar, IsComplex> {};
238
239template <typename Scalar>
240struct conj_retval {
241 typedef Scalar type;
242};
243
244/****************************************************************************
245 * Implementation of abs2 *
246 ****************************************************************************/
247
248template <typename Scalar, bool IsComplex>
249struct 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
254template <typename Scalar>
255struct 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
263template <typename Scalar>
264struct 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
271template <typename Scalar>
272struct abs2_retval {
273 typedef typename NumTraits<Scalar>::Real type;
274};
275
276/****************************************************************************
277 * Implementation of sqrt/rsqrt *
278 ****************************************************************************/
279
280template <typename Scalar>
281struct 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.
289template <typename ComplexT>
290EIGEN_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).
294template <typename T>
295struct 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
299template <typename Scalar>
300struct sqrt_retval {
301 typedef Scalar type;
302};
303
304// Default implementation relies on numext::sqrt, at bottom of file.
305template <typename T>
306struct rsqrt_impl;
307
308// Complex rsqrt defined in MathFunctionsImpl.h.
309template <typename ComplexT>
310EIGEN_DEVICE_FUNC ComplexT complex_rsqrt(const ComplexT& a_x);
311
312template <typename T>
313struct 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
319template <typename Scalar>
320struct rsqrt_retval {
321 typedef Scalar type;
322};
323
324/****************************************************************************
325 * Implementation of norm1 *
326 ****************************************************************************/
327
328template <typename Scalar, bool IsComplex>
329struct norm1_default_impl;
330
331template <typename Scalar>
332struct 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
340template <typename Scalar>
341struct 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
348template <typename Scalar>
349struct norm1_impl : norm1_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
350
351template <typename Scalar>
352struct norm1_retval {
353 typedef typename NumTraits<Scalar>::Real type;
354};
355
356/****************************************************************************
357 * Implementation of hypot *
358 ****************************************************************************/
359
360template <typename Scalar>
361struct hypot_impl;
362
363template <typename Scalar>
364struct hypot_retval {
365 typedef typename NumTraits<Scalar>::Real type;
366};
367
368/****************************************************************************
369 * Implementation of cast *
370 ****************************************************************************/
371
372template <typename OldType, typename NewType, typename EnableIf = void>
373struct cast_impl {
374 EIGEN_DEVICE_FUNC static inline NewType run(const OldType& x) { return static_cast<NewType>(x); }
375};
376
377template <typename OldType>
378struct 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.
384template <typename OldType, typename NewType>
385struct 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
395template <typename OldType, typename NewType>
396EIGEN_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
408template <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>
411struct arg_default_impl;
412
413template <typename Scalar>
414struct 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).
424template <typename Scalar>
425struct 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
432template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
433struct 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
440template <typename Scalar>
441struct 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
449template <typename Scalar>
450struct arg_impl : arg_default_impl<Scalar> {};
451
452template <typename Scalar>
453struct 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.
462namespace 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.
466template <typename Scalar>
467EIGEN_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
487template <typename Scalar>
488struct 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
496template <typename Scalar>
497struct expm1_retval {
498 typedef Scalar type;
499};
500
501/****************************************************************************
502 * Implementation of log *
503 ****************************************************************************/
504
505// Complex log defined in MathFunctionsImpl.h.
506template <typename ComplexT>
507EIGEN_DEVICE_FUNC ComplexT complex_log(const ComplexT& z);
508
509template <typename Scalar>
510struct 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
517template <typename Scalar>
518struct 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
526namespace 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
529template <typename Scalar>
530EIGEN_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
542template <typename Scalar>
543struct 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.
553template <typename RealScalar>
554struct 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
562template <typename Scalar>
563struct log1p_retval {
564 typedef Scalar type;
565};
566
567/****************************************************************************
568 * Implementation of pow *
569 ****************************************************************************/
570
571template <typename ScalarX, typename ScalarY,
572 bool IsInteger = NumTraits<ScalarX>::IsInteger && NumTraits<ScalarY>::IsInteger>
573struct 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
583template <typename ScalarX, typename ScalarY>
584struct 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
600enum { meta_floor_log2_terminate, meta_floor_log2_move_up, meta_floor_log2_move_down, meta_floor_log2_bogus };
601
602template <unsigned int n, int lower, int upper>
603struct 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
613template <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>
615struct meta_floor_log2 {};
616
617template <unsigned int n, int lower, int upper>
618struct 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
622template <unsigned int n, int lower, int upper>
623struct 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
627template <unsigned int n, int lower, int upper>
628struct meta_floor_log2<n, lower, upper, meta_floor_log2_terminate> {
629 enum { value = (n >= ((unsigned int)(1) << (lower + 1))) ? lower + 1 : lower };
630};
631
632template <unsigned int n, int lower, int upper>
633struct meta_floor_log2<n, lower, upper, meta_floor_log2_bogus> {
634 // no value, error at compile time
635};
636
637template <typename BitsType, typename EnableIf = void>
638struct 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.
677template <typename BitsType>
678EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
679 return count_bits_impl<BitsType>::clz(bits);
680}
681
682// Count trailing zeros.
683template <typename BitsType>
684EIGEN_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
690template <typename BitsType>
691struct 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
704template <typename BitsType>
705struct 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
719template <typename BitsType>
720struct 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
736template <typename BitsType>
737struct 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
755template <typename BitsType>
756struct 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
777template <typename BitsType>
778struct 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
791template <typename BitsType>
792int log2_ceil(const BitsType& x) {
793 return log_2_impl<BitsType>::run_ceil(x);
794}
795
796template <typename BitsType>
797int log2_floor(const BitsType& x) {
798 return log_2_impl<BitsType>::run_floor(x);
799}
800
801// Implementation of is* functions
802
803template <typename T>
804EIGEN_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>
807isfinite_impl(const T&) {
808 return true;
809}
810
811template <typename T>
812EIGEN_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>
816isfinite_impl(const T& x) {
817 EIGEN_USING_STD(isfinite);
818 return isfinite EIGEN_NOT_A_MACRO(x);
819}
820
821template <typename T>
822EIGEN_DEVICE_FUNC std::enable_if_t<!std::numeric_limits<T>::has_infinity, bool> isinf_impl(const T&) {
823 return false;
824}
825
826template <typename T>
827EIGEN_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
833template <typename T>
834EIGEN_DEVICE_FUNC
835std::enable_if_t<!(std::numeric_limits<T>::has_quiet_NaN || std::numeric_limits<T>::has_signaling_NaN), bool>
836isnan_impl(const T&) {
837 return false;
838}
839
840template <typename T>
841EIGEN_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>
844isnan_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
850template <typename T>
851EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x);
852template <typename T>
853EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x);
854template <typename T>
855EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x);
856template <typename T>
857EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS T ptanh_float(const T& a_x);
858
859/****************************************************************************
860 * Implementation of sign *
861 ****************************************************************************/
862template <typename Scalar, bool IsComplex = (NumTraits<Scalar>::IsComplex != 0),
863 bool IsInteger = (NumTraits<Scalar>::IsInteger != 0)>
864struct sign_impl {
865 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) { return Scalar((a > Scalar(0)) - (a < Scalar(0))); }
866};
867
868template <typename Scalar>
869struct 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
875template <typename Scalar, bool IsInteger>
876struct 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.
888template <>
889struct sign_impl<bool, false, true> {
890 EIGEN_DEVICE_FUNC static inline bool run(const bool& a) { return a; }
891};
892
893template <typename Scalar>
894struct 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
901template <typename Scalar, bool IsInteger = NumTraits<Scalar>::IsInteger>
902struct negate_impl {
903 static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar run(const Scalar& a) { return -a; }
904};
905
906template <typename Scalar>
907struct 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
912template <typename Scalar>
913struct negate_retval {
914 typedef Scalar type;
915};
916
917template <typename Scalar, bool IsInteger = NumTraits<typename unpacket_traits<Scalar>::type>::IsInteger>
918struct 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};
935template <typename Scalar>
936struct 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.
945namespace has_fma_detail {
946
947template <typename T, typename EnableIf = void>
948struct has_fma_impl : public std::false_type {};
949
950using std::fma;
951
952template <typename T>
953struct 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
959template <typename T>
960struct has_fma : public has_fma_detail::has_fma_impl<T> {};
961
962// Default implementation.
963template <typename T, typename Enable = void>
964struct 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.
969template <typename T>
970struct 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)
978template <>
979struct has_fma<float> : public true_type {};
980
981template <>
982struct 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
988template <>
989struct has_fma<double> : public true_type {};
990
991template <>
992struct 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.
1000template <typename Scalar, typename EnableIf = void>
1001struct 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// Use FMA if there is a single CPU instruction.
1008#ifdef EIGEN_VECTORIZE_FMA
1009template <typename Scalar>
1010struct madd_impl<Scalar, std::enable_if_t<has_fma<Scalar>::value>> {
1011 static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run(const Scalar& x, const Scalar& y, const Scalar& z) {
1012 return fma_impl<Scalar>::run(x, y, z);
1013 }
1014};
1015#endif
1016
1017} // end namespace internal
1018
1019/****************************************************************************
1020 * Generic math functions *
1021 ****************************************************************************/
1022
1023namespace numext {
1024
1025#if (!defined(EIGEN_GPUCC) || defined(EIGEN_CONSTEXPR_ARE_DEVICE_FUNC))
1026template <typename T>
1027EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y) {
1028 EIGEN_USING_STD(min)
1029 return min EIGEN_NOT_A_MACRO(x, y);
1030}
1031
1032template <typename T>
1033EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y) {
1034 EIGEN_USING_STD(max)
1035 return max EIGEN_NOT_A_MACRO(x, y);
1036}
1037#else
1038template <typename T>
1039EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y) {
1040 return y < x ? y : x;
1041}
1042template <>
1043EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float mini(const float& x, const float& y) {
1044 return fminf(x, y);
1045}
1046template <>
1047EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double mini(const double& x, const double& y) {
1048 return fmin(x, y);
1049}
1050
1051#ifndef EIGEN_GPU_COMPILE_PHASE
1052template <>
1053EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double mini(const long double& x, const long double& y) {
1054#if defined(EIGEN_HIPCC)
1055 // no "fminl" on HIP yet
1056 return (x < y) ? x : y;
1057#else
1058 return fminl(x, y);
1059#endif
1060}
1061#endif
1062
1063template <typename T>
1064EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y) {
1065 return x < y ? y : x;
1066}
1067template <>
1068EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float maxi(const float& x, const float& y) {
1069 return fmaxf(x, y);
1070}
1071template <>
1072EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double maxi(const double& x, const double& y) {
1073 return fmax(x, y);
1074}
1075#ifndef EIGEN_GPU_COMPILE_PHASE
1076template <>
1077EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double maxi(const long double& x, const long double& y) {
1078#if defined(EIGEN_HIPCC)
1079 // no "fmaxl" on HIP yet
1080 return (x > y) ? x : y;
1081#else
1082 return fmaxl(x, y);
1083#endif
1084}
1085#endif
1086#endif
1087
1088#if defined(SYCL_DEVICE_ONLY)
1089
1090#define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1091 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \
1092 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \
1093 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \
1094 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1095#define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1096 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \
1097 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \
1098 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \
1099 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1100#define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1101 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \
1102 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1103 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \
1104 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1105#define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1106 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \
1107 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1108 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \
1109 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1110#define SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(NAME, FUNC) \
1111 SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1112 SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC)
1113#define SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(NAME, FUNC) \
1114 SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1115 SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC)
1116#define SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(NAME, FUNC) \
1117 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1118 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_double)
1119#define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(NAME, FUNC) \
1120 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1121 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_double)
1122#define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(NAME, FUNC, RET_TYPE) \
1123 SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_float) \
1124 SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_double)
1125
1126#define SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1127 template <> \
1128 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE& x) { \
1129 return cl::sycl::FUNC(x); \
1130 }
1131
1132#define SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, TYPE) SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, TYPE, TYPE)
1133
1134#define SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE1, ARG_TYPE2) \
1135 template <> \
1136 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE1& x, const ARG_TYPE2& y) { \
1137 return cl::sycl::FUNC(x, y); \
1138 }
1139
1140#define SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1141 SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE, ARG_TYPE)
1142
1143#define SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, TYPE) SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, TYPE, TYPE)
1144
1145SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(mini, min)
1146SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(mini, fmin)
1147SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(maxi, max)
1148SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(maxi, fmax)
1149
1150#endif
1151
1152template <typename Scalar>
1153EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(arg, Scalar) arg(const Scalar& x) {
1154 return EIGEN_MATHFUNC_IMPL(arg, Scalar)::run(x);
1155}
1156
1157template <typename Scalar>
1158EIGEN_DEVICE_FUNC inline internal::add_const_on_value_type_t<EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar)> imag_ref(
1159 const Scalar& x) {
1160 return internal::imag_ref_impl<Scalar>::run(x);
1161}
1162
1163template <typename Scalar>
1164EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x) {
1165 return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(x);
1166}
1167
1168template <typename Scalar>
1169EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x) {
1170 return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x);
1171}
1172
1173template <typename Scalar>
1174EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(sign, Scalar) sign(const Scalar& x) {
1175 return EIGEN_MATHFUNC_IMPL(sign, Scalar)::run(x);
1176}
1177
1178template <typename Scalar>
1179EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(negate, Scalar) negate(const Scalar& x) {
1180 return EIGEN_MATHFUNC_IMPL(negate, Scalar)::run(x);
1181}
1182
1183template <typename Scalar>
1184EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x) {
1185 return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x);
1186}
1187
1188EIGEN_DEVICE_FUNC inline bool abs2(bool x) { return x; }
1189
1190template <typename T>
1191EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T absdiff(const T& x, const T& y) {
1192 return x > y ? x - y : y - x;
1193}
1194template <>
1195EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float absdiff(const float& x, const float& y) {
1196 return fabsf(x - y);
1197}
1198template <>
1199EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double absdiff(const double& x, const double& y) {
1200 return fabs(x - y);
1201}
1202
1203// HIP and CUDA do not support long double.
1204#ifndef EIGEN_GPU_COMPILE_PHASE
1205template <>
1206EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double absdiff(const long double& x, const long double& y) {
1207 return fabsl(x - y);
1208}
1209#endif
1210
1211template <typename Scalar>
1212EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x) {
1213 return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x);
1214}
1215
1216template <typename Scalar>
1217EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(hypot, Scalar) hypot(const Scalar& x, const Scalar& y) {
1218 return EIGEN_MATHFUNC_IMPL(hypot, Scalar)::run(x, y);
1219}
1220
1221#if defined(SYCL_DEVICE_ONLY)
1222SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(hypot, hypot)
1223#endif
1224
1225template <typename Scalar>
1226EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(log1p, Scalar) log1p(const Scalar& x) {
1227 return EIGEN_MATHFUNC_IMPL(log1p, Scalar)::run(x);
1228}
1229
1230#if defined(SYCL_DEVICE_ONLY)
1231SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log1p, log1p)
1232#endif
1233
1234#if defined(EIGEN_GPUCC)
1235template <>
1236EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float log1p(const float& x) {
1237 return ::log1pf(x);
1238}
1239
1240template <>
1241EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double log1p(const double& x) {
1242 return ::log1p(x);
1243}
1244#endif
1245
1246template <typename ScalarX, typename ScalarY>
1247EIGEN_DEVICE_FUNC inline typename internal::pow_impl<ScalarX, ScalarY>::result_type pow(const ScalarX& x,
1248 const ScalarY& y) {
1249 return internal::pow_impl<ScalarX, ScalarY>::run(x, y);
1250}
1251
1252#if defined(SYCL_DEVICE_ONLY)
1253SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(pow, pow)
1254#endif
1255
1256template <typename T>
1257EIGEN_DEVICE_FUNC bool(isnan)(const T& x) {
1258 return internal::isnan_impl(x);
1259}
1260template <typename T>
1261EIGEN_DEVICE_FUNC bool(isinf)(const T& x) {
1262 return internal::isinf_impl(x);
1263}
1264template <typename T>
1265EIGEN_DEVICE_FUNC bool(isfinite)(const T& x) {
1266 return internal::isfinite_impl(x);
1267}
1268
1269#if defined(SYCL_DEVICE_ONLY)
1270SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isnan, isnan, bool)
1271SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isinf, isinf, bool)
1272SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isfinite, isfinite, bool)
1273#endif
1274
1275template <typename Scalar>
1276EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar rint(const Scalar& x) {
1277 return internal::nearest_integer_impl<Scalar>::run_rint(x);
1278}
1279
1280template <typename Scalar>
1281EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar round(const Scalar& x) {
1282 return internal::nearest_integer_impl<Scalar>::run_round(x);
1283}
1284
1285template <typename Scalar>
1286EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(floor)(const Scalar& x) {
1287 return internal::nearest_integer_impl<Scalar>::run_floor(x);
1288}
1289
1290template <typename Scalar>
1291EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(ceil)(const Scalar& x) {
1292 return internal::nearest_integer_impl<Scalar>::run_ceil(x);
1293}
1294
1295template <typename Scalar>
1296EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(trunc)(const Scalar& x) {
1297 return internal::nearest_integer_impl<Scalar>::run_trunc(x);
1298}
1299
1300#if defined(SYCL_DEVICE_ONLY)
1301SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(round, round)
1302SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(floor, floor)
1303SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(ceil, ceil)
1304SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(trunc, trunc)
1305#endif
1306
1307#if defined(EIGEN_GPUCC)
1308template <>
1309EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float floor(const float& x) {
1310 return ::floorf(x);
1311}
1312template <>
1313EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double floor(const double& x) {
1314 return ::floor(x);
1315}
1316template <>
1317EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float ceil(const float& x) {
1318 return ::ceilf(x);
1319}
1320template <>
1321EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double ceil(const double& x) {
1322 return ::ceil(x);
1323}
1324template <>
1325EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float trunc(const float& x) {
1326 return ::truncf(x);
1327}
1328template <>
1329EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double trunc(const double& x) {
1330 return ::trunc(x);
1331}
1332#endif
1333
1334// Integer division with rounding up.
1335// T is assumed to be an integer type with a>=0, and b>0
1336template <typename T>
1337EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE constexpr T div_ceil(T a, T b) {
1338 using UnsignedT = typename internal::make_unsigned<T>::type;
1339 EIGEN_STATIC_ASSERT((NumTraits<T>::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES)
1340 // Note: explicitly declaring a and b as non-negative values allows the compiler to use better optimizations
1341 const UnsignedT ua = UnsignedT(a);
1342 const UnsignedT ub = UnsignedT(b);
1343 // Note: This form is used because it cannot overflow.
1344 return ua == 0 ? 0 : (ua - 1) / ub + 1;
1345}
1346
1347// Integer round down to nearest power of b
1348// T is assumed to be an integer type with a>=0, and b>0
1349template <typename T, typename U>
1350EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE constexpr T round_down(T a, U b) {
1351 using UnsignedT = typename internal::make_unsigned<T>::type;
1352 using UnsignedU = typename internal::make_unsigned<U>::type;
1353 EIGEN_STATIC_ASSERT((NumTraits<T>::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES)
1354 EIGEN_STATIC_ASSERT((NumTraits<U>::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES)
1355 // Note: explicitly declaring a and b as non-negative values allows the compiler to use better optimizations
1356 const UnsignedT ua = UnsignedT(a);
1357 const UnsignedU ub = UnsignedU(b);
1358 return ub * (ua / ub);
1359}
1360
1363constexpr int log2(int x) {
1364 unsigned int v(x);
1365 constexpr int table[32] = {0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30,
1366 8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31};
1367 v |= v >> 1;
1368 v |= v >> 2;
1369 v |= v >> 4;
1370 v |= v >> 8;
1371 v |= v >> 16;
1372 return table[(v * 0x07C4ACDDU) >> 27];
1373}
1374
1384template <typename Scalar>
1385EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE EIGEN_MATHFUNC_RETVAL(sqrt, Scalar) sqrt(const Scalar& x) {
1386 return EIGEN_MATHFUNC_IMPL(sqrt, Scalar)::run(x);
1387}
1388
1389// Boolean specialization, avoids implicit float to bool conversion (-Wimplicit-conversion-floating-point-to-bool).
1390template <>
1391EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS EIGEN_DEVICE_FUNC bool sqrt<bool>(const bool& x) {
1392 return x;
1393}
1394
1395#if defined(SYCL_DEVICE_ONLY)
1396SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sqrt, sqrt)
1397#endif
1398
1400template <typename T>
1401EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<!NumTraits<T>::IsComplex, T> cbrt(const T& x) {
1402 EIGEN_USING_STD(cbrt);
1403 return static_cast<T>(cbrt(x));
1404}
1405
1406template <typename T>
1407EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t<NumTraits<T>::IsComplex, T> cbrt(const T& x) {
1408 EIGEN_USING_STD(pow);
1409 return pow(x, typename NumTraits<T>::Real(1.0 / 3.0));
1410}
1411
1413template <typename T>
1414EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T rsqrt(const T& x) {
1415 return internal::rsqrt_impl<T>::run(x);
1416}
1417
1418template <typename T>
1419EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T log(const T& x) {
1420 return internal::log_impl<T>::run(x);
1421}
1422
1423#if defined(SYCL_DEVICE_ONLY)
1424SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log, log)
1425#endif
1426
1427#if defined(EIGEN_GPUCC)
1428template <>
1429EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float log(const float& x) {
1430 return ::logf(x);
1431}
1432
1433template <>
1434EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double log(const double& x) {
1435 return ::log(x);
1436}
1437#endif
1438
1439template <typename T>
1440EIGEN_DEVICE_FUNC
1441EIGEN_ALWAYS_INLINE std::enable_if_t<NumTraits<T>::IsSigned || NumTraits<T>::IsComplex, typename NumTraits<T>::Real>
1442abs(const T& x) {
1443 EIGEN_USING_STD(abs);
1444 return abs(x);
1445}
1446
1447template <typename T>
1448EIGEN_DEVICE_FUNC
1449EIGEN_ALWAYS_INLINE std::enable_if_t<!(NumTraits<T>::IsSigned || NumTraits<T>::IsComplex), typename NumTraits<T>::Real>
1450abs(const T& x) {
1451 return x;
1452}
1453
1454#if defined(SYCL_DEVICE_ONLY)
1455SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(abs, abs)
1456SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(abs, fabs)
1457#endif
1458
1459#if defined(EIGEN_GPUCC)
1460template <>
1461EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float abs(const float& x) {
1462 return ::fabsf(x);
1463}
1464
1465template <>
1466EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double abs(const double& x) {
1467 return ::fabs(x);
1468}
1469
1470template <>
1471EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float abs(const std::complex<float>& x) {
1472 return ::hypotf(x.real(), x.imag());
1473}
1474
1475template <>
1476EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double abs(const std::complex<double>& x) {
1477 return ::hypot(x.real(), x.imag());
1478}
1479#endif
1480
1481template <typename Scalar, bool IsInteger = NumTraits<Scalar>::IsInteger, bool IsSigned = NumTraits<Scalar>::IsSigned>
1482struct signbit_impl;
1483template <typename Scalar>
1484struct signbit_impl<Scalar, false, true> {
1485 static constexpr size_t Size = sizeof(Scalar);
1486 static constexpr size_t Shift = (CHAR_BIT * Size) - 1;
1487 using intSize_t = typename get_integer_by_size<Size>::signed_type;
1488 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static Scalar run(const Scalar& x) {
1489 intSize_t a = bit_cast<intSize_t, Scalar>(x);
1490 a = a >> Shift;
1491 Scalar result = bit_cast<Scalar, intSize_t>(a);
1492 return result;
1493 }
1494};
1495template <typename Scalar>
1496struct signbit_impl<Scalar, true, true> {
1497 static constexpr size_t Size = sizeof(Scalar);
1498 static constexpr size_t Shift = (CHAR_BIT * Size) - 1;
1499 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar run(const Scalar& x) { return x >> Shift; }
1500};
1501template <typename Scalar>
1502struct signbit_impl<Scalar, true, false> {
1503 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar run(const Scalar&) { return Scalar(0); }
1504};
1505template <typename Scalar>
1506EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar signbit(const Scalar& x) {
1507 return signbit_impl<Scalar>::run(x);
1508}
1509
1510template <typename T>
1511EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T exp(const T& x) {
1512 EIGEN_USING_STD(exp);
1513 return exp(x);
1514}
1515
1516// MSVC screws up some edge-cases for std::exp(complex).
1517#ifdef EIGEN_COMP_MSVC
1518template <typename RealScalar>
1519EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<RealScalar> exp(const std::complex<RealScalar>& x) {
1520 EIGEN_USING_STD(exp);
1521 // If z is (x,±∞) (for any finite x), the result is (NaN,NaN) and FE_INVALID is raised.
1522 // If z is (x,NaN) (for any finite x), the result is (NaN,NaN) and FE_INVALID may be raised.
1523 if ((isfinite)(real_ref(x)) && !(isfinite)(imag_ref(x))) {
1524 return std::complex<RealScalar>(NumTraits<RealScalar>::quiet_NaN(), NumTraits<RealScalar>::quiet_NaN());
1525 }
1526 // If z is (+∞,±∞), the result is (±∞,NaN) and FE_INVALID is raised (the sign of the real part is unspecified)
1527 // If z is (+∞,NaN), the result is (±∞,NaN) (the sign of the real part is unspecified)
1528 if ((real_ref(x) == NumTraits<RealScalar>::infinity() && !(isfinite)(imag_ref(x)))) {
1529 return std::complex<RealScalar>(NumTraits<RealScalar>::infinity(), NumTraits<RealScalar>::quiet_NaN());
1530 }
1531 return exp(x);
1532}
1533#endif
1534
1535#if defined(SYCL_DEVICE_ONLY)
1536SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp, exp)
1537#endif
1538
1539#if defined(EIGEN_GPUCC)
1540template <>
1541EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float exp(const float& x) {
1542 return ::expf(x);
1543}
1544
1545template <>
1546EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double exp(const double& x) {
1547 return ::exp(x);
1548}
1549
1550template <>
1551EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<float> exp(const std::complex<float>& x) {
1552 float com = ::expf(x.real());
1553 float res_real = com * ::cosf(x.imag());
1554 float res_imag = com * ::sinf(x.imag());
1555 return std::complex<float>(res_real, res_imag);
1556}
1557
1558template <>
1559EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<double> exp(const std::complex<double>& x) {
1560 double com = ::exp(x.real());
1561 double res_real = com * ::cos(x.imag());
1562 double res_imag = com * ::sin(x.imag());
1563 return std::complex<double>(res_real, res_imag);
1564}
1565#endif
1566
1567template <typename T>
1568EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T exp2(const T& x) {
1569 EIGEN_USING_STD(exp2);
1570 return exp2(x);
1571}
1572
1573// MSVC screws up some edge-cases for std::exp2(complex).
1574#ifdef EIGEN_COMP_MSVC
1575template <typename RealScalar>
1576EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<RealScalar> exp2(const std::complex<RealScalar>& x) {
1577 EIGEN_USING_STD(exp);
1578 // If z is (x,±∞) (for any finite x), the result is (NaN,NaN) and FE_INVALID is raised.
1579 // If z is (x,NaN) (for any finite x), the result is (NaN,NaN) and FE_INVALID may be raised.
1580 if ((isfinite)(real_ref(x)) && !(isfinite)(imag_ref(x))) {
1581 return std::complex<RealScalar>(NumTraits<RealScalar>::quiet_NaN(), NumTraits<RealScalar>::quiet_NaN());
1582 }
1583 // If z is (+∞,±∞), the result is (±∞,NaN) and FE_INVALID is raised (the sign of the real part is unspecified)
1584 // If z is (+∞,NaN), the result is (±∞,NaN) (the sign of the real part is unspecified)
1585 if ((real_ref(x) == NumTraits<RealScalar>::infinity() && !(isfinite)(imag_ref(x)))) {
1586 return std::complex<RealScalar>(NumTraits<RealScalar>::infinity(), NumTraits<RealScalar>::quiet_NaN());
1587 }
1588 return exp2(x);
1589}
1590#endif
1591
1592#if defined(SYCL_DEVICE_ONLY)
1593SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp2, exp2)
1594#endif
1595
1596#if defined(EIGEN_GPUCC)
1597template <>
1598EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float exp2(const float& x) {
1599 return ::exp2f(x);
1600}
1601
1602template <>
1603EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double exp2(const double& x) {
1604 return ::exp2(x);
1605}
1606
1607template <>
1608EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<float> exp2(const std::complex<float>& x) {
1609 float com = ::exp2f(x.real());
1610 float res_real = com * ::cosf(static_cast<float>(EIGEN_LN2) * x.imag());
1611 float res_imag = com * ::sinf(static_cast<float>(EIGEN_LN2) * x.imag());
1612 return std::complex<float>(res_real, res_imag);
1613}
1614
1615template <>
1616EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<double> exp2(const std::complex<double>& x) {
1617 double com = ::exp2(x.real());
1618 double res_real = com * ::cos(static_cast<double>(EIGEN_LN2) * x.imag());
1619 double res_imag = com * ::sin(static_cast<double>(EIGEN_LN2) * x.imag());
1620 return std::complex<double>(res_real, res_imag);
1621}
1622#endif
1623
1624template <typename Scalar>
1625EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(expm1, Scalar) expm1(const Scalar& x) {
1626 return EIGEN_MATHFUNC_IMPL(expm1, Scalar)::run(x);
1627}
1628
1629#if defined(SYCL_DEVICE_ONLY)
1630SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(expm1, expm1)
1631#endif
1632
1633#if defined(EIGEN_GPUCC)
1634template <>
1635EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float expm1(const float& x) {
1636 return ::expm1f(x);
1637}
1638
1639template <>
1640EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double expm1(const double& x) {
1641 return ::expm1(x);
1642}
1643#endif
1644
1645template <typename T>
1646EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T cos(const T& x) {
1647 EIGEN_USING_STD(cos);
1648 return cos(x);
1649}
1650
1651#if defined(SYCL_DEVICE_ONLY)
1652SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cos, cos)
1653#endif
1654
1655#if defined(EIGEN_GPUCC)
1656template <>
1657EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float cos(const float& x) {
1658 return ::cosf(x);
1659}
1660
1661template <>
1662EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double cos(const double& x) {
1663 return ::cos(x);
1664}
1665#endif
1666
1667template <typename T>
1668EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T sin(const T& x) {
1669 EIGEN_USING_STD(sin);
1670 return sin(x);
1671}
1672
1673#if defined(SYCL_DEVICE_ONLY)
1674SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sin, sin)
1675#endif
1676
1677#if defined(EIGEN_GPUCC)
1678template <>
1679EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float sin(const float& x) {
1680 return ::sinf(x);
1681}
1682
1683template <>
1684EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double sin(const double& x) {
1685 return ::sin(x);
1686}
1687#endif
1688
1689template <typename T>
1690EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T tan(const T& x) {
1691 EIGEN_USING_STD(tan);
1692 return tan(x);
1693}
1694
1695#if defined(SYCL_DEVICE_ONLY)
1696SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tan, tan)
1697#endif
1698
1699#if defined(EIGEN_GPUCC)
1700template <>
1701EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tan(const float& x) {
1702 return ::tanf(x);
1703}
1704
1705template <>
1706EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double tan(const double& x) {
1707 return ::tan(x);
1708}
1709#endif
1710
1711template <typename T>
1712EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acos(const T& x) {
1713 EIGEN_USING_STD(acos);
1714 return acos(x);
1715}
1716
1717template <typename T>
1718EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acosh(const T& x) {
1719 EIGEN_USING_STD(acosh);
1720 return static_cast<T>(acosh(x));
1721}
1722
1723#if defined(SYCL_DEVICE_ONLY)
1724SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acos, acos)
1725SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acosh, acosh)
1726#endif
1727
1728#if defined(EIGEN_GPUCC)
1729template <>
1730EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float acos(const float& x) {
1731 return ::acosf(x);
1732}
1733
1734template <>
1735EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double acos(const double& x) {
1736 return ::acos(x);
1737}
1738#endif
1739
1740template <typename T>
1741EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asin(const T& x) {
1742 EIGEN_USING_STD(asin);
1743 return asin(x);
1744}
1745
1746template <typename T>
1747EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asinh(const T& x) {
1748 EIGEN_USING_STD(asinh);
1749 return static_cast<T>(asinh(x));
1750}
1751
1752#if defined(SYCL_DEVICE_ONLY)
1753SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asin, asin)
1754SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asinh, asinh)
1755#endif
1756
1757#if defined(EIGEN_GPUCC)
1758template <>
1759EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float asin(const float& x) {
1760 return ::asinf(x);
1761}
1762
1763template <>
1764EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double asin(const double& x) {
1765 return ::asin(x);
1766}
1767#endif
1768
1769template <typename T>
1770EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atan(const T& x) {
1771 EIGEN_USING_STD(atan);
1772 return static_cast<T>(atan(x));
1773}
1774
1775template <typename T, std::enable_if_t<!NumTraits<T>::IsComplex, int> = 0>
1776EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atan2(const T& y, const T& x) {
1777 EIGEN_USING_STD(atan2);
1778 return static_cast<T>(atan2(y, x));
1779}
1780
1781template <typename T>
1782EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atanh(const T& x) {
1783 EIGEN_USING_STD(atanh);
1784 return static_cast<T>(atanh(x));
1785}
1786
1787#if defined(SYCL_DEVICE_ONLY)
1788SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atan, atan)
1789SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atanh, atanh)
1790#endif
1791
1792#if defined(EIGEN_GPUCC)
1793template <>
1794EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float atan(const float& x) {
1795 return ::atanf(x);
1796}
1797
1798template <>
1799EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double atan(const double& x) {
1800 return ::atan(x);
1801}
1802#endif
1803
1804template <typename T>
1805EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T cosh(const T& x) {
1806 EIGEN_USING_STD(cosh);
1807 return static_cast<T>(cosh(x));
1808}
1809
1810#if defined(SYCL_DEVICE_ONLY)
1811SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cosh, cosh)
1812#endif
1813
1814#if defined(EIGEN_GPUCC)
1815template <>
1816EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float cosh(const float& x) {
1817 return ::coshf(x);
1818}
1819
1820template <>
1821EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double cosh(const double& x) {
1822 return ::cosh(x);
1823}
1824#endif
1825
1826template <typename T>
1827EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T sinh(const T& x) {
1828 EIGEN_USING_STD(sinh);
1829 return static_cast<T>(sinh(x));
1830}
1831
1832#if defined(SYCL_DEVICE_ONLY)
1833SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sinh, sinh)
1834#endif
1835
1836#if defined(EIGEN_GPUCC)
1837template <>
1838EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float sinh(const float& x) {
1839 return ::sinhf(x);
1840}
1841
1842template <>
1843EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double sinh(const double& x) {
1844 return ::sinh(x);
1845}
1846#endif
1847
1848template <typename T>
1849EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T tanh(const T& x) {
1850 EIGEN_USING_STD(tanh);
1851 return tanh(x);
1852}
1853
1854#if (!defined(EIGEN_GPUCC)) && EIGEN_FAST_MATH && !defined(SYCL_DEVICE_ONLY)
1855EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tanh(float x) { return internal::ptanh_float(x); }
1856#endif
1857
1858#if defined(SYCL_DEVICE_ONLY)
1859SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tanh, tanh)
1860#endif
1861
1862#if defined(EIGEN_GPUCC)
1863template <>
1864EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tanh(const float& x) {
1865 return ::tanhf(x);
1866}
1867
1868template <>
1869EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double tanh(const double& x) {
1870 return ::tanh(x);
1871}
1872#endif
1873
1874template <typename T>
1875EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T fmod(const T& a, const T& b) {
1876 EIGEN_USING_STD(fmod);
1877 return fmod(a, b);
1878}
1879
1880#if defined(SYCL_DEVICE_ONLY)
1881SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(fmod, fmod)
1882#endif
1883
1884#if defined(EIGEN_GPUCC)
1885template <>
1886EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float fmod(const float& a, const float& b) {
1887 return ::fmodf(a, b);
1888}
1889
1890template <>
1891EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double fmod(const double& a, const double& b) {
1892 return ::fmod(a, b);
1893}
1894#endif
1895
1896#if defined(SYCL_DEVICE_ONLY)
1897#undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY
1898#undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY
1899#undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY
1900#undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1901#undef SYCL_SPECIALIZE_INTEGER_TYPES_BINARY
1902#undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1903#undef SYCL_SPECIALIZE_FLOATING_TYPES_BINARY
1904#undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY
1905#undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE
1906#undef SYCL_SPECIALIZE_GEN_UNARY_FUNC
1907#undef SYCL_SPECIALIZE_UNARY_FUNC
1908#undef SYCL_SPECIALIZE_GEN1_BINARY_FUNC
1909#undef SYCL_SPECIALIZE_GEN2_BINARY_FUNC
1910#undef SYCL_SPECIALIZE_BINARY_FUNC
1911#endif
1912
1913template <typename Scalar, typename Enable = std::enable_if_t<std::is_integral<Scalar>::value>>
1914EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar logical_shift_left(const Scalar& a, int n) {
1915 return a << n;
1916}
1917
1918template <typename Scalar, typename Enable = std::enable_if_t<std::is_integral<Scalar>::value>>
1919EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar logical_shift_right(const Scalar& a, int n) {
1920 using UnsignedScalar = typename numext::get_integer_by_size<sizeof(Scalar)>::unsigned_type;
1921 return bit_cast<Scalar, UnsignedScalar>(bit_cast<UnsignedScalar, Scalar>(a) >> n);
1922}
1923
1924template <typename Scalar, typename Enable = std::enable_if_t<std::is_integral<Scalar>::value>>
1925EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar arithmetic_shift_right(const Scalar& a, int n) {
1926 using SignedScalar = typename numext::get_integer_by_size<sizeof(Scalar)>::signed_type;
1927 return bit_cast<Scalar, SignedScalar>(bit_cast<SignedScalar, Scalar>(a) >> n);
1928}
1929
1930// Otherwise, rely on template implementation.
1931template <typename Scalar>
1932EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar fma(const Scalar& x, const Scalar& y, const Scalar& z) {
1933 return internal::fma_impl<Scalar>::run(x, y, z);
1934}
1935
1936// Multiply-add.
1937template <typename Scalar>
1938EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar madd(const Scalar& x, const Scalar& y, const Scalar& z) {
1939 return internal::madd_impl<Scalar>::run(x, y, z);
1940}
1941
1942} // end namespace numext
1943
1944namespace internal {
1945
1946template <typename T>
1947EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x) {
1948 return (numext::isfinite)(numext::real(x)) && (numext::isfinite)(numext::imag(x));
1949}
1950
1951template <typename T>
1952EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x) {
1953 return (numext::isnan)(numext::real(x)) || (numext::isnan)(numext::imag(x));
1954}
1955
1956template <typename T>
1957EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x) {
1958 return ((numext::isinf)(numext::real(x)) || (numext::isinf)(numext::imag(x))) && (!(numext::isnan)(x));
1959}
1960
1961/****************************************************************************
1962 * Implementation of fuzzy comparisons *
1963 ****************************************************************************/
1964
1965template <typename Scalar, bool IsComplex, bool IsInteger>
1966struct scalar_fuzzy_default_impl {};
1967
1968template <typename Scalar>
1969struct scalar_fuzzy_default_impl<Scalar, false, false> {
1970 typedef typename NumTraits<Scalar>::Real RealScalar;
1971 template <typename OtherScalar>
1972 EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
1973 const RealScalar& prec) {
1974 return numext::abs(x) <= numext::abs(y) * prec;
1975 }
1976 EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) {
1977 return numext::abs(x - y) <= numext::mini(numext::abs(x), numext::abs(y)) * prec;
1978 }
1979 EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar& prec) {
1980 return x <= y || isApprox(x, y, prec);
1981 }
1982};
1983
1984template <typename Scalar>
1985struct scalar_fuzzy_default_impl<Scalar, false, true> {
1986 typedef typename NumTraits<Scalar>::Real RealScalar;
1987 template <typename OtherScalar>
1988 EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const Scalar&, const RealScalar&) {
1989 return x == Scalar(0);
1990 }
1991 EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar&) { return x == y; }
1992 EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar&) {
1993 return x <= y;
1994 }
1995};
1996
1997template <typename Scalar>
1998struct scalar_fuzzy_default_impl<Scalar, true, false> {
1999 typedef typename NumTraits<Scalar>::Real RealScalar;
2000 template <typename OtherScalar>
2001 EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
2002 const RealScalar& prec) {
2003 return numext::abs2(x) <= numext::abs2(y) * prec * prec;
2004 }
2005 EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) {
2006 return numext::abs2(x - y) <= numext::mini(numext::abs2(x), numext::abs2(y)) * prec * prec;
2007 }
2008};
2009
2010template <typename Scalar>
2011struct scalar_fuzzy_impl
2012 : scalar_fuzzy_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {};
2013
2014template <typename Scalar, typename OtherScalar>
2015EIGEN_DEVICE_FUNC inline bool isMuchSmallerThan(
2016 const Scalar& x, const OtherScalar& y,
2017 const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
2018 return scalar_fuzzy_impl<Scalar>::template isMuchSmallerThan<OtherScalar>(x, y, precision);
2019}
2020
2021template <typename Scalar>
2022EIGEN_DEVICE_FUNC inline bool isApprox(
2023 const Scalar& x, const Scalar& y,
2024 const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
2025 return scalar_fuzzy_impl<Scalar>::isApprox(x, y, precision);
2026}
2027
2028template <typename Scalar>
2029EIGEN_DEVICE_FUNC inline bool isApproxOrLessThan(
2030 const Scalar& x, const Scalar& y,
2031 const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
2032 return scalar_fuzzy_impl<Scalar>::isApproxOrLessThan(x, y, precision);
2033}
2034
2035/******************************************
2036*** The special case of the bool type ***
2037******************************************/
2038
2039template <>
2040struct scalar_fuzzy_impl<bool> {
2041 typedef bool RealScalar;
2042
2043 template <typename OtherScalar>
2044 EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const bool& x, const bool&, const bool&) {
2045 return !x;
2046 }
2047
2048 EIGEN_DEVICE_FUNC static inline bool isApprox(bool x, bool y, bool) { return x == y; }
2049
2050 EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const bool& x, const bool& y, const bool&) {
2051 return (!x) || y;
2052 }
2053};
2054
2055} // end namespace internal
2056
2057// Default implementations that rely on other numext implementations
2058namespace internal {
2059
2060// Specialization for complex types that are not supported by std::expm1.
2061template <typename RealScalar>
2062struct expm1_impl<std::complex<RealScalar>> {
2063 EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
2064
2065 EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(const std::complex<RealScalar>& x) {
2066 RealScalar xr = x.real();
2067 RealScalar xi = x.imag();
2068 // expm1(z) = exp(z) - 1
2069 // = exp(x + i * y) - 1
2070 // = exp(x) * (cos(y) + i * sin(y)) - 1
2071 // = exp(x) * cos(y) - 1 + i * exp(x) * sin(y)
2072 // Imag(expm1(z)) = exp(x) * sin(y)
2073 // Real(expm1(z)) = exp(x) * cos(y) - 1
2074 // = exp(x) * cos(y) - 1.
2075 // = expm1(x) + exp(x) * (cos(y) - 1)
2076 // = expm1(x) + exp(x) * (2 * sin(y / 2) ** 2)
2077 RealScalar erm1 = numext::expm1<RealScalar>(xr);
2078 RealScalar er = erm1 + RealScalar(1.);
2079 RealScalar sin2 = numext::sin(xi / RealScalar(2.));
2080 sin2 = sin2 * sin2;
2081 RealScalar s = numext::sin(xi);
2082 RealScalar real_part = erm1 - RealScalar(2.) * er * sin2;
2083 return std::complex<RealScalar>(real_part, er * s);
2084 }
2085};
2086
2087template <typename T>
2088struct rsqrt_impl {
2089 EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE T run(const T& x) { return T(1) / numext::sqrt(x); }
2090};
2091
2092#if defined(EIGEN_GPU_COMPILE_PHASE)
2093template <typename T>
2094struct conj_impl<std::complex<T>, true> {
2095 EIGEN_DEVICE_FUNC static inline std::complex<T> run(const std::complex<T>& x) {
2096 return std::complex<T>(numext::real(x), -numext::imag(x));
2097 }
2098};
2099#endif
2100
2101} // end namespace internal
2102
2103} // end namespace Eigen
2104
2105#endif // EIGEN_MATHFUNCTIONS_H
Eigen::ArrayBase::pow
const GlobalUnaryPowReturnType< Derived, ScalarExponent > pow(const Eigen::ArrayBase< Derived > &x, const ScalarExponent &exponent)
Eigen
Namespace containing all symbols from the Eigen library.
Definition B01_Experimental.dox:1
Eigen::sqrt
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_sqrt_op< typename Derived::Scalar >, const Derived > sqrt(const Eigen::ArrayBase< Derived > &x)
Eigen::real
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_real_op< typename Derived::Scalar >, const Derived > real(const Eigen::ArrayBase< Derived > &x)
Eigen::expm1
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_expm1_op< typename Derived::Scalar >, const Derived > expm1(const Eigen::ArrayBase< Derived > &x)
Eigen::exp
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_exp_op< typename Derived::Scalar >, const Derived > exp(const Eigen::ArrayBase< Derived > &x)
Eigen::conj
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_conjugate_op< typename Derived::Scalar >, const Derived > conj(const Eigen::ArrayBase< Derived > &x)
Eigen::rsqrt
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_rsqrt_op< typename Derived::Scalar >, const Derived > rsqrt(const Eigen::ArrayBase< Derived > &x)
Eigen::abs
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_abs_op< typename Derived::Scalar >, const Derived > abs(const Eigen::ArrayBase< Derived > &x)
Eigen::log1p
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log1p_op< typename Derived::Scalar >, const Derived > log1p(const Eigen::ArrayBase< Derived > &x)
Eigen::imag
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_imag_op< typename Derived::Scalar >, const Derived > imag(const Eigen::ArrayBase< Derived > &x)
Eigen::log
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_log_op< typename Derived::Scalar >, const Derived > log(const Eigen::ArrayBase< Derived > &x)
Eigen::arg
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_arg_op< typename Derived::Scalar >, const Derived > arg(const Eigen::ArrayBase< Derived > &x)
Eigen::cbrt
const Eigen::CwiseUnaryOp< Eigen::internal::scalar_cbrt_op< typename Derived::Scalar >, const Derived > cbrt(const Eigen::ArrayBase< Derived > &x)