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