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