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