1 // This file is part of Eigen, a lightweight C++ template library
2 // for linear algebra.
3 //
4 // Copyright (C) 2006-2010 Benoit Jacob <[email protected]>
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 namespace 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
abs(long x)25 long abs(long x) { return (labs(x)); }
abs(double x)26 double abs(double x) { return (fabs(x)); }
abs(float x)27 float abs(float x) { return (fabsf(x)); }
abs(long double x)28 long double abs(long double x) { return (fabsl(x)); }
29 #endif
30
31 namespace internal {
32
33 /** \internal \class global_math_functions_filtering_base
34 *
35 * What it does:
36 * Defines a typedef 'type' as follows:
37 * - if type T has a member typedef Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl, then
38 * global_math_functions_filtering_base<T>::type is a typedef for it.
39 * - otherwise, global_math_functions_filtering_base<T>::type is a typedef for T.
40 *
41 * How it's used:
42 * To allow to defined the global math functions (like sin...) in certain cases, like the Array expressions.
43 * When you do sin(array1+array2), the object array1+array2 has a complicated expression type, all what you want to know
44 * is that it inherits ArrayBase. So we implement a partial specialization of sin_impl for ArrayBase<Derived>.
45 * So we must make sure to use sin_impl<ArrayBase<Derived> > and not sin_impl<Derived>, otherwise our partial specialization
46 * won't be used. How does sin know that? That's exactly what global_math_functions_filtering_base tells it.
47 *
48 * How it's implemented:
49 * SFINAE in the style of enable_if. Highly susceptible of breaking compilers. With GCC, it sure does work, but if you replace
50 * the typename dummy by an integer template parameter, it doesn't work anymore!
51 */
52
53 template<typename T, typename dummy = void>
54 struct global_math_functions_filtering_base
55 {
56 typedef T type;
57 };
58
59 template<typename T> struct always_void { typedef void type; };
60
61 template<typename T>
62 struct 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
77 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
78 struct 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
88 template<typename Scalar>
89 struct 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
100 template<typename Scalar> struct real_impl : real_default_impl<Scalar> {};
101
102 #if defined(EIGEN_GPU_COMPILE_PHASE)
103 template<typename T>
104 struct 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
115 template<typename Scalar>
116 struct real_retval
117 {
118 typedef typename NumTraits<Scalar>::Real type;
119 };
120
121 /****************************************************************************
122 * Implementation of imag *
123 ****************************************************************************/
124
125 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
126 struct 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
136 template<typename Scalar>
137 struct 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
148 template<typename Scalar> struct imag_impl : imag_default_impl<Scalar> {};
149
150 #if defined(EIGEN_GPU_COMPILE_PHASE)
151 template<typename T>
152 struct 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
163 template<typename Scalar>
164 struct imag_retval
165 {
166 typedef typename NumTraits<Scalar>::Real type;
167 };
168
169 /****************************************************************************
170 * Implementation of real_ref *
171 ****************************************************************************/
172
173 template<typename Scalar>
174 struct 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
189 template<typename Scalar>
190 struct real_ref_retval
191 {
192 typedef typename NumTraits<Scalar>::Real & type;
193 };
194
195 /****************************************************************************
196 * Implementation of imag_ref *
197 ****************************************************************************/
198
199 template<typename Scalar, bool IsComplex>
200 struct 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
215 template<typename Scalar>
216 struct 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
230 template<typename Scalar>
231 struct imag_ref_impl : imag_ref_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
232
233 template<typename Scalar>
234 struct imag_ref_retval
235 {
236 typedef typename NumTraits<Scalar>::Real & type;
237 };
238
239 /****************************************************************************
240 * Implementation of conj *
241 ****************************************************************************/
242
243 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
244 struct conj_default_impl
245 {
246 EIGEN_DEVICE_FUNC
247 static inline Scalar run(const Scalar& x)
248 {
249 return x;
250 }
251 };
252
253 template<typename Scalar>
254 struct 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
264 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
265 struct conj_impl : conj_default_impl<Scalar, IsComplex> {};
266
267 template<typename Scalar>
268 struct conj_retval
269 {
270 typedef Scalar type;
271 };
272
273 /****************************************************************************
274 * Implementation of abs2 *
275 ****************************************************************************/
276
277 template<typename Scalar,bool IsComplex>
278 struct 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
288 template<typename Scalar>
289 struct 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
299 template<typename Scalar>
300 struct 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
310 template<typename Scalar>
311 struct abs2_retval
312 {
313 typedef typename NumTraits<Scalar>::Real type;
314 };
315
316 /****************************************************************************
317 * Implementation of sqrt/rsqrt *
318 ****************************************************************************/
319
320 template<typename Scalar>
321 struct 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.
332 template<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).
336 template<typename T>
337 struct 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
346 template<typename Scalar>
347 struct sqrt_retval
348 {
349 typedef Scalar type;
350 };
351
352 // Default implementation relies on numext::sqrt, at bottom of file.
353 template<typename T>
354 struct rsqrt_impl;
355
356 // Complex rsqrt defined in MathFunctionsImpl.h.
357 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_rsqrt(const std::complex<T>& a_x);
358
359 template<typename T>
360 struct 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
369 template<typename Scalar>
370 struct rsqrt_retval
371 {
372 typedef Scalar type;
373 };
374
375 /****************************************************************************
376 * Implementation of norm1 *
377 ****************************************************************************/
378
379 template<typename Scalar, bool IsComplex>
380 struct norm1_default_impl;
381
382 template<typename Scalar>
383 struct 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
394 template<typename Scalar>
395 struct 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
405 template<typename Scalar>
406 struct norm1_impl : norm1_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
407
408 template<typename Scalar>
409 struct norm1_retval
410 {
411 typedef typename NumTraits<Scalar>::Real type;
412 };
413
414 /****************************************************************************
415 * Implementation of hypot *
416 ****************************************************************************/
417
418 template<typename Scalar> struct hypot_impl;
419
420 template<typename Scalar>
421 struct hypot_retval
422 {
423 typedef typename NumTraits<Scalar>::Real type;
424 };
425
426 /****************************************************************************
427 * Implementation of cast *
428 ****************************************************************************/
429
430 template<typename OldType, typename NewType, typename EnableIf = void>
431 struct 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.
442 template<typename OldType, typename NewType>
443 struct 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
458 template<typename OldType, typename NewType>
459 EIGEN_DEVICE_FUNC
460 inline NewType cast(const OldType& x)
461 {
462 return cast_impl<OldType, NewType>::run(x);
463 }
464
465 /****************************************************************************
466 * Implementation of round *
467 ****************************************************************************/
468
469 template<typename Scalar>
470 struct 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.
486 template<>
487 struct round_impl<float> {
488 EIGEN_DEVICE_FUNC
489 static inline float run(const float& x)
490 {
491 return ::roundf(x);
492 }
493 };
494 #else
495 template<typename Scalar>
496 struct 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
515 template<>
516 struct round_impl<float> : round_using_floor_ceil_impl<float> {};
517
518 template<>
519 struct round_impl<double> : round_using_floor_ceil_impl<double> {};
520 #endif // EIGEN_HAS_C99_MATH
521 #endif // !EIGEN_HAS_CXX11_MATH
522
523 template<typename Scalar>
524 struct round_retval
525 {
526 typedef Scalar type;
527 };
528
529 /****************************************************************************
530 * Implementation of rint *
531 ****************************************************************************/
532
533 template<typename Scalar>
534 struct 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
547 template<>
548 struct rint_impl<double> {
549 EIGEN_DEVICE_FUNC
550 static inline double run(const double& x)
551 {
552 return ::rint(x);
553 }
554 };
555 template<>
556 struct 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
565 template<typename Scalar>
566 struct 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
579 template<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 >
583 struct arg_default_impl;
584
585 template<typename Scalar>
586 struct 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 #if defined(EIGEN_HIP_DEVICE_COMPILE)
592 // HIP does not seem to have a native device side implementation for the math routine "arg"
593 using std::arg;
594 #else
595 EIGEN_USING_STD(arg);
596 #endif
597 return static_cast<RealScalar>(arg(x));
598 }
599 };
600
601 // Must be non-complex floating-point type (e.g. half/bfloat16).
602 template<typename Scalar>
603 struct arg_default_impl<Scalar, false> {
604 typedef typename NumTraits<Scalar>::Real RealScalar;
605 EIGEN_DEVICE_FUNC
606 static inline RealScalar run(const Scalar& x)
607 {
608 return (x < Scalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
609 }
610 };
611 #else
612 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
613 struct arg_default_impl
614 {
615 typedef typename NumTraits<Scalar>::Real RealScalar;
616 EIGEN_DEVICE_FUNC
617 static inline RealScalar run(const Scalar& x)
618 {
619 return (x < RealScalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
620 }
621 };
622
623 template<typename Scalar>
624 struct arg_default_impl<Scalar,true>
625 {
626 typedef typename NumTraits<Scalar>::Real RealScalar;
627 EIGEN_DEVICE_FUNC
628 static inline RealScalar run(const Scalar& x)
629 {
630 EIGEN_USING_STD(arg);
631 return arg(x);
632 }
633 };
634 #endif
635 template<typename Scalar> struct arg_impl : arg_default_impl<Scalar> {};
636
637 template<typename Scalar>
638 struct arg_retval
639 {
640 typedef typename NumTraits<Scalar>::Real type;
641 };
642
643 /****************************************************************************
644 * Implementation of expm1 *
645 ****************************************************************************/
646
647 // This implementation is based on GSL Math's expm1.
648 namespace std_fallback {
649 // fallback expm1 implementation in case there is no expm1(Scalar) function in namespace of Scalar,
650 // or that there is no suitable std::expm1 function available. Implementation
651 // attributed to Kahan. See: http://www.plunk.org/~hatch/rightway.php.
652 template<typename Scalar>
653 EIGEN_DEVICE_FUNC inline Scalar expm1(const Scalar& x) {
654 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
655 typedef typename NumTraits<Scalar>::Real RealScalar;
656
657 EIGEN_USING_STD(exp);
658 Scalar u = exp(x);
659 if (numext::equal_strict(u, Scalar(1))) {
660 return x;
661 }
662 Scalar um1 = u - RealScalar(1);
663 if (numext::equal_strict(um1, Scalar(-1))) {
664 return RealScalar(-1);
665 }
666
667 EIGEN_USING_STD(log);
668 Scalar logu = log(u);
669 return numext::equal_strict(u, logu) ? u : (u - RealScalar(1)) * x / logu;
670 }
671 }
672
673 template<typename Scalar>
674 struct expm1_impl {
675 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x)
676 {
677 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
678 #if EIGEN_HAS_CXX11_MATH
679 using std::expm1;
680 #else
681 using std_fallback::expm1;
682 #endif
683 return expm1(x);
684 }
685 };
686
687 template<typename Scalar>
688 struct expm1_retval
689 {
690 typedef Scalar type;
691 };
692
693 /****************************************************************************
694 * Implementation of log *
695 ****************************************************************************/
696
697 // Complex log defined in MathFunctionsImpl.h.
698 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_log(const std::complex<T>& z);
699
700 template<typename Scalar>
701 struct log_impl {
702 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x)
703 {
704 EIGEN_USING_STD(log);
705 return static_cast<Scalar>(log(x));
706 }
707 };
708
709 template<typename Scalar>
710 struct log_impl<std::complex<Scalar> > {
711 EIGEN_DEVICE_FUNC static inline std::complex<Scalar> run(const std::complex<Scalar>& z)
712 {
713 return complex_log(z);
714 }
715 };
716
717 /****************************************************************************
718 * Implementation of log1p *
719 ****************************************************************************/
720
721 namespace std_fallback {
722 // fallback log1p implementation in case there is no log1p(Scalar) function in namespace of Scalar,
723 // or that there is no suitable std::log1p function available
724 template<typename Scalar>
725 EIGEN_DEVICE_FUNC inline Scalar log1p(const Scalar& x) {
726 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
727 typedef typename NumTraits<Scalar>::Real RealScalar;
728 EIGEN_USING_STD(log);
729 Scalar x1p = RealScalar(1) + x;
730 Scalar log_1p = log_impl<Scalar>::run(x1p);
731 const bool is_small = numext::equal_strict(x1p, Scalar(1));
732 const bool is_inf = numext::equal_strict(x1p, log_1p);
733 return (is_small || is_inf) ? x : x * (log_1p / (x1p - RealScalar(1)));
734 }
735 }
736
737 template<typename Scalar>
738 struct log1p_impl {
739 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x)
740 {
741 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
742 #if EIGEN_HAS_CXX11_MATH
743 using std::log1p;
744 #else
745 using std_fallback::log1p;
746 #endif
747 return log1p(x);
748 }
749 };
750
751 // Specialization for complex types that are not supported by std::log1p.
752 template <typename RealScalar>
753 struct log1p_impl<std::complex<RealScalar> > {
754 EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(
755 const std::complex<RealScalar>& x) {
756 EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
757 return std_fallback::log1p(x);
758 }
759 };
760
761 template<typename Scalar>
762 struct log1p_retval
763 {
764 typedef Scalar type;
765 };
766
767 /****************************************************************************
768 * Implementation of pow *
769 ****************************************************************************/
770
771 template<typename ScalarX,typename ScalarY, bool IsInteger = NumTraits<ScalarX>::IsInteger&&NumTraits<ScalarY>::IsInteger>
772 struct pow_impl
773 {
774 //typedef Scalar retval;
775 typedef typename ScalarBinaryOpTraits<ScalarX,ScalarY,internal::scalar_pow_op<ScalarX,ScalarY> >::ReturnType result_type;
776 static EIGEN_DEVICE_FUNC inline result_type run(const ScalarX& x, const ScalarY& y)
777 {
778 EIGEN_USING_STD(pow);
779 return pow(x, y);
780 }
781 };
782
783 template<typename ScalarX,typename ScalarY>
784 struct pow_impl<ScalarX,ScalarY, true>
785 {
786 typedef ScalarX result_type;
787 static EIGEN_DEVICE_FUNC inline ScalarX run(ScalarX x, ScalarY y)
788 {
789 ScalarX res(1);
790 eigen_assert(!NumTraits<ScalarY>::IsSigned || y >= 0);
791 if(y & 1) res *= x;
792 y >>= 1;
793 while(y)
794 {
795 x *= x;
796 if(y&1) res *= x;
797 y >>= 1;
798 }
799 return res;
800 }
801 };
802
803 /****************************************************************************
804 * Implementation of random *
805 ****************************************************************************/
806
807 template<typename Scalar,
808 bool IsComplex,
809 bool IsInteger>
810 struct random_default_impl {};
811
812 template<typename Scalar>
813 struct random_impl : random_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {};
814
815 template<typename Scalar>
816 struct random_retval
817 {
818 typedef Scalar type;
819 };
820
821 template<typename Scalar> inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(const Scalar& x, const Scalar& y);
822 template<typename Scalar> inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random();
823
824 template<typename Scalar>
825 struct random_default_impl<Scalar, false, false>
826 {
827 static inline Scalar run(const Scalar& x, const Scalar& y)
828 {
829 return x + (y-x) * Scalar(std::rand()) / Scalar(RAND_MAX);
830 }
831 static inline Scalar run()
832 {
833 return run(Scalar(NumTraits<Scalar>::IsSigned ? -1 : 0), Scalar(1));
834 }
835 };
836
837 enum {
838 meta_floor_log2_terminate,
839 meta_floor_log2_move_up,
840 meta_floor_log2_move_down,
841 meta_floor_log2_bogus
842 };
843
844 template<unsigned int n, int lower, int upper> struct meta_floor_log2_selector
845 {
846 enum { middle = (lower + upper) / 2,
847 value = (upper <= lower + 1) ? int(meta_floor_log2_terminate)
848 : (n < (1 << middle)) ? int(meta_floor_log2_move_down)
849 : (n==0) ? int(meta_floor_log2_bogus)
850 : int(meta_floor_log2_move_up)
851 };
852 };
853
854 template<unsigned int n,
855 int lower = 0,
856 int upper = sizeof(unsigned int) * CHAR_BIT - 1,
857 int selector = meta_floor_log2_selector<n, lower, upper>::value>
858 struct meta_floor_log2 {};
859
860 template<unsigned int n, int lower, int upper>
861 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_down>
862 {
863 enum { value = meta_floor_log2<n, lower, meta_floor_log2_selector<n, lower, upper>::middle>::value };
864 };
865
866 template<unsigned int n, int lower, int upper>
867 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_up>
868 {
869 enum { value = meta_floor_log2<n, meta_floor_log2_selector<n, lower, upper>::middle, upper>::value };
870 };
871
872 template<unsigned int n, int lower, int upper>
873 struct meta_floor_log2<n, lower, upper, meta_floor_log2_terminate>
874 {
875 enum { value = (n >= ((unsigned int)(1) << (lower+1))) ? lower+1 : lower };
876 };
877
878 template<unsigned int n, int lower, int upper>
879 struct meta_floor_log2<n, lower, upper, meta_floor_log2_bogus>
880 {
881 // no value, error at compile time
882 };
883
884 template<typename Scalar>
885 struct random_default_impl<Scalar, false, true>
886 {
887 static inline Scalar run(const Scalar& x, const Scalar& y)
888 {
889 if (y <= x)
890 return x;
891 // ScalarU is the unsigned counterpart of Scalar, possibly Scalar itself.
892 typedef typename make_unsigned<Scalar>::type ScalarU;
893 // ScalarX is the widest of ScalarU and unsigned int.
894 // We'll deal only with ScalarX and unsigned int below thus avoiding signed
895 // types and arithmetic and signed overflows (which are undefined behavior).
896 typedef typename conditional<(ScalarU(-1) > unsigned(-1)), ScalarU, unsigned>::type ScalarX;
897 // The following difference doesn't overflow, provided our integer types are two's
898 // complement and have the same number of padding bits in signed and unsigned variants.
899 // This is the case in most modern implementations of C++.
900 ScalarX range = ScalarX(y) - ScalarX(x);
901 ScalarX offset = 0;
902 ScalarX divisor = 1;
903 ScalarX multiplier = 1;
904 const unsigned rand_max = RAND_MAX;
905 if (range <= rand_max) divisor = (rand_max + 1) / (range + 1);
906 else multiplier = 1 + range / (rand_max + 1);
907 // Rejection sampling.
908 do {
909 offset = (unsigned(std::rand()) * multiplier) / divisor;
910 } while (offset > range);
911 return Scalar(ScalarX(x) + offset);
912 }
913
914 static inline Scalar run()
915 {
916 #ifdef EIGEN_MAKING_DOCS
917 return run(Scalar(NumTraits<Scalar>::IsSigned ? -10 : 0), Scalar(10));
918 #else
919 enum { rand_bits = meta_floor_log2<(unsigned int)(RAND_MAX)+1>::value,
920 scalar_bits = sizeof(Scalar) * CHAR_BIT,
921 shift = EIGEN_PLAIN_ENUM_MAX(0, int(rand_bits) - int(scalar_bits)),
922 offset = NumTraits<Scalar>::IsSigned ? (1 << (EIGEN_PLAIN_ENUM_MIN(rand_bits,scalar_bits)-1)) : 0
923 };
924 return Scalar((std::rand() >> shift) - offset);
925 #endif
926 }
927 };
928
929 template<typename Scalar>
930 struct random_default_impl<Scalar, true, false>
931 {
932 static inline Scalar run(const Scalar& x, const Scalar& y)
933 {
934 return Scalar(random(x.real(), y.real()),
935 random(x.imag(), y.imag()));
936 }
937 static inline Scalar run()
938 {
939 typedef typename NumTraits<Scalar>::Real RealScalar;
940 return Scalar(random<RealScalar>(), random<RealScalar>());
941 }
942 };
943
944 template<typename Scalar>
945 inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(const Scalar& x, const Scalar& y)
946 {
947 return EIGEN_MATHFUNC_IMPL(random, Scalar)::run(x, y);
948 }
949
950 template<typename Scalar>
951 inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random()
952 {
953 return EIGEN_MATHFUNC_IMPL(random, Scalar)::run();
954 }
955
956 // Implementation of is* functions
957
958 // std::is* do not work with fast-math and gcc, std::is* are available on MSVC 2013 and newer, as well as in clang.
959 #if (EIGEN_HAS_CXX11_MATH && !(EIGEN_COMP_GNUC_STRICT && __FINITE_MATH_ONLY__)) || (EIGEN_COMP_MSVC>=1800) || (EIGEN_COMP_CLANG)
960 #define EIGEN_USE_STD_FPCLASSIFY 1
961 #else
962 #define EIGEN_USE_STD_FPCLASSIFY 0
963 #endif
964
965 template<typename T>
966 EIGEN_DEVICE_FUNC
967 typename internal::enable_if<internal::is_integral<T>::value,bool>::type
968 isnan_impl(const T&) { return false; }
969
970 template<typename T>
971 EIGEN_DEVICE_FUNC
972 typename internal::enable_if<internal::is_integral<T>::value,bool>::type
973 isinf_impl(const T&) { return false; }
974
975 template<typename T>
976 EIGEN_DEVICE_FUNC
977 typename internal::enable_if<internal::is_integral<T>::value,bool>::type
978 isfinite_impl(const T&) { return true; }
979
980 template<typename T>
981 EIGEN_DEVICE_FUNC
982 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type
983 isfinite_impl(const T& x)
984 {
985 #if defined(EIGEN_GPU_COMPILE_PHASE)
986 return (::isfinite)(x);
987 #elif EIGEN_USE_STD_FPCLASSIFY
988 using std::isfinite;
989 return isfinite EIGEN_NOT_A_MACRO (x);
990 #else
991 return x<=NumTraits<T>::highest() && x>=NumTraits<T>::lowest();
992 #endif
993 }
994
995 template<typename T>
996 EIGEN_DEVICE_FUNC
997 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type
998 isinf_impl(const T& x)
999 {
1000 #if defined(EIGEN_GPU_COMPILE_PHASE)
1001 return (::isinf)(x);
1002 #elif EIGEN_USE_STD_FPCLASSIFY
1003 using std::isinf;
1004 return isinf EIGEN_NOT_A_MACRO (x);
1005 #else
1006 return x>NumTraits<T>::highest() || x<NumTraits<T>::lowest();
1007 #endif
1008 }
1009
1010 template<typename T>
1011 EIGEN_DEVICE_FUNC
1012 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type
1013 isnan_impl(const T& x)
1014 {
1015 #if defined(EIGEN_GPU_COMPILE_PHASE)
1016 return (::isnan)(x);
1017 #elif EIGEN_USE_STD_FPCLASSIFY
1018 using std::isnan;
1019 return isnan EIGEN_NOT_A_MACRO (x);
1020 #else
1021 return x != x;
1022 #endif
1023 }
1024
1025 #if (!EIGEN_USE_STD_FPCLASSIFY)
1026
1027 #if EIGEN_COMP_MSVC
1028
1029 template<typename T> EIGEN_DEVICE_FUNC bool isinf_msvc_helper(T x)
1030 {
1031 return _fpclass(x)==_FPCLASS_NINF || _fpclass(x)==_FPCLASS_PINF;
1032 }
1033
1034 //MSVC defines a _isnan builtin function, but for double only
1035 EIGEN_DEVICE_FUNC inline bool isnan_impl(const long double& x) { return _isnan(x)!=0; }
1036 EIGEN_DEVICE_FUNC inline bool isnan_impl(const double& x) { return _isnan(x)!=0; }
1037 EIGEN_DEVICE_FUNC inline bool isnan_impl(const float& x) { return _isnan(x)!=0; }
1038
1039 EIGEN_DEVICE_FUNC inline bool isinf_impl(const long double& x) { return isinf_msvc_helper(x); }
1040 EIGEN_DEVICE_FUNC inline bool isinf_impl(const double& x) { return isinf_msvc_helper(x); }
1041 EIGEN_DEVICE_FUNC inline bool isinf_impl(const float& x) { return isinf_msvc_helper(x); }
1042
1043 #elif (defined __FINITE_MATH_ONLY__ && __FINITE_MATH_ONLY__ && EIGEN_COMP_GNUC)
1044
1045 #if EIGEN_GNUC_AT_LEAST(5,0)
1046 #define EIGEN_TMP_NOOPT_ATTRIB EIGEN_DEVICE_FUNC inline __attribute__((optimize("no-finite-math-only")))
1047 #else
1048 // NOTE the inline qualifier and noinline attribute are both needed: the former is to avoid linking issue (duplicate symbol),
1049 // while the second prevent too aggressive optimizations in fast-math mode:
1050 #define EIGEN_TMP_NOOPT_ATTRIB EIGEN_DEVICE_FUNC inline __attribute__((noinline,optimize("no-finite-math-only")))
1051 #endif
1052
1053 template<> EIGEN_TMP_NOOPT_ATTRIB bool isnan_impl(const long double& x) { return __builtin_isnan(x); }
1054 template<> EIGEN_TMP_NOOPT_ATTRIB bool isnan_impl(const double& x) { return __builtin_isnan(x); }
1055 template<> EIGEN_TMP_NOOPT_ATTRIB bool isnan_impl(const float& x) { return __builtin_isnan(x); }
1056 template<> EIGEN_TMP_NOOPT_ATTRIB bool isinf_impl(const double& x) { return __builtin_isinf(x); }
1057 template<> EIGEN_TMP_NOOPT_ATTRIB bool isinf_impl(const float& x) { return __builtin_isinf(x); }
1058 template<> EIGEN_TMP_NOOPT_ATTRIB bool isinf_impl(const long double& x) { return __builtin_isinf(x); }
1059
1060 #undef EIGEN_TMP_NOOPT_ATTRIB
1061
1062 #endif
1063
1064 #endif
1065
1066 // The following overload are defined at the end of this file
1067 template<typename T> EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x);
1068 template<typename T> EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x);
1069 template<typename T> EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x);
1070
1071 template<typename T> T generic_fast_tanh_float(const T& a_x);
1072 } // end namespace internal
1073
1074 /****************************************************************************
1075 * Generic math functions *
1076 ****************************************************************************/
1077
1078 namespace numext {
1079
1080 #if (!defined(EIGEN_GPUCC) || defined(EIGEN_CONSTEXPR_ARE_DEVICE_FUNC))
1081 template<typename T>
1082 EIGEN_DEVICE_FUNC
1083 EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y)
1084 {
1085 EIGEN_USING_STD(min)
1086 return min EIGEN_NOT_A_MACRO (x,y);
1087 }
1088
1089 template<typename T>
1090 EIGEN_DEVICE_FUNC
1091 EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y)
1092 {
1093 EIGEN_USING_STD(max)
1094 return max EIGEN_NOT_A_MACRO (x,y);
1095 }
1096 #else
1097 template<typename T>
1098 EIGEN_DEVICE_FUNC
1099 EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y)
1100 {
1101 return y < x ? y : x;
1102 }
1103 template<>
1104 EIGEN_DEVICE_FUNC
1105 EIGEN_ALWAYS_INLINE float mini(const float& x, const float& y)
1106 {
1107 return fminf(x, y);
1108 }
1109 template<>
1110 EIGEN_DEVICE_FUNC
1111 EIGEN_ALWAYS_INLINE double mini(const double& x, const double& y)
1112 {
1113 return fmin(x, y);
1114 }
1115 template<>
1116 EIGEN_DEVICE_FUNC
1117 EIGEN_ALWAYS_INLINE long double mini(const long double& x, const long double& y)
1118 {
1119 #if defined(EIGEN_HIPCC)
1120 // no "fminl" on HIP yet
1121 return (x < y) ? x : y;
1122 #else
1123 return fminl(x, y);
1124 #endif
1125 }
1126
1127 template<typename T>
1128 EIGEN_DEVICE_FUNC
1129 EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y)
1130 {
1131 return x < y ? y : x;
1132 }
1133 template<>
1134 EIGEN_DEVICE_FUNC
1135 EIGEN_ALWAYS_INLINE float maxi(const float& x, const float& y)
1136 {
1137 return fmaxf(x, y);
1138 }
1139 template<>
1140 EIGEN_DEVICE_FUNC
1141 EIGEN_ALWAYS_INLINE double maxi(const double& x, const double& y)
1142 {
1143 return fmax(x, y);
1144 }
1145 template<>
1146 EIGEN_DEVICE_FUNC
1147 EIGEN_ALWAYS_INLINE long double maxi(const long double& x, const long double& y)
1148 {
1149 #if defined(EIGEN_HIPCC)
1150 // no "fmaxl" on HIP yet
1151 return (x > y) ? x : y;
1152 #else
1153 return fmaxl(x, y);
1154 #endif
1155 }
1156 #endif
1157
1158 #if defined(SYCL_DEVICE_ONLY)
1159
1160
1161 #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1162 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \
1163 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \
1164 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \
1165 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1166 #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1167 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \
1168 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \
1169 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \
1170 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1171 #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1172 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \
1173 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1174 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \
1175 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1176 #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1177 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \
1178 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1179 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \
1180 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1181 #define SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(NAME, FUNC) \
1182 SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1183 SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC)
1184 #define SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(NAME, FUNC) \
1185 SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1186 SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC)
1187 #define SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(NAME, FUNC) \
1188 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1189 SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC,cl::sycl::cl_double)
1190 #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(NAME, FUNC) \
1191 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1192 SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC,cl::sycl::cl_double)
1193 #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(NAME, FUNC, RET_TYPE) \
1194 SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_float) \
1195 SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_double)
1196
1197 #define SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1198 template<> \
1199 EIGEN_DEVICE_FUNC \
1200 EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE& x) { \
1201 return cl::sycl::FUNC(x); \
1202 }
1203
1204 #define SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, TYPE) \
1205 SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, TYPE, TYPE)
1206
1207 #define SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE1, ARG_TYPE2) \
1208 template<> \
1209 EIGEN_DEVICE_FUNC \
1210 EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE1& x, const ARG_TYPE2& y) { \
1211 return cl::sycl::FUNC(x, y); \
1212 }
1213
1214 #define SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1215 SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE, ARG_TYPE)
1216
1217 #define SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, TYPE) \
1218 SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, TYPE, TYPE)
1219
1220 SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(mini, min)
1221 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(mini, fmin)
1222 SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(maxi, max)
1223 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(maxi, fmax)
1224
1225 #endif
1226
1227
1228 template<typename Scalar>
1229 EIGEN_DEVICE_FUNC
1230 inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x)
1231 {
1232 return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x);
1233 }
1234
1235 template<typename Scalar>
1236 EIGEN_DEVICE_FUNC
1237 inline typename internal::add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) >::type real_ref(const Scalar& x)
1238 {
1239 return internal::real_ref_impl<Scalar>::run(x);
1240 }
1241
1242 template<typename Scalar>
1243 EIGEN_DEVICE_FUNC
1244 inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x)
1245 {
1246 return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x);
1247 }
1248
1249 template<typename Scalar>
1250 EIGEN_DEVICE_FUNC
1251 inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x)
1252 {
1253 return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x);
1254 }
1255
1256 template<typename Scalar>
1257 EIGEN_DEVICE_FUNC
1258 inline EIGEN_MATHFUNC_RETVAL(arg, Scalar) arg(const Scalar& x)
1259 {
1260 return EIGEN_MATHFUNC_IMPL(arg, Scalar)::run(x);
1261 }
1262
1263 template<typename Scalar>
1264 EIGEN_DEVICE_FUNC
1265 inline typename internal::add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) >::type imag_ref(const Scalar& x)
1266 {
1267 return internal::imag_ref_impl<Scalar>::run(x);
1268 }
1269
1270 template<typename Scalar>
1271 EIGEN_DEVICE_FUNC
1272 inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x)
1273 {
1274 return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(x);
1275 }
1276
1277 template<typename Scalar>
1278 EIGEN_DEVICE_FUNC
1279 inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x)
1280 {
1281 return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x);
1282 }
1283
1284 template<typename Scalar>
1285 EIGEN_DEVICE_FUNC
1286 inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x)
1287 {
1288 return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x);
1289 }
1290
1291 EIGEN_DEVICE_FUNC
1292 inline bool abs2(bool x) { return x; }
1293
1294 template<typename T>
1295 EIGEN_DEVICE_FUNC
1296 EIGEN_ALWAYS_INLINE T absdiff(const T& x, const T& y)
1297 {
1298 return x > y ? x - y : y - x;
1299 }
1300 template<>
1301 EIGEN_DEVICE_FUNC
1302 EIGEN_ALWAYS_INLINE float absdiff(const float& x, const float& y)
1303 {
1304 return fabsf(x - y);
1305 }
1306 template<>
1307 EIGEN_DEVICE_FUNC
1308 EIGEN_ALWAYS_INLINE double absdiff(const double& x, const double& y)
1309 {
1310 return fabs(x - y);
1311 }
1312
1313 #if !defined(EIGEN_GPUCC)
1314 // HIP and CUDA do not support long double.
1315 template<>
1316 EIGEN_DEVICE_FUNC
1317 EIGEN_ALWAYS_INLINE long double absdiff(const long double& x, const long double& y) {
1318 return fabsl(x - y);
1319 }
1320 #endif
1321
1322 template<typename Scalar>
1323 EIGEN_DEVICE_FUNC
1324 inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x)
1325 {
1326 return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x);
1327 }
1328
1329 template<typename Scalar>
1330 EIGEN_DEVICE_FUNC
1331 inline EIGEN_MATHFUNC_RETVAL(hypot, Scalar) hypot(const Scalar& x, const Scalar& y)
1332 {
1333 return EIGEN_MATHFUNC_IMPL(hypot, Scalar)::run(x, y);
1334 }
1335
1336 #if defined(SYCL_DEVICE_ONLY)
1337 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(hypot, hypot)
1338 #endif
1339
1340 template<typename Scalar>
1341 EIGEN_DEVICE_FUNC
1342 inline EIGEN_MATHFUNC_RETVAL(log1p, Scalar) log1p(const Scalar& x)
1343 {
1344 return EIGEN_MATHFUNC_IMPL(log1p, Scalar)::run(x);
1345 }
1346
1347 #if defined(SYCL_DEVICE_ONLY)
1348 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log1p, log1p)
1349 #endif
1350
1351 #if defined(EIGEN_GPUCC)
1352 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1353 float log1p(const float &x) { return ::log1pf(x); }
1354
1355 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1356 double log1p(const double &x) { return ::log1p(x); }
1357 #endif
1358
1359 template<typename ScalarX,typename ScalarY>
1360 EIGEN_DEVICE_FUNC
1361 inline typename internal::pow_impl<ScalarX,ScalarY>::result_type pow(const ScalarX& x, const ScalarY& y)
1362 {
1363 return internal::pow_impl<ScalarX,ScalarY>::run(x, y);
1364 }
1365
1366 #if defined(SYCL_DEVICE_ONLY)
1367 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(pow, pow)
1368 #endif
1369
1370 template<typename T> EIGEN_DEVICE_FUNC bool (isnan) (const T &x) { return internal::isnan_impl(x); }
1371 template<typename T> EIGEN_DEVICE_FUNC bool (isinf) (const T &x) { return internal::isinf_impl(x); }
1372 template<typename T> EIGEN_DEVICE_FUNC bool (isfinite)(const T &x) { return internal::isfinite_impl(x); }
1373
1374 #if defined(SYCL_DEVICE_ONLY)
1375 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isnan, isnan, bool)
1376 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isinf, isinf, bool)
1377 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isfinite, isfinite, bool)
1378 #endif
1379
1380 template<typename Scalar>
1381 EIGEN_DEVICE_FUNC
1382 inline EIGEN_MATHFUNC_RETVAL(rint, Scalar) rint(const Scalar& x)
1383 {
1384 return EIGEN_MATHFUNC_IMPL(rint, Scalar)::run(x);
1385 }
1386
1387 template<typename Scalar>
1388 EIGEN_DEVICE_FUNC
1389 inline EIGEN_MATHFUNC_RETVAL(round, Scalar) round(const Scalar& x)
1390 {
1391 return EIGEN_MATHFUNC_IMPL(round, Scalar)::run(x);
1392 }
1393
1394 #if defined(SYCL_DEVICE_ONLY)
1395 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(round, round)
1396 #endif
1397
1398 template<typename T>
1399 EIGEN_DEVICE_FUNC
1400 T (floor)(const T& x)
1401 {
1402 EIGEN_USING_STD(floor)
1403 return floor(x);
1404 }
1405
1406 #if defined(SYCL_DEVICE_ONLY)
1407 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(floor, floor)
1408 #endif
1409
1410 #if defined(EIGEN_GPUCC)
1411 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1412 float floor(const float &x) { return ::floorf(x); }
1413
1414 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1415 double floor(const double &x) { return ::floor(x); }
1416 #endif
1417
1418 template<typename T>
1419 EIGEN_DEVICE_FUNC
1420 T (ceil)(const T& x)
1421 {
1422 EIGEN_USING_STD(ceil);
1423 return ceil(x);
1424 }
1425
1426 #if defined(SYCL_DEVICE_ONLY)
1427 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(ceil, ceil)
1428 #endif
1429
1430 #if defined(EIGEN_GPUCC)
1431 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1432 float ceil(const float &x) { return ::ceilf(x); }
1433
1434 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1435 double ceil(const double &x) { return ::ceil(x); }
1436 #endif
1437
1438
1439 /** Log base 2 for 32 bits positive integers.
1440 * Conveniently returns 0 for x==0. */
1441 inline int log2(int x)
1442 {
1443 eigen_assert(x>=0);
1444 unsigned int v(x);
1445 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 };
1446 v |= v >> 1;
1447 v |= v >> 2;
1448 v |= v >> 4;
1449 v |= v >> 8;
1450 v |= v >> 16;
1451 return table[(v * 0x07C4ACDDU) >> 27];
1452 }
1453
1454 /** \returns the square root of \a x.
1455 *
1456 * It is essentially equivalent to
1457 * \code using std::sqrt; return sqrt(x); \endcode
1458 * but slightly faster for float/double and some compilers (e.g., gcc), thanks to
1459 * specializations when SSE is enabled.
1460 *
1461 * It's usage is justified in performance critical functions, like norm/normalize.
1462 */
1463 template<typename Scalar>
1464 EIGEN_DEVICE_FUNC
1465 EIGEN_ALWAYS_INLINE EIGEN_MATHFUNC_RETVAL(sqrt, Scalar) sqrt(const Scalar& x)
1466 {
1467 return EIGEN_MATHFUNC_IMPL(sqrt, Scalar)::run(x);
1468 }
1469
1470 // Boolean specialization, avoids implicit float to bool conversion (-Wimplicit-conversion-floating-point-to-bool).
1471 template<>
1472 EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS EIGEN_DEVICE_FUNC
1473 bool sqrt<bool>(const bool &x) { return x; }
1474
1475 #if defined(SYCL_DEVICE_ONLY)
1476 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sqrt, sqrt)
1477 #endif
1478
1479 /** \returns the reciprocal square root of \a x. **/
1480 template<typename T>
1481 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1482 T rsqrt(const T& x)
1483 {
1484 return internal::rsqrt_impl<T>::run(x);
1485 }
1486
1487 template<typename T>
1488 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1489 T log(const T &x) {
1490 return internal::log_impl<T>::run(x);
1491 }
1492
1493 #if defined(SYCL_DEVICE_ONLY)
1494 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log, log)
1495 #endif
1496
1497
1498 #if defined(EIGEN_GPUCC)
1499 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1500 float log(const float &x) { return ::logf(x); }
1501
1502 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1503 double log(const double &x) { return ::log(x); }
1504 #endif
1505
1506 template<typename T>
1507 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1508 typename internal::enable_if<NumTraits<T>::IsSigned || NumTraits<T>::IsComplex,typename NumTraits<T>::Real>::type
1509 abs(const T &x) {
1510 EIGEN_USING_STD(abs);
1511 return abs(x);
1512 }
1513
1514 template<typename T>
1515 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1516 typename internal::enable_if<!(NumTraits<T>::IsSigned || NumTraits<T>::IsComplex),typename NumTraits<T>::Real>::type
1517 abs(const T &x) {
1518 return x;
1519 }
1520
1521 #if defined(SYCL_DEVICE_ONLY)
1522 SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(abs, abs)
1523 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(abs, fabs)
1524 #endif
1525
1526 #if defined(EIGEN_GPUCC)
1527 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1528 float abs(const float &x) { return ::fabsf(x); }
1529
1530 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1531 double abs(const double &x) { return ::fabs(x); }
1532
1533 template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1534 float abs(const std::complex<float>& x) {
1535 return ::hypotf(x.real(), x.imag());
1536 }
1537
1538 template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1539 double abs(const std::complex<double>& x) {
1540 return ::hypot(x.real(), x.imag());
1541 }
1542 #endif
1543
1544 template<typename T>
1545 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1546 T exp(const T &x) {
1547 EIGEN_USING_STD(exp);
1548 return exp(x);
1549 }
1550
1551 #if defined(SYCL_DEVICE_ONLY)
1552 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp, exp)
1553 #endif
1554
1555 #if defined(EIGEN_GPUCC)
1556 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1557 float exp(const float &x) { return ::expf(x); }
1558
1559 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1560 double exp(const double &x) { return ::exp(x); }
1561
1562 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1563 std::complex<float> exp(const std::complex<float>& x) {
1564 float com = ::expf(x.real());
1565 float res_real = com * ::cosf(x.imag());
1566 float res_imag = com * ::sinf(x.imag());
1567 return std::complex<float>(res_real, res_imag);
1568 }
1569
1570 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1571 std::complex<double> exp(const std::complex<double>& x) {
1572 double com = ::exp(x.real());
1573 double res_real = com * ::cos(x.imag());
1574 double res_imag = com * ::sin(x.imag());
1575 return std::complex<double>(res_real, res_imag);
1576 }
1577 #endif
1578
1579 template<typename Scalar>
1580 EIGEN_DEVICE_FUNC
1581 inline EIGEN_MATHFUNC_RETVAL(expm1, Scalar) expm1(const Scalar& x)
1582 {
1583 return EIGEN_MATHFUNC_IMPL(expm1, Scalar)::run(x);
1584 }
1585
1586 #if defined(SYCL_DEVICE_ONLY)
1587 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(expm1, expm1)
1588 #endif
1589
1590 #if defined(EIGEN_GPUCC)
1591 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1592 float expm1(const float &x) { return ::expm1f(x); }
1593
1594 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1595 double expm1(const double &x) { return ::expm1(x); }
1596 #endif
1597
1598 template<typename T>
1599 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1600 T cos(const T &x) {
1601 EIGEN_USING_STD(cos);
1602 return cos(x);
1603 }
1604
1605 #if defined(SYCL_DEVICE_ONLY)
1606 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cos,cos)
1607 #endif
1608
1609 #if defined(EIGEN_GPUCC)
1610 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1611 float cos(const float &x) { return ::cosf(x); }
1612
1613 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1614 double cos(const double &x) { return ::cos(x); }
1615 #endif
1616
1617 template<typename T>
1618 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1619 T sin(const T &x) {
1620 EIGEN_USING_STD(sin);
1621 return sin(x);
1622 }
1623
1624 #if defined(SYCL_DEVICE_ONLY)
1625 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sin, sin)
1626 #endif
1627
1628 #if defined(EIGEN_GPUCC)
1629 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1630 float sin(const float &x) { return ::sinf(x); }
1631
1632 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1633 double sin(const double &x) { return ::sin(x); }
1634 #endif
1635
1636 template<typename T>
1637 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1638 T tan(const T &x) {
1639 EIGEN_USING_STD(tan);
1640 return tan(x);
1641 }
1642
1643 #if defined(SYCL_DEVICE_ONLY)
1644 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tan, tan)
1645 #endif
1646
1647 #if defined(EIGEN_GPUCC)
1648 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1649 float tan(const float &x) { return ::tanf(x); }
1650
1651 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1652 double tan(const double &x) { return ::tan(x); }
1653 #endif
1654
1655 template<typename T>
1656 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1657 T acos(const T &x) {
1658 EIGEN_USING_STD(acos);
1659 return acos(x);
1660 }
1661
1662 #if EIGEN_HAS_CXX11_MATH
1663 template<typename T>
1664 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1665 T acosh(const T &x) {
1666 EIGEN_USING_STD(acosh);
1667 return static_cast<T>(acosh(x));
1668 }
1669 #endif
1670
1671 #if defined(SYCL_DEVICE_ONLY)
1672 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acos, acos)
1673 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acosh, acosh)
1674 #endif
1675
1676 #if defined(EIGEN_GPUCC)
1677 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1678 float acos(const float &x) { return ::acosf(x); }
1679
1680 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1681 double acos(const double &x) { return ::acos(x); }
1682 #endif
1683
1684 template<typename T>
1685 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1686 T asin(const T &x) {
1687 EIGEN_USING_STD(asin);
1688 return asin(x);
1689 }
1690
1691 #if EIGEN_HAS_CXX11_MATH
1692 template<typename T>
1693 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1694 T asinh(const T &x) {
1695 EIGEN_USING_STD(asinh);
1696 return static_cast<T>(asinh(x));
1697 }
1698 #endif
1699
1700 #if defined(SYCL_DEVICE_ONLY)
1701 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asin, asin)
1702 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asinh, asinh)
1703 #endif
1704
1705 #if defined(EIGEN_GPUCC)
1706 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1707 float asin(const float &x) { return ::asinf(x); }
1708
1709 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1710 double asin(const double &x) { return ::asin(x); }
1711 #endif
1712
1713 template<typename T>
1714 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1715 T atan(const T &x) {
1716 EIGEN_USING_STD(atan);
1717 return static_cast<T>(atan(x));
1718 }
1719
1720 #if EIGEN_HAS_CXX11_MATH
1721 template<typename T>
1722 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1723 T atanh(const T &x) {
1724 EIGEN_USING_STD(atanh);
1725 return static_cast<T>(atanh(x));
1726 }
1727 #endif
1728
1729 #if defined(SYCL_DEVICE_ONLY)
1730 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atan, atan)
1731 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atanh, atanh)
1732 #endif
1733
1734 #if defined(EIGEN_GPUCC)
1735 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1736 float atan(const float &x) { return ::atanf(x); }
1737
1738 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1739 double atan(const double &x) { return ::atan(x); }
1740 #endif
1741
1742
1743 template<typename T>
1744 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1745 T cosh(const T &x) {
1746 EIGEN_USING_STD(cosh);
1747 return static_cast<T>(cosh(x));
1748 }
1749
1750 #if defined(SYCL_DEVICE_ONLY)
1751 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cosh, cosh)
1752 #endif
1753
1754 #if defined(EIGEN_GPUCC)
1755 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1756 float cosh(const float &x) { return ::coshf(x); }
1757
1758 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1759 double cosh(const double &x) { return ::cosh(x); }
1760 #endif
1761
1762 template<typename T>
1763 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1764 T sinh(const T &x) {
1765 EIGEN_USING_STD(sinh);
1766 return static_cast<T>(sinh(x));
1767 }
1768
1769 #if defined(SYCL_DEVICE_ONLY)
1770 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sinh, sinh)
1771 #endif
1772
1773 #if defined(EIGEN_GPUCC)
1774 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1775 float sinh(const float &x) { return ::sinhf(x); }
1776
1777 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1778 double sinh(const double &x) { return ::sinh(x); }
1779 #endif
1780
1781 template<typename T>
1782 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1783 T tanh(const T &x) {
1784 EIGEN_USING_STD(tanh);
1785 return tanh(x);
1786 }
1787
1788 #if (!defined(EIGEN_GPUCC)) && EIGEN_FAST_MATH && !defined(SYCL_DEVICE_ONLY)
1789 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1790 float tanh(float x) { return internal::generic_fast_tanh_float(x); }
1791 #endif
1792
1793 #if defined(SYCL_DEVICE_ONLY)
1794 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tanh, tanh)
1795 #endif
1796
1797 #if defined(EIGEN_GPUCC)
1798 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1799 float tanh(const float &x) { return ::tanhf(x); }
1800
1801 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1802 double tanh(const double &x) { return ::tanh(x); }
1803 #endif
1804
1805 template <typename T>
1806 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1807 T fmod(const T& a, const T& b) {
1808 EIGEN_USING_STD(fmod);
1809 return fmod(a, b);
1810 }
1811
1812 #if defined(SYCL_DEVICE_ONLY)
1813 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(fmod, fmod)
1814 #endif
1815
1816 #if defined(EIGEN_GPUCC)
1817 template <>
1818 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1819 float fmod(const float& a, const float& b) {
1820 return ::fmodf(a, b);
1821 }
1822
1823 template <>
1824 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1825 double fmod(const double& a, const double& b) {
1826 return ::fmod(a, b);
1827 }
1828 #endif
1829
1830 #if defined(SYCL_DEVICE_ONLY)
1831 #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY
1832 #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY
1833 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY
1834 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1835 #undef SYCL_SPECIALIZE_INTEGER_TYPES_BINARY
1836 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1837 #undef SYCL_SPECIALIZE_FLOATING_TYPES_BINARY
1838 #undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY
1839 #undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE
1840 #undef SYCL_SPECIALIZE_GEN_UNARY_FUNC
1841 #undef SYCL_SPECIALIZE_UNARY_FUNC
1842 #undef SYCL_SPECIALIZE_GEN1_BINARY_FUNC
1843 #undef SYCL_SPECIALIZE_GEN2_BINARY_FUNC
1844 #undef SYCL_SPECIALIZE_BINARY_FUNC
1845 #endif
1846
1847 } // end namespace numext
1848
1849 namespace internal {
1850
1851 template<typename T>
1852 EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x)
1853 {
1854 return (numext::isfinite)(numext::real(x)) && (numext::isfinite)(numext::imag(x));
1855 }
1856
1857 template<typename T>
1858 EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x)
1859 {
1860 return (numext::isnan)(numext::real(x)) || (numext::isnan)(numext::imag(x));
1861 }
1862
1863 template<typename T>
1864 EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x)
1865 {
1866 return ((numext::isinf)(numext::real(x)) || (numext::isinf)(numext::imag(x))) && (!(numext::isnan)(x));
1867 }
1868
1869 /****************************************************************************
1870 * Implementation of fuzzy comparisons *
1871 ****************************************************************************/
1872
1873 template<typename Scalar,
1874 bool IsComplex,
1875 bool IsInteger>
1876 struct scalar_fuzzy_default_impl {};
1877
1878 template<typename Scalar>
1879 struct scalar_fuzzy_default_impl<Scalar, false, false>
1880 {
1881 typedef typename NumTraits<Scalar>::Real RealScalar;
1882 template<typename OtherScalar> EIGEN_DEVICE_FUNC
1883 static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y, const RealScalar& prec)
1884 {
1885 return numext::abs(x) <= numext::abs(y) * prec;
1886 }
1887 EIGEN_DEVICE_FUNC
1888 static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec)
1889 {
1890 return numext::abs(x - y) <= numext::mini(numext::abs(x), numext::abs(y)) * prec;
1891 }
1892 EIGEN_DEVICE_FUNC
1893 static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar& prec)
1894 {
1895 return x <= y || isApprox(x, y, prec);
1896 }
1897 };
1898
1899 template<typename Scalar>
1900 struct scalar_fuzzy_default_impl<Scalar, false, true>
1901 {
1902 typedef typename NumTraits<Scalar>::Real RealScalar;
1903 template<typename OtherScalar> EIGEN_DEVICE_FUNC
1904 static inline bool isMuchSmallerThan(const Scalar& x, const Scalar&, const RealScalar&)
1905 {
1906 return x == Scalar(0);
1907 }
1908 EIGEN_DEVICE_FUNC
1909 static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar&)
1910 {
1911 return x == y;
1912 }
1913 EIGEN_DEVICE_FUNC
1914 static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar&)
1915 {
1916 return x <= y;
1917 }
1918 };
1919
1920 template<typename Scalar>
1921 struct scalar_fuzzy_default_impl<Scalar, true, false>
1922 {
1923 typedef typename NumTraits<Scalar>::Real RealScalar;
1924 template<typename OtherScalar> EIGEN_DEVICE_FUNC
1925 static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y, const RealScalar& prec)
1926 {
1927 return numext::abs2(x) <= numext::abs2(y) * prec * prec;
1928 }
1929 EIGEN_DEVICE_FUNC
1930 static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec)
1931 {
1932 return numext::abs2(x - y) <= numext::mini(numext::abs2(x), numext::abs2(y)) * prec * prec;
1933 }
1934 };
1935
1936 template<typename Scalar>
1937 struct scalar_fuzzy_impl : scalar_fuzzy_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {};
1938
1939 template<typename Scalar, typename OtherScalar> EIGEN_DEVICE_FUNC
1940 inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
1941 const typename NumTraits<Scalar>::Real &precision = NumTraits<Scalar>::dummy_precision())
1942 {
1943 return scalar_fuzzy_impl<Scalar>::template isMuchSmallerThan<OtherScalar>(x, y, precision);
1944 }
1945
1946 template<typename Scalar> EIGEN_DEVICE_FUNC
1947 inline bool isApprox(const Scalar& x, const Scalar& y,
1948 const typename NumTraits<Scalar>::Real &precision = NumTraits<Scalar>::dummy_precision())
1949 {
1950 return scalar_fuzzy_impl<Scalar>::isApprox(x, y, precision);
1951 }
1952
1953 template<typename Scalar> EIGEN_DEVICE_FUNC
1954 inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y,
1955 const typename NumTraits<Scalar>::Real &precision = NumTraits<Scalar>::dummy_precision())
1956 {
1957 return scalar_fuzzy_impl<Scalar>::isApproxOrLessThan(x, y, precision);
1958 }
1959
1960 /******************************************
1961 *** The special case of the bool type ***
1962 ******************************************/
1963
1964 template<> struct random_impl<bool>
1965 {
1966 static inline bool run()
1967 {
1968 return random<int>(0,1)==0 ? false : true;
1969 }
1970
1971 static inline bool run(const bool& a, const bool& b)
1972 {
1973 return random<int>(a, b)==0 ? false : true;
1974 }
1975 };
1976
1977 template<> struct scalar_fuzzy_impl<bool>
1978 {
1979 typedef bool RealScalar;
1980
1981 template<typename OtherScalar> EIGEN_DEVICE_FUNC
1982 static inline bool isMuchSmallerThan(const bool& x, const bool&, const bool&)
1983 {
1984 return !x;
1985 }
1986
1987 EIGEN_DEVICE_FUNC
1988 static inline bool isApprox(bool x, bool y, bool)
1989 {
1990 return x == y;
1991 }
1992
1993 EIGEN_DEVICE_FUNC
1994 static inline bool isApproxOrLessThan(const bool& x, const bool& y, const bool&)
1995 {
1996 return (!x) || y;
1997 }
1998
1999 };
2000
2001 } // end namespace internal
2002
2003 // Default implementations that rely on other numext implementations
2004 namespace internal {
2005
2006 // Specialization for complex types that are not supported by std::expm1.
2007 template <typename RealScalar>
2008 struct expm1_impl<std::complex<RealScalar> > {
2009 EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(
2010 const std::complex<RealScalar>& x) {
2011 EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
2012 RealScalar xr = x.real();
2013 RealScalar xi = x.imag();
2014 // expm1(z) = exp(z) - 1
2015 // = exp(x + i * y) - 1
2016 // = exp(x) * (cos(y) + i * sin(y)) - 1
2017 // = exp(x) * cos(y) - 1 + i * exp(x) * sin(y)
2018 // Imag(expm1(z)) = exp(x) * sin(y)
2019 // Real(expm1(z)) = exp(x) * cos(y) - 1
2020 // = exp(x) * cos(y) - 1.
2021 // = expm1(x) + exp(x) * (cos(y) - 1)
2022 // = expm1(x) + exp(x) * (2 * sin(y / 2) ** 2)
2023 RealScalar erm1 = numext::expm1<RealScalar>(xr);
2024 RealScalar er = erm1 + RealScalar(1.);
2025 RealScalar sin2 = numext::sin(xi / RealScalar(2.));
2026 sin2 = sin2 * sin2;
2027 RealScalar s = numext::sin(xi);
2028 RealScalar real_part = erm1 - RealScalar(2.) * er * sin2;
2029 return std::complex<RealScalar>(real_part, er * s);
2030 }
2031 };
2032
2033 template<typename T>
2034 struct rsqrt_impl {
2035 EIGEN_DEVICE_FUNC
2036 static EIGEN_ALWAYS_INLINE T run(const T& x) {
2037 return T(1)/numext::sqrt(x);
2038 }
2039 };
2040
2041 #if defined(EIGEN_GPU_COMPILE_PHASE)
2042 template<typename T>
2043 struct conj_impl<std::complex<T>, true>
2044 {
2045 EIGEN_DEVICE_FUNC
2046 static inline std::complex<T> run(const std::complex<T>& x)
2047 {
2048 return std::complex<T>(numext::real(x), -numext::imag(x));
2049 }
2050 };
2051 #endif
2052
2053 } // end namespace internal
2054
2055 } // end namespace Eigen
2056
2057 #endif // EIGEN_MATHFUNCTIONS_H
2058