blob: f907d1e9bc4bf6819dd733d7abb1aea679bfd4bd [file] [log] [blame]
// This file is part of Eigen, a lightweight C++ template library
// for linear algebra.
//
// Copyright (C) 2006-2010 Benoit Jacob <jacob.benoit.1@gmail.com>
// 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<T>::type is a typedef for it.
* - otherwise, global_math_functions_filtering_base<T>::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<Derived>.
* So we must make sure to use sin_impl<ArrayBase<Derived> > and not sin_impl<Derived>, 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 <typename T, typename dummy = void>
struct global_math_functions_filtering_base {
typedef T type;
};
template <typename T>
struct always_void {
typedef void type;
};
template <typename T>
struct global_math_functions_filtering_base<
T, typename always_void<typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl>::type> {
typedef typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl type;
};
#define EIGEN_MATHFUNC_IMPL(func, scalar) \
Eigen::internal::func##_impl<typename Eigen::internal::global_math_functions_filtering_base<scalar>::type>
#define EIGEN_MATHFUNC_RETVAL(func, scalar) \
typename Eigen::internal::func##_retval< \
typename Eigen::internal::global_math_functions_filtering_base<scalar>::type>::type
/****************************************************************************
* Implementation of real *
****************************************************************************/
template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
struct real_default_impl {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x; }
};
template <typename Scalar>
struct real_default_impl<Scalar, true> {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
using std::real;
return real(x);
}
};
template <typename Scalar>
struct real_impl : real_default_impl<Scalar> {};
#if defined(EIGEN_GPU_COMPILE_PHASE)
template <typename T>
struct real_impl<std::complex<T>> {
typedef T RealScalar;
EIGEN_DEVICE_FUNC static inline T run(const std::complex<T>& x) { return x.real(); }
};
#endif
template <typename Scalar>
struct real_retval {
typedef typename NumTraits<Scalar>::Real type;
};
/****************************************************************************
* Implementation of imag *
****************************************************************************/
template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
struct imag_default_impl {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar&) { return RealScalar(0); }
};
template <typename Scalar>
struct imag_default_impl<Scalar, true> {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
using std::imag;
return imag(x);
}
};
template <typename Scalar>
struct imag_impl : imag_default_impl<Scalar> {};
#if defined(EIGEN_GPU_COMPILE_PHASE)
template <typename T>
struct imag_impl<std::complex<T>> {
typedef T RealScalar;
EIGEN_DEVICE_FUNC static inline T run(const std::complex<T>& x) { return x.imag(); }
};
#endif
template <typename Scalar>
struct imag_retval {
typedef typename NumTraits<Scalar>::Real type;
};
/****************************************************************************
* Implementation of real_ref *
****************************************************************************/
template <typename Scalar>
struct real_ref_impl {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast<RealScalar*>(&x)[0]; }
EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) {
return reinterpret_cast<const RealScalar*>(&x)[0];
}
};
template <typename Scalar>
struct real_ref_retval {
typedef typename NumTraits<Scalar>::Real& type;
};
/****************************************************************************
* Implementation of imag_ref *
****************************************************************************/
template <typename Scalar, bool IsComplex>
struct imag_ref_default_impl {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar& run(Scalar& x) { return reinterpret_cast<RealScalar*>(&x)[1]; }
EIGEN_DEVICE_FUNC static inline const RealScalar& run(const Scalar& x) {
return reinterpret_cast<const RealScalar*>(&x)[1];
}
};
template <typename Scalar>
struct imag_ref_default_impl<Scalar, false> {
EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static inline Scalar run(Scalar&) { return Scalar(0); }
EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR static inline const Scalar run(const Scalar&) { return Scalar(0); }
};
template <typename Scalar>
struct imag_ref_impl : imag_ref_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
template <typename Scalar>
struct imag_ref_retval {
typedef typename NumTraits<Scalar>::Real& type;
};
/****************************************************************************
* Implementation of conj *
****************************************************************************/
template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
struct conj_default_impl {
EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) { return x; }
};
template <typename Scalar>
struct conj_default_impl<Scalar, true> {
EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
using std::conj;
return conj(x);
}
};
template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
struct conj_impl : conj_default_impl<Scalar, IsComplex> {};
template <typename Scalar>
struct conj_retval {
typedef Scalar type;
};
/****************************************************************************
* Implementation of abs2 *
****************************************************************************/
template <typename Scalar, bool IsComplex>
struct abs2_impl_default {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x * x; }
};
template <typename Scalar>
struct abs2_impl_default<Scalar, true> // IsComplex
{
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) { return x.real() * x.real() + x.imag() * x.imag(); }
};
template <typename Scalar>
struct abs2_impl {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
return abs2_impl_default<Scalar, NumTraits<Scalar>::IsComplex>::run(x);
}
};
template <typename Scalar>
struct abs2_retval {
typedef typename NumTraits<Scalar>::Real type;
};
/****************************************************************************
* Implementation of sqrt/rsqrt *
****************************************************************************/
template <typename Scalar>
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 <typename T>
EIGEN_DEVICE_FUNC std::complex<T> complex_sqrt(const std::complex<T>& a_x);
// Custom implementation is faster than `std::sqrt`, works on
// GPU, and correctly handles special cases (unlike MSVC).
template <typename T>
struct sqrt_impl<std::complex<T>> {
EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x) {
return complex_sqrt<T>(x);
}
};
template <typename Scalar>
struct sqrt_retval {
typedef Scalar type;
};
// Default implementation relies on numext::sqrt, at bottom of file.
template <typename T>
struct rsqrt_impl;
// Complex rsqrt defined in MathFunctionsImpl.h.
template <typename T>
EIGEN_DEVICE_FUNC std::complex<T> complex_rsqrt(const std::complex<T>& a_x);
template <typename T>
struct rsqrt_impl<std::complex<T>> {
EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x) {
return complex_rsqrt<T>(x);
}
};
template <typename Scalar>
struct rsqrt_retval {
typedef Scalar type;
};
/****************************************************************************
* Implementation of norm1 *
****************************************************************************/
template <typename Scalar, bool IsComplex>
struct norm1_default_impl;
template <typename Scalar>
struct norm1_default_impl<Scalar, true> {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
EIGEN_USING_STD(abs);
return abs(x.real()) + abs(x.imag());
}
};
template <typename Scalar>
struct norm1_default_impl<Scalar, false> {
EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
EIGEN_USING_STD(abs);
return abs(x);
}
};
template <typename Scalar>
struct norm1_impl : norm1_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
template <typename Scalar>
struct norm1_retval {
typedef typename NumTraits<Scalar>::Real type;
};
/****************************************************************************
* Implementation of hypot *
****************************************************************************/
template <typename Scalar>
struct hypot_impl;
template <typename Scalar>
struct hypot_retval {
typedef typename NumTraits<Scalar>::Real type;
};
/****************************************************************************
* Implementation of cast *
****************************************************************************/
template <typename OldType, typename NewType, typename EnableIf = void>
struct cast_impl {
EIGEN_DEVICE_FUNC static inline NewType run(const OldType& x) { return static_cast<NewType>(x); }
};
template <typename OldType>
struct cast_impl<OldType, bool> {
EIGEN_DEVICE_FUNC static inline bool run(const OldType& x) { return x != OldType(0); }
};
// Casting from S -> Complex<T> leads to an implicit conversion from S to T,
// generating warnings on clang. Here we explicitly cast the real component.
template <typename OldType, typename NewType>
struct cast_impl<OldType, NewType,
typename std::enable_if_t<!NumTraits<OldType>::IsComplex && NumTraits<NewType>::IsComplex>> {
EIGEN_DEVICE_FUNC static inline NewType run(const OldType& x) {
typedef typename NumTraits<NewType>::Real NewReal;
return static_cast<NewType>(static_cast<NewReal>(x));
}
};
// here, for once, we're plainly returning NewType: we don't want cast to do weird things.
template <typename OldType, typename NewType>
EIGEN_DEVICE_FUNC inline NewType cast(const OldType& x) {
return cast_impl<OldType, NewType>::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 <typename Scalar, bool HasStdImpl = NumTraits<Scalar>::IsComplex || is_integral<Scalar>::value ||
is_same<Scalar, float>::value || is_same<Scalar, double>::value ||
is_same<Scalar, long double>::value>
struct arg_default_impl;
template <typename Scalar>
struct arg_default_impl<Scalar, true> {
typedef typename NumTraits<Scalar>::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<RealScalar>(arg(x));
}
};
// Must be non-complex floating-point type (e.g. half/bfloat16).
template <typename Scalar>
struct arg_default_impl<Scalar, false> {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
return (x < Scalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
}
};
#else
template <typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
struct arg_default_impl {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
return (x < RealScalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
}
};
template <typename Scalar>
struct arg_default_impl<Scalar, true> {
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_DEVICE_FUNC static inline RealScalar run(const Scalar& x) {
EIGEN_USING_STD(arg);
return arg(x);
}
};
#endif
template <typename Scalar>
struct arg_impl : arg_default_impl<Scalar> {};
template <typename Scalar>
struct arg_retval {
typedef typename NumTraits<Scalar>::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 <typename Scalar>
EIGEN_DEVICE_FUNC inline Scalar expm1(const Scalar& x) {
EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
typedef typename NumTraits<Scalar>::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 <typename Scalar>
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 <typename Scalar>
struct expm1_retval {
typedef Scalar type;
};
/****************************************************************************
* Implementation of log *
****************************************************************************/
// Complex log defined in MathFunctionsImpl.h.
template <typename T>
EIGEN_DEVICE_FUNC std::complex<T> complex_log(const std::complex<T>& z);
template <typename Scalar>
struct log_impl {
EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) {
EIGEN_USING_STD(log);
return static_cast<Scalar>(log(x));
}
};
template <typename Scalar>
struct log_impl<std::complex<Scalar>> {
EIGEN_DEVICE_FUNC static inline std::complex<Scalar> run(const std::complex<Scalar>& 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 <typename Scalar>
EIGEN_DEVICE_FUNC inline Scalar log1p(const Scalar& x) {
EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
typedef typename NumTraits<Scalar>::Real RealScalar;
EIGEN_USING_STD(log);
Scalar x1p = RealScalar(1) + x;
Scalar log_1p = log_impl<Scalar>::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 <typename Scalar>
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 <typename RealScalar>
struct log1p_impl<std::complex<RealScalar>> {
EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(const std::complex<RealScalar>& x) {
return std_fallback::log1p(x);
}
};
template <typename Scalar>
struct log1p_retval {
typedef Scalar type;
};
/****************************************************************************
* Implementation of pow *
****************************************************************************/
template <typename ScalarX, typename ScalarY,
bool IsInteger = NumTraits<ScalarX>::IsInteger && NumTraits<ScalarY>::IsInteger>
struct pow_impl {
// typedef Scalar retval;
typedef typename ScalarBinaryOpTraits<ScalarX, ScalarY, internal::scalar_pow_op<ScalarX, ScalarY>>::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 <typename ScalarX, typename ScalarY>
struct pow_impl<ScalarX, ScalarY, true> {
typedef ScalarX result_type;
static EIGEN_DEVICE_FUNC inline ScalarX run(ScalarX x, ScalarY y) {
ScalarX res(1);
eigen_assert(!NumTraits<ScalarY>::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 <unsigned int n, int lower, int upper>
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 <unsigned int n, int lower = 0, int upper = sizeof(unsigned int) * CHAR_BIT - 1,
int selector = meta_floor_log2_selector<n, lower, upper>::value>
struct meta_floor_log2 {};
template <unsigned int n, int lower, int upper>
struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_down> {
enum { value = meta_floor_log2<n, lower, meta_floor_log2_selector<n, lower, upper>::middle>::value };
};
template <unsigned int n, int lower, int upper>
struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_up> {
enum { value = meta_floor_log2<n, meta_floor_log2_selector<n, lower, upper>::middle, upper>::value };
};
template <unsigned int n, int lower, int upper>
struct meta_floor_log2<n, lower, upper, meta_floor_log2_terminate> {
enum { value = (n >= ((unsigned int)(1) << (lower + 1))) ? lower + 1 : lower };
};
template <unsigned int n, int lower, int upper>
struct meta_floor_log2<n, lower, upper, meta_floor_log2_bogus> {
// no value, error at compile time
};
template <typename BitsType, typename EnableIf = void>
struct count_bits_impl {
static_assert(std::is_integral<BitsType>::value && std::is_unsigned<BitsType>::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 <typename BitsType>
EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
return count_bits_impl<BitsType>::clz(bits);
}
// Count trailing zeros.
template <typename BitsType>
EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
return count_bits_impl<BitsType>::ctz(bits);
}
#if EIGEN_COMP_GNUC || EIGEN_COMP_CLANG
template <typename BitsType>
struct count_bits_impl<
BitsType, std::enable_if_t<std::is_integral<BitsType>::value && sizeof(BitsType) <= sizeof(unsigned int)>> {
static constexpr int kNumBits = static_cast<int>(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<unsigned int>(bits)) - kLeadingBitsOffset;
}
static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
return bits == 0 ? kNumBits : __builtin_ctz(static_cast<unsigned int>(bits));
}
};
template <typename BitsType>
struct count_bits_impl<BitsType,
std::enable_if_t<std::is_integral<BitsType>::value && sizeof(unsigned int) < sizeof(BitsType) &&
sizeof(BitsType) <= sizeof(unsigned long)>> {
static constexpr int kNumBits = static_cast<int>(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<unsigned long>(bits)) - kLeadingBitsOffset;
}
static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
return bits == 0 ? kNumBits : __builtin_ctzl(static_cast<unsigned long>(bits));
}
};
template <typename BitsType>
struct count_bits_impl<BitsType,
std::enable_if_t<std::is_integral<BitsType>::value && sizeof(unsigned long) < sizeof(BitsType) &&
sizeof(BitsType) <= sizeof(unsigned long long)>> {
static constexpr int kNumBits = static_cast<int>(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<unsigned long long>(bits)) - kLeadingBitsOffset;
}
static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
return bits == 0 ? kNumBits : __builtin_ctzll(static_cast<unsigned long long>(bits));
}
};
#elif EIGEN_COMP_MSVC
template <typename BitsType>
struct count_bits_impl<
BitsType, std::enable_if_t<std::is_integral<BitsType>::value && sizeof(BitsType) <= sizeof(unsigned long)>> {
static constexpr int kNumBits = static_cast<int>(sizeof(BitsType) * CHAR_BIT);
static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
unsigned long out;
_BitScanReverse(&out, static_cast<unsigned long>(bits));
return bits == 0 ? kNumBits : (kNumBits - 1) - static_cast<int>(out);
}
static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
unsigned long out;
_BitScanForward(&out, static_cast<unsigned long>(bits));
return bits == 0 ? kNumBits : static_cast<int>(out);
}
};
#ifdef _WIN64
template <typename BitsType>
struct count_bits_impl<BitsType,
std::enable_if_t<std::is_integral<BitsType>::value && sizeof(unsigned long) < sizeof(BitsType) &&
sizeof(BitsType) <= sizeof(__int64)>> {
static constexpr int kNumBits = static_cast<int>(sizeof(BitsType) * CHAR_BIT);
static EIGEN_DEVICE_FUNC inline int clz(BitsType bits) {
unsigned long out;
_BitScanReverse64(&out, static_cast<unsigned __int64>(bits));
return bits == 0 ? kNumBits : (kNumBits - 1) - static_cast<int>(out);
}
static EIGEN_DEVICE_FUNC inline int ctz(BitsType bits) {
unsigned long out;
_BitScanForward64(&out, static_cast<unsigned __int64>(bits));
return bits == 0 ? kNumBits : static_cast<int>(out);
}
};
#endif // _WIN64
#endif // EIGEN_COMP_GNUC || EIGEN_COMP_CLANG
template <typename BitsType>
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 <typename BitsType>
int log2_ceil(const BitsType& x) {
return log_2_impl<BitsType>::run_ceil(x);
}
template <typename BitsType>
int log2_floor(const BitsType& x) {
return log_2_impl<BitsType>::run_floor(x);
}
// Implementation of is* functions
template <typename T>
EIGEN_DEVICE_FUNC std::enable_if_t<!(std::numeric_limits<T>::has_infinity || std::numeric_limits<T>::has_quiet_NaN ||
std::numeric_limits<T>::has_signaling_NaN),
bool>
isfinite_impl(const T&) {
return true;
}
template <typename T>
EIGEN_DEVICE_FUNC std::enable_if_t<(std::numeric_limits<T>::has_infinity || std::numeric_limits<T>::has_quiet_NaN ||
std::numeric_limits<T>::has_signaling_NaN) &&
(!NumTraits<T>::IsComplex),
bool>
isfinite_impl(const T& x) {
EIGEN_USING_STD(isfinite);
return isfinite EIGEN_NOT_A_MACRO(x);
}
template <typename T>
EIGEN_DEVICE_FUNC std::enable_if_t<!std::numeric_limits<T>::has_infinity, bool> isinf_impl(const T&) {
return false;
}
template <typename T>
EIGEN_DEVICE_FUNC std::enable_if_t<(std::numeric_limits<T>::has_infinity && !NumTraits<T>::IsComplex), bool> isinf_impl(
const T& x) {
EIGEN_USING_STD(isinf);
return isinf EIGEN_NOT_A_MACRO(x);
}
template <typename T>
EIGEN_DEVICE_FUNC
std::enable_if_t<!(std::numeric_limits<T>::has_quiet_NaN || std::numeric_limits<T>::has_signaling_NaN), bool>
isnan_impl(const T&) {
return false;
}
template <typename T>
EIGEN_DEVICE_FUNC std::enable_if_t<
(std::numeric_limits<T>::has_quiet_NaN || std::numeric_limits<T>::has_signaling_NaN) && (!NumTraits<T>::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 <typename T>
EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x);
template <typename T>
EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x);
template <typename T>
EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x);
template <typename T>
EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS T ptanh_float(const T& a_x);
/****************************************************************************
* Implementation of sign *
****************************************************************************/
template <typename Scalar, bool IsComplex = (NumTraits<Scalar>::IsComplex != 0),
bool IsInteger = (NumTraits<Scalar>::IsInteger != 0)>
struct sign_impl {
EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) { return Scalar((a > Scalar(0)) - (a < Scalar(0))); }
};
template <typename Scalar>
struct sign_impl<Scalar, false, false> {
EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) {
return (isnan_impl<Scalar>)(a) ? a : Scalar((a > Scalar(0)) - (a < Scalar(0)));
}
};
template <typename Scalar, bool IsInteger>
struct sign_impl<Scalar, true, IsInteger> {
EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& a) {
using real_type = typename NumTraits<Scalar>::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(a.real() * aa, a.imag() * aa);
}
};
// The sign function for bool is the identity.
template <>
struct sign_impl<bool, false, true> {
EIGEN_DEVICE_FUNC static inline bool run(const bool& a) { return a; }
};
template <typename Scalar>
struct sign_retval {
typedef Scalar type;
};
template <typename Scalar, bool IsInteger = NumTraits<typename unpacket_traits<Scalar>::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);
}
};
template <typename Scalar>
struct nearest_integer_impl<Scalar, true> {
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; }
};
} // end namespace internal
/****************************************************************************
* Generic math functions *
****************************************************************************/
namespace numext {
#if (!defined(EIGEN_GPUCC) || defined(EIGEN_CONSTEXPR_ARE_DEVICE_FUNC))
template <typename T>
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 <typename T>
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 <typename T>
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 <typename T>
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 <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x) {
return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline internal::add_const_on_value_type_t<EIGEN_MATHFUNC_RETVAL(real_ref, Scalar)> real_ref(
const Scalar& x) {
return internal::real_ref_impl<Scalar>::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x) {
return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x) {
return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(arg, Scalar) arg(const Scalar& x) {
return EIGEN_MATHFUNC_IMPL(arg, Scalar)::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline internal::add_const_on_value_type_t<EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar)> imag_ref(
const Scalar& x) {
return internal::imag_ref_impl<Scalar>::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x) {
return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x) {
return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(sign, Scalar) sign(const Scalar& x) {
return EIGEN_MATHFUNC_IMPL(sign, Scalar)::run(x);
}
template <typename Scalar>
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 <typename T>
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 <typename Scalar>
EIGEN_DEVICE_FUNC inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x) {
return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x);
}
template <typename Scalar>
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 <typename Scalar>
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 <typename ScalarX, typename ScalarY>
EIGEN_DEVICE_FUNC inline typename internal::pow_impl<ScalarX, ScalarY>::result_type pow(const ScalarX& x,
const ScalarY& y) {
return internal::pow_impl<ScalarX, ScalarY>::run(x, y);
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(pow, pow)
#endif
template <typename T>
EIGEN_DEVICE_FUNC bool(isnan)(const T& x) {
return internal::isnan_impl(x);
}
template <typename T>
EIGEN_DEVICE_FUNC bool(isinf)(const T& x) {
return internal::isinf_impl(x);
}
template <typename T>
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 <typename Scalar>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar rint(const Scalar& x) {
return internal::nearest_integer_impl<Scalar>::run_rint(x);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar round(const Scalar& x) {
return internal::nearest_integer_impl<Scalar>::run_round(x);
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(round, round)
#endif
template <typename Scalar>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(floor)(const Scalar& x) {
return internal::nearest_integer_impl<Scalar>::run_floor(x);
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(floor, floor)
#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);
}
#endif
template <typename Scalar>
EIGEN_DEVICE_FUNC EIGEN_STRONG_INLINE Scalar(ceil)(const Scalar& x) {
return internal::nearest_integer_impl<Scalar>::run_ceil(x);
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(ceil, ceil)
#endif
#if defined(EIGEN_GPUCC)
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);
}
#endif
// Integer division with rounding up.
// T is assumed to be an integer type with a>=0, and b>0
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE EIGEN_CONSTEXPR T div_ceil(T a, T b) {
EIGEN_STATIC_ASSERT((NumTraits<T>::IsInteger), THIS FUNCTION IS FOR INTEGER TYPES)
eigen_assert(a >= 0);
eigen_assert(b > 0);
// Note: This form is used because it cannot overflow.
return a == 0 ? 0 : (a - 1) / b + 1;
}
/** Log base 2 for 32 bits positive integers.
* Conveniently returns 0 for x==0. */
inline int log2(int x) {
eigen_assert(x >= 0);
unsigned int v(x);
static const int table[32] = {0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30,
8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31};
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 <typename Scalar>
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<bool>(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 <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T cbrt(const T& x) {
EIGEN_USING_STD(cbrt);
return static_cast<T>(cbrt(x));
}
/** \returns the reciprocal square root of \a x. **/
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T rsqrt(const T& x) {
return internal::rsqrt_impl<T>::run(x);
}
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T log(const T& x) {
return internal::log_impl<T>::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 <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
std::enable_if_t<NumTraits<T>::IsSigned || NumTraits<T>::IsComplex, typename NumTraits<T>::Real>
abs(const T& x) {
EIGEN_USING_STD(abs);
return abs(x);
}
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
std::enable_if_t<!(NumTraits<T>::IsSigned || NumTraits<T>::IsComplex), typename NumTraits<T>::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<float>& x) {
return ::hypotf(x.real(), x.imag());
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double abs(const std::complex<double>& x) {
return ::hypot(x.real(), x.imag());
}
#endif
template <typename Scalar, bool IsInteger = NumTraits<Scalar>::IsInteger, bool IsSigned = NumTraits<Scalar>::IsSigned>
struct signbit_impl;
template <typename Scalar>
struct signbit_impl<Scalar, false, true> {
static constexpr size_t Size = sizeof(Scalar);
static constexpr size_t Shift = (CHAR_BIT * Size) - 1;
using intSize_t = typename get_integer_by_size<Size>::signed_type;
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static Scalar run(const Scalar& x) {
intSize_t a = bit_cast<intSize_t, Scalar>(x);
a = a >> Shift;
Scalar result = bit_cast<Scalar, intSize_t>(a);
return result;
}
};
template <typename Scalar>
struct signbit_impl<Scalar, true, true> {
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 <typename Scalar>
struct signbit_impl<Scalar, true, false> {
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar run(const Scalar&) { return Scalar(0); }
};
template <typename Scalar>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE static constexpr Scalar signbit(const Scalar& x) {
return signbit_impl<Scalar>::run(x);
}
template <typename T>
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 <typename RealScalar>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<RealScalar> exp(const std::complex<RealScalar>& 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<RealScalar>(NumTraits<RealScalar>::quiet_NaN(), NumTraits<RealScalar>::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<RealScalar>::infinity() && !(isfinite)(imag_ref(x)))) {
return std::complex<RealScalar>(NumTraits<RealScalar>::infinity(), NumTraits<RealScalar>::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<float> exp(const std::complex<float>& x) {
float com = ::expf(x.real());
float res_real = com * ::cosf(x.imag());
float res_imag = com * ::sinf(x.imag());
return std::complex<float>(res_real, res_imag);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE std::complex<double> exp(const std::complex<double>& x) {
double com = ::exp(x.real());
double res_real = com * ::cos(x.imag());
double res_imag = com * ::sin(x.imag());
return std::complex<double>(res_real, res_imag);
}
#endif
template <typename Scalar>
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 <typename T>
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 <typename T>
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 <typename T>
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 <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acos(const T& x) {
EIGEN_USING_STD(acos);
return acos(x);
}
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T acosh(const T& x) {
EIGEN_USING_STD(acosh);
return static_cast<T>(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 <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asin(const T& x) {
EIGEN_USING_STD(asin);
return asin(x);
}
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T asinh(const T& x) {
EIGEN_USING_STD(asinh);
return static_cast<T>(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 <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atan(const T& x) {
EIGEN_USING_STD(atan);
return static_cast<T>(atan(x));
}
template <typename T, std::enable_if_t<!NumTraits<T>::IsComplex, int> = 0>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atan2(const T& y, const T& x) {
EIGEN_USING_STD(atan2);
return static_cast<T>(atan2(y, x));
}
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T atanh(const T& x) {
EIGEN_USING_STD(atanh);
return static_cast<T>(atanh(x));
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atan, atan)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atanh, atanh)
#endif
#if defined(EIGEN_GPUCC)
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float atan(const float& x) {
return ::atanf(x);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double atan(const double& x) {
return ::atan(x);
}
#endif
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T cosh(const T& x) {
EIGEN_USING_STD(cosh);
return static_cast<T>(cosh(x));
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cosh, cosh)
#endif
#if defined(EIGEN_GPUCC)
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float cosh(const float& x) {
return ::coshf(x);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double cosh(const double& x) {
return ::cosh(x);
}
#endif
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T sinh(const T& x) {
EIGEN_USING_STD(sinh);
return static_cast<T>(sinh(x));
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sinh, sinh)
#endif
#if defined(EIGEN_GPUCC)
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float sinh(const float& x) {
return ::sinhf(x);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double sinh(const double& x) {
return ::sinh(x);
}
#endif
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T tanh(const T& x) {
EIGEN_USING_STD(tanh);
return tanh(x);
}
#if (!defined(EIGEN_GPUCC)) && EIGEN_FAST_MATH && !defined(SYCL_DEVICE_ONLY)
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tanh(float x) { return internal::ptanh_float(x); }
#endif
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tanh, tanh)
#endif
#if defined(EIGEN_GPUCC)
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float tanh(const float& x) {
return ::tanhf(x);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double tanh(const double& x) {
return ::tanh(x);
}
#endif
template <typename T>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE T fmod(const T& a, const T& b) {
EIGEN_USING_STD(fmod);
return fmod(a, b);
}
#if defined(SYCL_DEVICE_ONLY)
SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(fmod, fmod)
#endif
#if defined(EIGEN_GPUCC)
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE float fmod(const float& a, const float& b) {
return ::fmodf(a, b);
}
template <>
EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE double fmod(const double& a, const double& b) {
return ::fmod(a, b);
}
#endif
#if defined(SYCL_DEVICE_ONLY)
#undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY
#undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY
#undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY
#undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
#undef SYCL_SPECIALIZE_INTEGER_TYPES_BINARY
#undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
#undef SYCL_SPECIALIZE_FLOATING_TYPES_BINARY
#undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY
#undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE
#undef SYCL_SPECIALIZE_GEN_UNARY_FUNC
#undef SYCL_SPECIALIZE_UNARY_FUNC
#undef SYCL_SPECIALIZE_GEN1_BINARY_FUNC
#undef SYCL_SPECIALIZE_GEN2_BINARY_FUNC
#undef SYCL_SPECIALIZE_BINARY_FUNC
#endif
} // end namespace numext
namespace internal {
template <typename T>
EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x) {
return (numext::isfinite)(numext::real(x)) && (numext::isfinite)(numext::imag(x));
}
template <typename T>
EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x) {
return (numext::isnan)(numext::real(x)) || (numext::isnan)(numext::imag(x));
}
template <typename T>
EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x) {
return ((numext::isinf)(numext::real(x)) || (numext::isinf)(numext::imag(x))) && (!(numext::isnan)(x));
}
/****************************************************************************
* Implementation of fuzzy comparisons *
****************************************************************************/
template <typename Scalar, bool IsComplex, bool IsInteger>
struct scalar_fuzzy_default_impl {};
template <typename Scalar>
struct scalar_fuzzy_default_impl<Scalar, false, false> {
typedef typename NumTraits<Scalar>::Real RealScalar;
template <typename OtherScalar>
EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
const RealScalar& prec) {
return numext::abs(x) <= numext::abs(y) * prec;
}
EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) {
return numext::abs(x - y) <= numext::mini(numext::abs(x), numext::abs(y)) * prec;
}
EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar& prec) {
return x <= y || isApprox(x, y, prec);
}
};
template <typename Scalar>
struct scalar_fuzzy_default_impl<Scalar, false, true> {
typedef typename NumTraits<Scalar>::Real RealScalar;
template <typename OtherScalar>
EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const Scalar&, const RealScalar&) {
return x == Scalar(0);
}
EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar&) { return x == y; }
EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar&) {
return x <= y;
}
};
template <typename Scalar>
struct scalar_fuzzy_default_impl<Scalar, true, false> {
typedef typename NumTraits<Scalar>::Real RealScalar;
template <typename OtherScalar>
EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
const RealScalar& prec) {
return numext::abs2(x) <= numext::abs2(y) * prec * prec;
}
EIGEN_DEVICE_FUNC static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec) {
return numext::abs2(x - y) <= numext::mini(numext::abs2(x), numext::abs2(y)) * prec * prec;
}
};
template <typename Scalar>
struct scalar_fuzzy_impl
: scalar_fuzzy_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {};
template <typename Scalar, typename OtherScalar>
EIGEN_DEVICE_FUNC inline bool isMuchSmallerThan(
const Scalar& x, const OtherScalar& y,
const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
return scalar_fuzzy_impl<Scalar>::template isMuchSmallerThan<OtherScalar>(x, y, precision);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline bool isApprox(
const Scalar& x, const Scalar& y,
const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
return scalar_fuzzy_impl<Scalar>::isApprox(x, y, precision);
}
template <typename Scalar>
EIGEN_DEVICE_FUNC inline bool isApproxOrLessThan(
const Scalar& x, const Scalar& y,
const typename NumTraits<Scalar>::Real& precision = NumTraits<Scalar>::dummy_precision()) {
return scalar_fuzzy_impl<Scalar>::isApproxOrLessThan(x, y, precision);
}
/******************************************
*** The special case of the bool type ***
******************************************/
template <>
struct scalar_fuzzy_impl<bool> {
typedef bool RealScalar;
template <typename OtherScalar>
EIGEN_DEVICE_FUNC static inline bool isMuchSmallerThan(const bool& x, const bool&, const bool&) {
return !x;
}
EIGEN_DEVICE_FUNC static inline bool isApprox(bool x, bool y, bool) { return x == y; }
EIGEN_DEVICE_FUNC static inline bool isApproxOrLessThan(const bool& x, const bool& y, const bool&) {
return (!x) || y;
}
};
} // end namespace internal
// Default implementations that rely on other numext implementations
namespace internal {
// Specialization for complex types that are not supported by std::expm1.
template <typename RealScalar>
struct expm1_impl<std::complex<RealScalar>> {
EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(const std::complex<RealScalar>& x) {
RealScalar xr = x.real();
RealScalar xi = x.imag();
// expm1(z) = exp(z) - 1
// = exp(x + i * y) - 1
// = exp(x) * (cos(y) + i * sin(y)) - 1
// = exp(x) * cos(y) - 1 + i * exp(x) * sin(y)
// Imag(expm1(z)) = exp(x) * sin(y)
// Real(expm1(z)) = exp(x) * cos(y) - 1
// = exp(x) * cos(y) - 1.
// = expm1(x) + exp(x) * (cos(y) - 1)
// = expm1(x) + exp(x) * (2 * sin(y / 2) ** 2)
RealScalar erm1 = numext::expm1<RealScalar>(xr);
RealScalar er = erm1 + RealScalar(1.);
RealScalar sin2 = numext::sin(xi / RealScalar(2.));
sin2 = sin2 * sin2;
RealScalar s = numext::sin(xi);
RealScalar real_part = erm1 - RealScalar(2.) * er * sin2;
return std::complex<RealScalar>(real_part, er * s);
}
};
template <typename T>
struct rsqrt_impl {
EIGEN_DEVICE_FUNC static EIGEN_ALWAYS_INLINE T run(const T& x) { return T(1) / numext::sqrt(x); }
};
#if defined(EIGEN_GPU_COMPILE_PHASE)
template <typename T>
struct conj_impl<std::complex<T>, true> {
EIGEN_DEVICE_FUNC static inline std::complex<T> run(const std::complex<T>& x) {
return std::complex<T>(numext::real(x), -numext::imag(x));
}
};
#endif
} // end namespace internal
} // end namespace Eigen
#endif // EIGEN_MATHFUNCTIONS_H