// This file is part of Eigen, a lightweight C++ template library // for linear algebra. // // Copyright (C) 2006-2010 Benoit Jacob // Copyright (c) 2021, NVIDIA CORPORATION. All rights reserved. // // This Source Code Form is subject to the terms of the Mozilla // Public License v. 2.0. If a copy of the MPL was not distributed // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. #ifndef EIGEN_MATHFUNCTIONS_H #define EIGEN_MATHFUNCTIONS_H // TODO this should better be moved to NumTraits // Source: WolframAlpha #define EIGEN_PI 3.141592653589793238462643383279502884197169399375105820974944592307816406L #define EIGEN_LOG2E 1.442695040888963407359924681001892137426645954152985934135449406931109219L #define EIGEN_LN2 0.693147180559945309417232121458176568075500134360255254120680009493393621L // IWYU pragma: private #include "./InternalHeaderCheck.h" namespace Eigen { namespace internal { /** \internal \class global_math_functions_filtering_base * * What it does: * Defines a typedef 'type' as follows: * - if type T has a member typedef Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl, then * global_math_functions_filtering_base::type is a typedef for it. * - otherwise, global_math_functions_filtering_base::type is a typedef for T. * * How it's used: * To allow to defined the global math functions (like sin...) in certain cases, like the Array expressions. * When you do sin(array1+array2), the object array1+array2 has a complicated expression type, all what you want to know * is that it inherits ArrayBase. So we implement a partial specialization of sin_impl for ArrayBase. * So we must make sure to use sin_impl > and not sin_impl, otherwise our partial * specialization won't be used. How does sin know that? That's exactly what global_math_functions_filtering_base tells * it. * * How it's implemented: * SFINAE in the style of enable_if. Highly susceptible of breaking compilers. With GCC, it sure does work, but if you * replace the typename dummy by an integer template parameter, it doesn't work anymore! */ template struct global_math_functions_filtering_base { typedef T type; }; template struct always_void { typedef void type; }; template struct global_math_functions_filtering_base< T, typename always_void::type> { typedef typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl type; }; #define EIGEN_MATHFUNC_IMPL(func, scalar) \ Eigen::internal::func##_impl::type> #define EIGEN_MATHFUNC_RETVAL(func, scalar) \ typename Eigen::internal::func##_retval< \ typename Eigen::internal::global_math_functions_filtering_base::type>::type /**************************************************************************** * Implementation of real * ****************************************************************************/ template ::IsComplex> struct real_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x; } }; template struct real_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { using std::real; return real(x); } }; template struct real_impl : real_default_impl {}; #if defined(EIGEN_GPU_COMPILE_PHASE) template struct real_impl> { typedef T RealScalar; EIGEN_DEVICE_FUNC static inline T run(const std::complex& x) { return x.real(); } }; #endif template struct real_retval { typedef typename NumTraits::Real type; }; /**************************************************************************** * Implementation of imag * ****************************************************************************/ template ::IsComplex> struct imag_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar&) { return RealScalar(0); } }; template struct imag_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { using std::imag; return imag(x); } }; template struct imag_impl : imag_default_impl {}; #if defined(EIGEN_GPU_COMPILE_PHASE) template struct imag_impl> { typedef T RealScalar; EIGEN_DEVICE_FUNC static inline T run(const std::complex& x) { return x.imag(); } }; #endif template struct imag_retval { typedef typename NumTraits::Real type; }; /**************************************************************************** * Implementation of real_ref * ****************************************************************************/ template struct real_ref_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast(&x)[0]; } EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) { return reinterpret_cast(&x)[0]; } }; template struct real_ref_retval { typedef typename NumTraits::Real& type; }; /**************************************************************************** * Implementation of imag_ref * ****************************************************************************/ template struct imag_ref_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast(&x)[1]; } EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) { return reinterpret_cast(&x)[1]; } }; template struct imag_ref_default_impl { EIGEN_DEVICE_FUNC constexpr static Scalar run(Scalar&) { return Scalar(0); } EIGEN_DEVICE_FUNC constexpr static const Scalar run(const Scalar&) { return Scalar(0); } }; template struct imag_ref_impl : imag_ref_default_impl::IsComplex> {}; template struct imag_ref_retval { typedef typename NumTraits::Real& type; }; } // namespace internal namespace numext { template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline internal::add_const_on_value_type_t real_ref( const Scalar& x) { return internal::real_ref_impl::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x) { return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar select(const Scalar& mask, const Scalar& a, const Scalar& b) { return numext::is_exactly_zero(mask) ? b : a; } } // namespace numext namespace internal { /**************************************************************************** * Implementation of conj * ****************************************************************************/ template ::IsComplex> struct conj_default_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { return x; } }; template struct conj_default_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { using std::conj; return conj(x); } }; template ::IsComplex> struct conj_impl : conj_default_impl {}; template struct conj_retval { typedef Scalar type; }; /**************************************************************************** * Implementation of abs2 * ****************************************************************************/ template struct abs2_impl_default { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x * x; } }; template struct abs2_impl_default // IsComplex { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return numext::real(x) * numext::real(x) + numext::imag(x) * numext::imag(x); } }; template struct abs2_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return abs2_impl_default::IsComplex>::run(x); } }; template struct abs2_retval { typedef typename NumTraits::Real type; }; /**************************************************************************** * Implementation of sqrt/rsqrt * ****************************************************************************/ template struct sqrt_impl { EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE Scalar run(const Scalar& x) { EIGEN_USING_STD(sqrt); return sqrt(x); } }; // Complex sqrt defined in MathFunctionsImpl.h. template EIGEN_DEVICE_FUNC ComplexT complex_sqrt(const ComplexT& a_x); // Custom implementation is faster than `std::sqrt`, works on // GPU, and correctly handles special cases (unlike MSVC). template struct sqrt_impl> { EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE std::complex run(const std::complex& x) { return complex_sqrt(x); } }; template struct sqrt_retval { typedef Scalar type; }; // Default implementation relies on numext::sqrt, at bottom of file. template struct rsqrt_impl; // Complex rsqrt defined in MathFunctionsImpl.h. template EIGEN_DEVICE_FUNC ComplexT complex_rsqrt(const ComplexT& a_x); template struct rsqrt_impl> { EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE std::complex run(const std::complex& x) { return complex_rsqrt(x); } }; template struct rsqrt_retval { typedef Scalar type; }; /**************************************************************************** * Implementation of norm1 * ****************************************************************************/ template struct norm1_default_impl; template struct norm1_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { EIGEN_USING_STD(abs); return abs(numext::real(x)) + abs(numext::imag(x)); } }; template struct norm1_default_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { EIGEN_USING_STD(abs); return abs(x); } }; template struct norm1_impl : norm1_default_impl::IsComplex> {}; template struct norm1_retval { typedef typename NumTraits::Real type; }; /**************************************************************************** * Implementation of hypot * ****************************************************************************/ template struct hypot_impl; template struct hypot_retval { typedef typename NumTraits::Real type; }; /**************************************************************************** * Implementation of cast * ****************************************************************************/ template struct cast_impl { EIGEN_DEVICE_FUNC static inline NewType run(const OldType& x) { return static_cast(x); } }; template struct cast_impl { EIGEN_DEVICE_FUNC static inline bool run(const OldType& x) { return x != OldType(0); } }; // Casting from S -> Complex leads to an implicit conversion from S to T, // generating warnings on clang. Here we explicitly cast the real component. template struct cast_impl::IsComplex && NumTraits::IsComplex>> { EIGEN_DEVICE_FUNC static inline NewType run(const OldType& x) { typedef typename NumTraits::Real NewReal; return static_cast(static_cast(x)); } }; // here, for once, we're plainly returning NewType: we don't want cast to do weird things. template EIGEN_DEVICE_FUNC inline NewType cast(const OldType& x) { return cast_impl::run(x); } /**************************************************************************** * Implementation of arg * ****************************************************************************/ // Visual Studio 2017 has a bug where arg(float) returns 0 for negative inputs. // This seems to be fixed in VS 2019. #if (!EIGEN_COMP_MSVC || EIGEN_COMP_MSVC >= 1920) // std::arg is only defined for types of std::complex, or integer types or float/double/long double template ::IsComplex || is_integral::value || is_same::value || is_same::value || is_same::value> struct arg_default_impl; template struct arg_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { // There is no official ::arg on device in CUDA/HIP, so we always need to use std::arg. using std::arg; return static_cast(arg(x)); } }; // Must be non-complex floating-point type (e.g. half/bfloat16). template struct arg_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return (x < Scalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0); } }; #else template ::IsComplex> struct arg_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return (x < RealScalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0); } }; template struct arg_default_impl { typedef typename NumTraits::Real RealScalar; EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { EIGEN_USING_STD(arg); return arg(x); } }; #endif template struct arg_impl : arg_default_impl {}; template struct arg_retval { typedef typename NumTraits::Real type; }; /**************************************************************************** * Implementation of expm1 * ****************************************************************************/ // This implementation is based on GSL Math's expm1. namespace std_fallback { // fallback expm1 implementation in case there is no expm1(Scalar) function in namespace of Scalar, // or that there is no suitable std::expm1 function available. Implementation // attributed to Kahan. See: http://www.plunk.org/~hatch/rightway.php. template EIGEN_DEVICE_FUNC inline Scalar expm1(const Scalar& x) { EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) typedef typename NumTraits::Real RealScalar; EIGEN_USING_STD(exp); Scalar u = exp(x); if (numext::equal_strict(u, Scalar(1))) { return x; } Scalar um1 = u - RealScalar(1); if (numext::equal_strict(um1, Scalar(-1))) { return RealScalar(-1); } EIGEN_USING_STD(log); Scalar logu = log(u); return numext::equal_strict(u, logu) ? u : (u - RealScalar(1)) * x / logu; } } // namespace std_fallback template struct expm1_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) EIGEN_USING_STD(expm1); return expm1(x); } }; template struct expm1_retval { typedef Scalar type; }; /**************************************************************************** * Implementation of log * ****************************************************************************/ // Complex log defined in MathFunctionsImpl.h. template EIGEN_DEVICE_FUNC ComplexT complex_log(const ComplexT& z); template struct log_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { EIGEN_USING_STD(log); return static_cast(log(x)); } }; template struct log_impl> { EIGEN_DEVICE_FUNC static inline std::complex run(const std::complex& z) { return complex_log(z); } }; /**************************************************************************** * Implementation of log1p * ****************************************************************************/ namespace std_fallback { // fallback log1p implementation in case there is no log1p(Scalar) function in namespace of Scalar, // or that there is no suitable std::log1p function available template EIGEN_DEVICE_FUNC inline Scalar log1p(const Scalar& x) { EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) typedef typename NumTraits::Real RealScalar; EIGEN_USING_STD(log); Scalar x1p = RealScalar(1) + x; Scalar log_1p = log_impl::run(x1p); const bool is_small = numext::equal_strict(x1p, Scalar(1)); const bool is_inf = numext::equal_strict(x1p, log_1p); return (is_small || is_inf) ? x : x * (log_1p / (x1p - RealScalar(1))); } } // namespace std_fallback template struct log1p_impl { EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { EIGEN_USING_STD(log1p); return log1p(x); } }; // Specialization for complex types that are not supported by std::log1p. template struct log1p_impl> { EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar) EIGEN_DEVICE_FUNC static inline std::complex run(const std::complex& x) { return std_fallback::log1p(x); } }; template struct log1p_retval { typedef Scalar type; }; /**************************************************************************** * Implementation of pow * ****************************************************************************/ template ::IsInteger && NumTraits::IsInteger> struct pow_impl { // typedef Scalar retval; typedef typename ScalarBinaryOpTraits>::ReturnType result_type; static EIGEN_DEVICE_FUNC inline result_type run(const ScalarX& x, const ScalarY& y) { EIGEN_USING_STD(pow); return pow(x, y); } }; template struct pow_impl { typedef ScalarX result_type; static EIGEN_DEVICE_FUNC inline ScalarX run(ScalarX x, ScalarY y) { ScalarX res(1); eigen_assert(!NumTraits::IsSigned || y >= 0); if (y & 1) res *= x; y >>= 1; while (y) { x *= x; if (y & 1) res *= x; y >>= 1; } return res; } }; enum { meta_floor_log2_terminate, meta_floor_log2_move_up, meta_floor_log2_move_down, meta_floor_log2_bogus }; template struct meta_floor_log2_selector { enum { middle = (lower + upper) / 2, value = (upper <= lower + 1) ? int(meta_floor_log2_terminate) : (n < (1 << middle)) ? int(meta_floor_log2_move_down) : (n == 0) ? int(meta_floor_log2_bogus) : int(meta_floor_log2_move_up) }; }; template ::value> struct meta_floor_log2 {}; template struct meta_floor_log2 { enum { value = meta_floor_log2::middle>::value }; }; template struct meta_floor_log2 { enum { value = meta_floor_log2::middle, upper>::value }; }; template struct meta_floor_log2 { enum { value = (n >= ((unsigned int)(1) << (lower + 1))) ? lower + 1 : lower }; }; template struct meta_floor_log2 { // no value, error at compile time }; template struct count_bits_impl { static_assert(std::is_integral::value && std::is_unsigned::value, "BitsType must be an unsigned integer"); static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) { int n = CHAR_BIT * sizeof(BitsType); int shift = n / 2; while (bits > 0 && shift > 0) { BitsType y = bits >> shift; if (y > 0) { n -= shift; bits = y; } shift /= 2; } if (shift == 0) { --n; } return n; } static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) { int n = CHAR_BIT * sizeof(BitsType); int shift = n / 2; while (bits > 0 && shift > 0) { BitsType y = bits << shift; if (y > 0) { n -= shift; bits = y; } shift /= 2; } if (shift == 0) { --n; } return n; } }; // Count leading zeros. template EIGEN_DEVICE_FUNC inline int clz(BitsType bits) { return count_bits_impl::clz(bits); } // Count trailing zeros. template EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) { return count_bits_impl::ctz(bits); } #if EIGEN_COMP_GNUC || EIGEN_COMP_CLANG template struct count_bits_impl< BitsType, std::enable_if_t::value && sizeof(BitsType) <= sizeof(unsigned int)>> { static constexpr int kNumBits = static_cast(sizeof(BitsType) * CHAR_BIT); static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) { static constexpr int kLeadingBitsOffset = (sizeof(unsigned int) - sizeof(BitsType)) * CHAR_BIT; return bits == 0 ? kNumBits : __builtin_clz(static_cast(bits)) - kLeadingBitsOffset; } static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) { return bits == 0 ? kNumBits : __builtin_ctz(static_cast(bits)); } }; template struct count_bits_impl::value && sizeof(unsigned int) < sizeof(BitsType) && sizeof(BitsType) <= sizeof(unsigned long)>> { static constexpr int kNumBits = static_cast(sizeof(BitsType) * CHAR_BIT); static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) { static constexpr int kLeadingBitsOffset = (sizeof(unsigned long) - sizeof(BitsType)) * CHAR_BIT; return bits == 0 ? kNumBits : __builtin_clzl(static_cast(bits)) - kLeadingBitsOffset; } static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) { return bits == 0 ? kNumBits : __builtin_ctzl(static_cast(bits)); } }; template struct count_bits_impl::value && sizeof(unsigned long) < sizeof(BitsType) && sizeof(BitsType) <= sizeof(unsigned long long)>> { static constexpr int kNumBits = static_cast(sizeof(BitsType) * CHAR_BIT); static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) { static constexpr int kLeadingBitsOffset = (sizeof(unsigned long long) - sizeof(BitsType)) * CHAR_BIT; return bits == 0 ? kNumBits : __builtin_clzll(static_cast(bits)) - kLeadingBitsOffset; } static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) { return bits == 0 ? kNumBits : __builtin_ctzll(static_cast(bits)); } }; #elif EIGEN_COMP_MSVC template struct count_bits_impl< BitsType, std::enable_if_t::value && sizeof(BitsType) <= sizeof(unsigned long)>> { static constexpr int kNumBits = static_cast(sizeof(BitsType) * CHAR_BIT); static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) { unsigned long out; _BitScanReverse(&out, static_cast(bits)); return bits == 0 ? kNumBits : (kNumBits - 1) - static_cast(out); } static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) { unsigned long out; _BitScanForward(&out, static_cast(bits)); return bits == 0 ? kNumBits : static_cast(out); } }; #ifdef _WIN64 template struct count_bits_impl::value && sizeof(unsigned long) < sizeof(BitsType) && sizeof(BitsType) <= sizeof(__int64)>> { static constexpr int kNumBits = static_cast(sizeof(BitsType) * CHAR_BIT); static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) { unsigned long out; _BitScanReverse64(&out, static_cast(bits)); return bits == 0 ? kNumBits : (kNumBits - 1) - static_cast(out); } static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) { unsigned long out; _BitScanForward64(&out, static_cast(bits)); return bits == 0 ? kNumBits : static_cast(out); } }; #endif // _WIN64 #endif // EIGEN_COMP_GNUC || EIGEN_COMP_CLANG template struct log_2_impl { static constexpr int kTotalBits = sizeof(BitsType) * CHAR_BIT; static EIGEN_DEVICE_FUNC inline int run_ceil(const BitsType& x) { const int n = kTotalBits - clz(x); bool power_of_two = (x & (x - 1)) == 0; return x == 0 ? 0 : power_of_two ? (n - 1) : n; } static EIGEN_DEVICE_FUNC inline int run_floor(const BitsType& x) { const int n = kTotalBits - clz(x); return x == 0 ? 0 : n - 1; } }; template int log2_ceil(const BitsType& x) { return log_2_impl::run_ceil(x); } template int log2_floor(const BitsType& x) { return log_2_impl::run_floor(x); } // Implementation of is* functions template EIGEN_DEVICE_FUNC std::enable_if_t::has_infinity || std::numeric_limits::has_quiet_NaN || std::numeric_limits::has_signaling_NaN), bool> isfinite_impl(const T&) { return true; } template EIGEN_DEVICE_FUNC std::enable_if_t<(std::numeric_limits::has_infinity || std::numeric_limits::has_quiet_NaN || std::numeric_limits::has_signaling_NaN) && (!NumTraits::IsComplex), bool> isfinite_impl(const T& x) { EIGEN_USING_STD(isfinite); return isfinite EIGEN_NOT_A_MACRO(x); } template EIGEN_DEVICE_FUNC std::enable_if_t::has_infinity, bool> isinf_impl(const T&) { return false; } template EIGEN_DEVICE_FUNC std::enable_if_t<(std::numeric_limits::has_infinity && !NumTraits::IsComplex), bool> isinf_impl( const T& x) { EIGEN_USING_STD(isinf); return isinf EIGEN_NOT_A_MACRO(x); } template EIGEN_DEVICE_FUNC std::enable_if_t::has_quiet_NaN || std::numeric_limits::has_signaling_NaN), bool> isnan_impl(const T&) { return false; } template EIGEN_DEVICE_FUNC std::enable_if_t< (std::numeric_limits::has_quiet_NaN || std::numeric_limits::has_signaling_NaN) && (!NumTraits::IsComplex), bool> isnan_impl(const T& x) { EIGEN_USING_STD(isnan); return isnan EIGEN_NOT_A_MACRO(x); } // The following overload are defined at the end of this file template EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex& x); template EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex& x); template EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex& x); template EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS T ptanh_float(const T& a_x); /**************************************************************************** * Implementation of sign * ****************************************************************************/ template ::IsComplex != 0), bool IsInteger = (NumTraits::IsInteger != 0)> struct sign_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) { return Scalar((a > Scalar(0)) - (a < Scalar(0))); } }; template struct sign_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) { return (isnan_impl)(a) ? a : Scalar((a > Scalar(0)) - (a < Scalar(0))); } }; template struct sign_impl { EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) { using real_type = typename NumTraits::Real; EIGEN_USING_STD(abs); real_type aa = abs(a); if (aa == real_type(0)) return Scalar(0); aa = real_type(1) / aa; return Scalar(numext::real(a) * aa, numext::imag(a) * aa); } }; // The sign function for bool is the identity. template <> struct sign_impl { EIGEN_DEVICE_FUNC static inline bool run(const bool& a) { return a; } }; template struct sign_retval { typedef Scalar type; }; // suppress "unary minus operator applied to unsigned type, result still unsigned" warnings on MSVC // note: `0 - a` is distinct from `-a` when Scalar is a floating point type and `a` is zero template ::IsInteger> struct negate_impl { static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar run(const Scalar& a) { return -a; } }; template struct negate_impl { EIGEN_STATIC_ASSERT((!is_same::value), NEGATE IS NOT DEFINED FOR BOOLEAN TYPES) static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE Scalar run(const Scalar& a) { return Scalar(0) - a; } }; template struct negate_retval { typedef Scalar type; }; template ::type>::IsInteger> struct nearest_integer_impl { static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_floor(const Scalar& x) { EIGEN_USING_STD(floor) return floor(x); } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_ceil(const Scalar& x) { EIGEN_USING_STD(ceil) return ceil(x); } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_rint(const Scalar& x) { EIGEN_USING_STD(rint) return rint(x); } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_round(const Scalar& x) { EIGEN_USING_STD(round) return round(x); } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_trunc(const Scalar& x) { EIGEN_USING_STD(trunc) return trunc(x); } }; template struct nearest_integer_impl { static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_floor(const Scalar& x) { return x; } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_ceil(const Scalar& x) { return x; } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_rint(const Scalar& x) { return x; } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_round(const Scalar& x) { return x; } static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run_trunc(const Scalar& x) { return x; } }; // Extra namespace to prevent leaking std::fma into Eigen::internal. namespace has_fma_detail { template struct has_fma_impl : public std::false_type {}; using std::fma; template struct has_fma_impl< T, std::enable_if_t(), std::declval(), std::declval()))>::value>> : public std::true_type {}; } // namespace has_fma_detail template struct has_fma : public has_fma_detail::has_fma_impl {}; // Default implementation. template struct fma_impl { static_assert(has_fma::value, "No function fma(...) for type. Please provide an implementation."); }; // STD or ADL version if it exists. template struct fma_impl::value>> { static EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T run(const T& a, const T& b, const T& c) { using std::fma; return fma(a, b, c); } }; #if defined(EIGEN_GPUCC) template <> struct has_fma : public true_type {}; template <> struct fma_impl { static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE float run(const float& a, const float& b, const float& c) { return ::fmaf(a, b, c); } }; template <> struct has_fma : public true_type {}; template <> struct fma_impl { static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE double run(const double& a, const double& b, const double& c) { return ::fma(a, b, c); } }; #endif // Basic multiply-add. template struct madd_impl { static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run(const Scalar& x, const Scalar& y, const Scalar& z) { return x * y + z; } }; #if EIGEN_SCALAR_MADD_USE_FMA template struct madd_impl::value>> { static EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar run(const Scalar& x, const Scalar& y, const Scalar& z) { return fma_impl::run(x, y, z); } }; #endif } // end namespace internal /**************************************************************************** * Generic math functions * ****************************************************************************/ namespace numext { #if (!defined(EIGEN_GPUCC) || defined(EIGEN_CONSTEXPR_ARE_DEVICE_FUNC)) template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y) { EIGEN_USING_STD(min) return min EIGEN_NOT_A_MACRO(x, y); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y) { EIGEN_USING_STD(max) return max EIGEN_NOT_A_MACRO(x, y); } #else template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y) { return y < x ? y : x; } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float mini(const float& x, const float& y) { return fminf(x, y); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double mini(const double& x, const double& y) { return fmin(x, y); } #ifndef EIGEN_GPU_COMPILE_PHASE template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double mini(const long double& x, const long double& y) { #if defined(EIGEN_HIPCC) // no "fminl" on HIP yet return (x < y) ? x : y; #else return fminl(x, y); #endif } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y) { return x < y ? y : x; } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float maxi(const float& x, const float& y) { return fmaxf(x, y); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double maxi(const double& x, const double& y) { return fmax(x, y); } #ifndef EIGEN_GPU_COMPILE_PHASE template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double maxi(const long double& x, const long double& y) { #if defined(EIGEN_HIPCC) // no "fmaxl" on HIP yet return (x > y) ? x : y; #else return fmaxl(x, y); #endif } #endif #endif #if defined(SYCL_DEVICE_ONLY) #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_long) #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_char) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_short) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_int) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_long) #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong) #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uint) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong) #define SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(NAME, FUNC) \ SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \ SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) #define SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(NAME, FUNC) \ SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \ SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) #define SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(NAME, FUNC) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \ SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_double) #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(NAME, FUNC) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \ SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_double) #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(NAME, FUNC, RET_TYPE) \ SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_float) \ SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_double) #define SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE& x) { \ return cl::sycl::FUNC(x); \ } #define SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, TYPE) SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, TYPE, TYPE) #define SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE1, ARG_TYPE2) \ template <> \ EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE1& x, const ARG_TYPE2& y) { \ return cl::sycl::FUNC(x, y); \ } #define SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \ SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE, ARG_TYPE) #define SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, TYPE) SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, TYPE, TYPE) SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(mini, min) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(mini, fmin) SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(maxi, max) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(maxi, fmax) #endif template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(arg, Scalar) arg(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(arg, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline internal::add_const_on_value_type_t imag_ref( const Scalar& x) { return internal::imag_ref_impl::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x) { return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(sign, Scalar) sign(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(sign, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(negate, Scalar) negate(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(negate, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x); } EIGEN_DEVICE_FUNC inline bool abs2(bool x) { return x; } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T absdiff(const T& x, const T& y) { return x > y ? x - y : y - x; } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float absdiff(const float& x, const float& y) { return fabsf(x - y); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double absdiff(const double& x, const double& y) { return fabs(x - y); } // HIP and CUDA do not support long double. #ifndef EIGEN_GPU_COMPILE_PHASE template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE long double absdiff(const long double& x, const long double& y) { return fabsl(x - y); } #endif template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x); } template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(hypot, Scalar) hypot(const Scalar& x, const Scalar& y) { return EIGEN_MATHFUNC_IMPL(hypot, Scalar)::run(x, y); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(hypot, hypot) #endif template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(log1p, Scalar) log1p(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(log1p, Scalar)::run(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log1p, log1p) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float log1p(const float& x) { return ::log1pf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double log1p(const double& x) { return ::log1p(x); } #endif template EIGEN_DEVICE_FUNC inline typename internal::pow_impl::result_type pow(const ScalarX& x, const ScalarY& y) { return internal::pow_impl::run(x, y); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(pow, pow) #endif template EIGEN_DEVICE_FUNC bool(isnan)(const T& x) { return internal::isnan_impl(x); } template EIGEN_DEVICE_FUNC bool(isinf)(const T& x) { return internal::isinf_impl(x); } template EIGEN_DEVICE_FUNC bool(isfinite)(const T& x) { return internal::isfinite_impl(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isnan, isnan, bool) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isinf, isinf, bool) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isfinite, isfinite, bool) #endif template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar rint(const Scalar& x) { return internal::nearest_integer_impl::run_rint(x); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar round(const Scalar& x) { return internal::nearest_integer_impl::run_round(x); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(floor)(const Scalar& x) { return internal::nearest_integer_impl::run_floor(x); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(ceil)(const Scalar& x) { return internal::nearest_integer_impl::run_ceil(x); } template EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(trunc)(const Scalar& x) { return internal::nearest_integer_impl::run_trunc(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(round, round) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(floor, floor) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(ceil, ceil) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(trunc, trunc) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float floor(const float& x) { return ::floorf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double floor(const double& x) { return ::floor(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float ceil(const float& x) { return ::ceilf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double ceil(const double& x) { return ::ceil(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float trunc(const float& x) { return ::truncf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double trunc(const double& x) { return ::trunc(x); } #endif // Integer division with rounding up. // T is assumed to be an integer type with a>=0, and b>0 template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE constexpr T div_ceil(T a, T b) { using UnsignedT = typename internal::make_unsigned::type; EIGEN_STATIC_ASSERT((NumTraits::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES) // Note: explicitly declaring a and b as non-negative values allows the compiler to use better optimizations const UnsignedT ua = UnsignedT(a); const UnsignedT ub = UnsignedT(b); // Note: This form is used because it cannot overflow. return ua == 0 ? 0 : (ua - 1) / ub + 1; } // Integer round down to nearest power of b // T is assumed to be an integer type with a>=0, and b>0 template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE constexpr T round_down(T a, U b) { using UnsignedT = typename internal::make_unsigned::type; using UnsignedU = typename internal::make_unsigned::type; EIGEN_STATIC_ASSERT((NumTraits::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES) EIGEN_STATIC_ASSERT((NumTraits::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES) // Note: explicitly declaring a and b as non-negative values allows the compiler to use better optimizations const UnsignedT ua = UnsignedT(a); const UnsignedU ub = UnsignedU(b); return ub * (ua / ub); } /** Log base 2 for 32 bits positive integers. * Conveniently returns 0 for x==0. */ constexpr int log2(int x) { unsigned int v(x); constexpr int table[32] = {0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30, 8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31}; v |= v >> 1; v |= v >> 2; v |= v >> 4; v |= v >> 8; v |= v >> 16; return table[(v * 0x07C4ACDDU) >> 27]; } /** \returns the square root of \a x. * * It is essentially equivalent to * \code using std::sqrt; return sqrt(x); \endcode * but slightly faster for float/double and some compilers (e.g., gcc), thanks to * specializations when SSE is enabled. * * It's usage is justified in performance critical functions, like norm/normalize. */ template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE EIGEN_MATHFUNC_RETVAL(sqrt, Scalar) sqrt(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(sqrt, Scalar)::run(x); } // Boolean specialization, avoids implicit float to bool conversion (-Wimplicit-conversion-floating-point-to-bool). template <> EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS EIGEN_DEVICE_FUNC bool sqrt(const bool& x) { return x; } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sqrt, sqrt) #endif /** \returns the cube root of \a x. **/ template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t::IsComplex, T> cbrt(const T& x) { EIGEN_USING_STD(cbrt); return static_cast(cbrt(x)); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t::IsComplex, T> cbrt(const T& x) { EIGEN_USING_STD(pow); return pow(x, typename NumTraits::Real(1.0 / 3.0)); } /** \returns the reciprocal square root of \a x. **/ template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T rsqrt(const T& x) { return internal::rsqrt_impl::run(x); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T log(const T& x) { return internal::log_impl::run(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log, log) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float log(const float& x) { return ::logf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double log(const double& x) { return ::log(x); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t::IsSigned || NumTraits::IsComplex, typename NumTraits::Real> abs(const T& x) { EIGEN_USING_STD(abs); return abs(x); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::enable_if_t::IsSigned || NumTraits::IsComplex), typename NumTraits::Real> abs(const T& x) { return x; } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(abs, abs) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(abs, fabs) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float abs(const float& x) { return ::fabsf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double abs(const double& x) { return ::fabs(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float abs(const std::complex& x) { return ::hypotf(x.real(), x.imag()); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double abs(const std::complex& x) { return ::hypot(x.real(), x.imag()); } #endif template ::IsInteger, bool IsSigned = NumTraits::IsSigned> struct signbit_impl; template struct signbit_impl { static constexpr size_t Size = sizeof(Scalar); static constexpr size_t Shift = (CHAR_BIT * Size) - 1; using intSize_t = typename get_integer_by_size::signed_type; EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static Scalar run(const Scalar& x) { intSize_t a = bit_cast(x); a = a >> Shift; Scalar result = bit_cast(a); return result; } }; template struct signbit_impl { static constexpr size_t Size = sizeof(Scalar); static constexpr size_t Shift = (CHAR_BIT * Size) - 1; EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar run(const Scalar& x) { return x >> Shift; } }; template struct signbit_impl { EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar run(const Scalar&) { return Scalar(0); } }; template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar signbit(const Scalar& x) { return signbit_impl::run(x); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T exp(const T& x) { EIGEN_USING_STD(exp); return exp(x); } // MSVC screws up some edge-cases for std::exp(complex). #ifdef EIGEN_COMP_MSVC template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex exp(const std::complex& x) { EIGEN_USING_STD(exp); // If z is (x,±∞) (for any finite x), the result is (NaN,NaN) and FE_INVALID is raised. // If z is (x,NaN) (for any finite x), the result is (NaN,NaN) and FE_INVALID may be raised. if ((isfinite)(real_ref(x)) && !(isfinite)(imag_ref(x))) { return std::complex(NumTraits::quiet_NaN(), NumTraits::quiet_NaN()); } // If z is (+∞,±∞), the result is (±∞,NaN) and FE_INVALID is raised (the sign of the real part is unspecified) // If z is (+∞,NaN), the result is (±∞,NaN) (the sign of the real part is unspecified) if ((real_ref(x) == NumTraits::infinity() && !(isfinite)(imag_ref(x)))) { return std::complex(NumTraits::infinity(), NumTraits::quiet_NaN()); } return exp(x); } #endif #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp, exp) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float exp(const float& x) { return ::expf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double exp(const double& x) { return ::exp(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex exp(const std::complex& x) { float com = ::expf(x.real()); float res_real = com * ::cosf(x.imag()); float res_imag = com * ::sinf(x.imag()); return std::complex(res_real, res_imag); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex exp(const std::complex& x) { double com = ::exp(x.real()); double res_real = com * ::cos(x.imag()); double res_imag = com * ::sin(x.imag()); return std::complex(res_real, res_imag); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T exp2(const T& x) { EIGEN_USING_STD(exp2); return exp2(x); } // MSVC screws up some edge-cases for std::exp2(complex). #ifdef EIGEN_COMP_MSVC template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex exp2(const std::complex& x) { EIGEN_USING_STD(exp); // If z is (x,±∞) (for any finite x), the result is (NaN,NaN) and FE_INVALID is raised. // If z is (x,NaN) (for any finite x), the result is (NaN,NaN) and FE_INVALID may be raised. if ((isfinite)(real_ref(x)) && !(isfinite)(imag_ref(x))) { return std::complex(NumTraits::quiet_NaN(), NumTraits::quiet_NaN()); } // If z is (+∞,±∞), the result is (±∞,NaN) and FE_INVALID is raised (the sign of the real part is unspecified) // If z is (+∞,NaN), the result is (±∞,NaN) (the sign of the real part is unspecified) if ((real_ref(x) == NumTraits::infinity() && !(isfinite)(imag_ref(x)))) { return std::complex(NumTraits::infinity(), NumTraits::quiet_NaN()); } return exp2(x); } #endif #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp2, exp2) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float exp2(const float& x) { return ::exp2f(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double exp2(const double& x) { return ::exp2(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex exp2(const std::complex& x) { float com = ::exp2f(x.real()); float res_real = com * ::cosf(static_cast(EIGEN_LN2) * x.imag()); float res_imag = com * ::sinf(static_cast(EIGEN_LN2) * x.imag()); return std::complex(res_real, res_imag); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex exp2(const std::complex& x) { double com = ::exp2(x.real()); double res_real = com * ::cos(static_cast(EIGEN_LN2) * x.imag()); double res_imag = com * ::sin(static_cast(EIGEN_LN2) * x.imag()); return std::complex(res_real, res_imag); } #endif template EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(expm1, Scalar) expm1(const Scalar& x) { return EIGEN_MATHFUNC_IMPL(expm1, Scalar)::run(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(expm1, expm1) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float expm1(const float& x) { return ::expm1f(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double expm1(const double& x) { return ::expm1(x); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T cos(const T& x) { EIGEN_USING_STD(cos); return cos(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cos, cos) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float cos(const float& x) { return ::cosf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double cos(const double& x) { return ::cos(x); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T sin(const T& x) { EIGEN_USING_STD(sin); return sin(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sin, sin) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float sin(const float& x) { return ::sinf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double sin(const double& x) { return ::sin(x); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T tan(const T& x) { EIGEN_USING_STD(tan); return tan(x); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tan, tan) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tan(const float& x) { return ::tanf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double tan(const double& x) { return ::tan(x); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acos(const T& x) { EIGEN_USING_STD(acos); return acos(x); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acosh(const T& x) { EIGEN_USING_STD(acosh); return static_cast(acosh(x)); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acos, acos) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acosh, acosh) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float acos(const float& x) { return ::acosf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double acos(const double& x) { return ::acos(x); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asin(const T& x) { EIGEN_USING_STD(asin); return asin(x); } template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asinh(const T& x) { EIGEN_USING_STD(asinh); return static_cast(asinh(x)); } #if defined(SYCL_DEVICE_ONLY) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asin, asin) SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asinh, asinh) #endif #if defined(EIGEN_GPUCC) template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float asin(const float& x) { return ::asinf(x); } template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double asin(const double& x) { return ::asin(x); } #endif template EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atan(const T& x) { EIGEN_USING_STD(atan); return static_cast