Back to home page

EIC code displayed by LXR
 
 

    

File indexing completed on 2025-01-18 09:56:15 

0001 // This file is part of Eigen, a lightweight C++ template library
0002 // for linear algebra.
0003 //
0004 // Copyright (C) 2006-2010 Benoit Jacob <jacob.benoit.1@gmail.com>
0005 // Copyright (c) 2021, NVIDIA CORPORATION. All rights reserved.
0006 //
0007 // This Source Code Form is subject to the terms of the Mozilla
0008 // Public License v. 2.0. If a copy of the MPL was not distributed
0009 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/.
0010 
0011 #ifndef EIGEN_MATHFUNCTIONS_H
0012 #define EIGEN_MATHFUNCTIONS_H
0013 
0014 // TODO this should better be moved to NumTraits
0015 // Source: WolframAlpha
0016 #define EIGEN_PI    3.141592653589793238462643383279502884197169399375105820974944592307816406L
0017 #define EIGEN_LOG2E 1.442695040888963407359924681001892137426645954152985934135449406931109219L
0018 #define EIGEN_LN2   0.693147180559945309417232121458176568075500134360255254120680009493393621L
0019 
0020 namespace Eigen {
0021 
0022 // On WINCE, std::abs is defined for int only, so let's defined our own overloads:
0023 // This issue has been confirmed with MSVC 2008 only, but the issue might exist for more recent versions too.
0024 #if EIGEN_OS_WINCE && EIGEN_COMP_MSVC && EIGEN_COMP_MSVC<=1500
0025 long        abs(long        x) { return (labs(x));  }
0026 double      abs(double      x) { return (fabs(x));  }
0027 float       abs(float       x) { return (fabsf(x)); }
0028 long double abs(long double x) { return (fabsl(x)); }
0029 #endif
0030 
0031 namespace internal {
0032 
0033 /** \internal \class global_math_functions_filtering_base
0034   *
0035   * What it does:
0036   * Defines a typedef 'type' as follows:
0037   * - if type T has a member typedef Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl, then
0038   *   global_math_functions_filtering_base<T>::type is a typedef for it.
0039   * - otherwise, global_math_functions_filtering_base<T>::type is a typedef for T.
0040   *
0041   * How it's used:
0042   * To allow to defined the global math functions (like sin...) in certain cases, like the Array expressions.
0043   * When you do sin(array1+array2), the object array1+array2 has a complicated expression type, all what you want to know
0044   * is that it inherits ArrayBase. So we implement a partial specialization of sin_impl for ArrayBase<Derived>.
0045   * So we must make sure to use sin_impl<ArrayBase<Derived> > and not sin_impl<Derived>, otherwise our partial specialization
0046   * won't be used. How does sin know that? That's exactly what global_math_functions_filtering_base tells it.
0047   *
0048   * How it's implemented:
0049   * SFINAE in the style of enable_if. Highly susceptible of breaking compilers. With GCC, it sure does work, but if you replace
0050   * the typename dummy by an integer template parameter, it doesn't work anymore!
0051   */
0052 
0053 template<typename T, typename dummy = void>
0054 struct global_math_functions_filtering_base
0055 {
0056   typedef T type;
0057 };
0058 
0059 template<typename T> struct always_void { typedef void type; };
0060 
0061 template<typename T>
0062 struct global_math_functions_filtering_base
0063   <T,
0064    typename always_void<typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl>::type
0065   >
0066 {
0067   typedef typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl type;
0068 };
0069 
0070 #define EIGEN_MATHFUNC_IMPL(func, scalar) Eigen::internal::func##_impl<typename Eigen::internal::global_math_functions_filtering_base<scalar>::type>
0071 #define EIGEN_MATHFUNC_RETVAL(func, scalar) typename Eigen::internal::func##_retval<typename Eigen::internal::global_math_functions_filtering_base<scalar>::type>::type
0072 
0073 /****************************************************************************
0074 * Implementation of real                                                 *
0075 ****************************************************************************/
0076 
0077 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
0078 struct real_default_impl
0079 {
0080   typedef typename NumTraits<Scalar>::Real RealScalar;
0081   EIGEN_DEVICE_FUNC
0082   static inline RealScalar run(const Scalar& x)
0083   {
0084     return x;
0085   }
0086 };
0087 
0088 template<typename Scalar>
0089 struct real_default_impl<Scalar,true>
0090 {
0091   typedef typename NumTraits<Scalar>::Real RealScalar;
0092   EIGEN_DEVICE_FUNC
0093   static inline RealScalar run(const Scalar& x)
0094   {
0095     using std::real;
0096     return real(x);
0097   }
0098 };
0099 
0100 template<typename Scalar> struct real_impl : real_default_impl<Scalar> {};
0101 
0102 #if defined(EIGEN_GPU_COMPILE_PHASE)
0103 template<typename T>
0104 struct real_impl<std::complex<T> >
0105 {
0106   typedef T RealScalar;
0107   EIGEN_DEVICE_FUNC
0108   static inline T run(const std::complex<T>& x)
0109   {
0110     return x.real();
0111   }
0112 };
0113 #endif
0114 
0115 template<typename Scalar>
0116 struct real_retval
0117 {
0118   typedef typename NumTraits<Scalar>::Real type;
0119 };
0120 
0121 /****************************************************************************
0122 * Implementation of imag                                                 *
0123 ****************************************************************************/
0124 
0125 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
0126 struct imag_default_impl
0127 {
0128   typedef typename NumTraits<Scalar>::Real RealScalar;
0129   EIGEN_DEVICE_FUNC
0130   static inline RealScalar run(const Scalar&)
0131   {
0132     return RealScalar(0);
0133   }
0134 };
0135 
0136 template<typename Scalar>
0137 struct imag_default_impl<Scalar,true>
0138 {
0139   typedef typename NumTraits<Scalar>::Real RealScalar;
0140   EIGEN_DEVICE_FUNC
0141   static inline RealScalar run(const Scalar& x)
0142   {
0143     using std::imag;
0144     return imag(x);
0145   }
0146 };
0147 
0148 template<typename Scalar> struct imag_impl : imag_default_impl<Scalar> {};
0149 
0150 #if defined(EIGEN_GPU_COMPILE_PHASE)
0151 template<typename T>
0152 struct imag_impl<std::complex<T> >
0153 {
0154   typedef T RealScalar;
0155   EIGEN_DEVICE_FUNC
0156   static inline T run(const std::complex<T>& x)
0157   {
0158     return x.imag();
0159   }
0160 };
0161 #endif
0162 
0163 template<typename Scalar>
0164 struct imag_retval
0165 {
0166   typedef typename NumTraits<Scalar>::Real type;
0167 };
0168 
0169 /****************************************************************************
0170 * Implementation of real_ref                                             *
0171 ****************************************************************************/
0172 
0173 template<typename Scalar>
0174 struct real_ref_impl
0175 {
0176   typedef typename NumTraits<Scalar>::Real RealScalar;
0177   EIGEN_DEVICE_FUNC
0178   static inline RealScalar& run(Scalar& x)
0179   {
0180     return reinterpret_cast<RealScalar*>(&x)[0];
0181   }
0182   EIGEN_DEVICE_FUNC
0183   static inline const RealScalar& run(const Scalar& x)
0184   {
0185     return reinterpret_cast<const RealScalar*>(&x)[0];
0186   }
0187 };
0188 
0189 template<typename Scalar>
0190 struct real_ref_retval
0191 {
0192   typedef typename NumTraits<Scalar>::Real & type;
0193 };
0194 
0195 /****************************************************************************
0196 * Implementation of imag_ref                                             *
0197 ****************************************************************************/
0198 
0199 template<typename Scalar, bool IsComplex>
0200 struct imag_ref_default_impl
0201 {
0202   typedef typename NumTraits<Scalar>::Real RealScalar;
0203   EIGEN_DEVICE_FUNC
0204   static inline RealScalar& run(Scalar& x)
0205   {
0206     return reinterpret_cast<RealScalar*>(&x)[1];
0207   }
0208   EIGEN_DEVICE_FUNC
0209   static inline const RealScalar& run(const Scalar& x)
0210   {
0211     return reinterpret_cast<RealScalar*>(&x)[1];
0212   }
0213 };
0214 
0215 template<typename Scalar>
0216 struct imag_ref_default_impl<Scalar, false>
0217 {
0218   EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
0219   static inline Scalar run(Scalar&)
0220   {
0221     return Scalar(0);
0222   }
0223   EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR
0224   static inline const Scalar run(const Scalar&)
0225   {
0226     return Scalar(0);
0227   }
0228 };
0229 
0230 template<typename Scalar>
0231 struct imag_ref_impl : imag_ref_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
0232 
0233 template<typename Scalar>
0234 struct imag_ref_retval
0235 {
0236   typedef typename NumTraits<Scalar>::Real & type;
0237 };
0238 
0239 /****************************************************************************
0240 * Implementation of conj                                                 *
0241 ****************************************************************************/
0242 
0243 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
0244 struct conj_default_impl
0245 {
0246   EIGEN_DEVICE_FUNC
0247   static inline Scalar run(const Scalar& x)
0248   {
0249     return x;
0250   }
0251 };
0252 
0253 template<typename Scalar>
0254 struct conj_default_impl<Scalar,true>
0255 {
0256   EIGEN_DEVICE_FUNC
0257   static inline Scalar run(const Scalar& x)
0258   {
0259     using std::conj;
0260     return conj(x);
0261   }
0262 };
0263 
0264 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
0265 struct conj_impl : conj_default_impl<Scalar, IsComplex> {};
0266 
0267 template<typename Scalar>
0268 struct conj_retval
0269 {
0270   typedef Scalar type;
0271 };
0272 
0273 /****************************************************************************
0274 * Implementation of abs2                                                 *
0275 ****************************************************************************/
0276 
0277 template<typename Scalar,bool IsComplex>
0278 struct abs2_impl_default
0279 {
0280   typedef typename NumTraits<Scalar>::Real RealScalar;
0281   EIGEN_DEVICE_FUNC
0282   static inline RealScalar run(const Scalar& x)
0283   {
0284     return x*x;
0285   }
0286 };
0287 
0288 template<typename Scalar>
0289 struct abs2_impl_default<Scalar, true> // IsComplex
0290 {
0291   typedef typename NumTraits<Scalar>::Real RealScalar;
0292   EIGEN_DEVICE_FUNC
0293   static inline RealScalar run(const Scalar& x)
0294   {
0295     return x.real()*x.real() + x.imag()*x.imag();
0296   }
0297 };
0298 
0299 template<typename Scalar>
0300 struct abs2_impl
0301 {
0302   typedef typename NumTraits<Scalar>::Real RealScalar;
0303   EIGEN_DEVICE_FUNC
0304   static inline RealScalar run(const Scalar& x)
0305   {
0306     return abs2_impl_default<Scalar,NumTraits<Scalar>::IsComplex>::run(x);
0307   }
0308 };
0309 
0310 template<typename Scalar>
0311 struct abs2_retval
0312 {
0313   typedef typename NumTraits<Scalar>::Real type;
0314 };
0315 
0316 /****************************************************************************
0317 * Implementation of sqrt/rsqrt                                             *
0318 ****************************************************************************/
0319 
0320 template<typename Scalar>
0321 struct sqrt_impl
0322 {
0323   EIGEN_DEVICE_FUNC
0324   static EIGEN_ALWAYS_INLINE Scalar run(const Scalar& x)
0325   {
0326     EIGEN_USING_STD(sqrt);
0327     return sqrt(x);
0328   }
0329 };
0330 
0331 // Complex sqrt defined in MathFunctionsImpl.h.
0332 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_sqrt(const std::complex<T>& a_x);
0333 
0334 // Custom implementation is faster than `std::sqrt`, works on
0335 // GPU, and correctly handles special cases (unlike MSVC).
0336 template<typename T>
0337 struct sqrt_impl<std::complex<T> >
0338 {
0339   EIGEN_DEVICE_FUNC
0340   static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x)
0341   {
0342     return complex_sqrt<T>(x);
0343   }
0344 };
0345 
0346 template<typename Scalar>
0347 struct sqrt_retval
0348 {
0349   typedef Scalar type;
0350 };
0351 
0352 // Default implementation relies on numext::sqrt, at bottom of file.
0353 template<typename T>
0354 struct rsqrt_impl;
0355 
0356 // Complex rsqrt defined in MathFunctionsImpl.h.
0357 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_rsqrt(const std::complex<T>& a_x);
0358 
0359 template<typename T>
0360 struct rsqrt_impl<std::complex<T> >
0361 {
0362   EIGEN_DEVICE_FUNC
0363   static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x)
0364   {
0365     return complex_rsqrt<T>(x);
0366   }
0367 };
0368 
0369 template<typename Scalar>
0370 struct rsqrt_retval
0371 {
0372   typedef Scalar type;
0373 };
0374 
0375 /****************************************************************************
0376 * Implementation of norm1                                                *
0377 ****************************************************************************/
0378 
0379 template<typename Scalar, bool IsComplex>
0380 struct norm1_default_impl;
0381 
0382 template<typename Scalar>
0383 struct norm1_default_impl<Scalar,true>
0384 {
0385   typedef typename NumTraits<Scalar>::Real RealScalar;
0386   EIGEN_DEVICE_FUNC
0387   static inline RealScalar run(const Scalar& x)
0388   {
0389     EIGEN_USING_STD(abs);
0390     return abs(x.real()) + abs(x.imag());
0391   }
0392 };
0393 
0394 template<typename Scalar>
0395 struct norm1_default_impl<Scalar, false>
0396 {
0397   EIGEN_DEVICE_FUNC
0398   static inline Scalar run(const Scalar& x)
0399   {
0400     EIGEN_USING_STD(abs);
0401     return abs(x);
0402   }
0403 };
0404 
0405 template<typename Scalar>
0406 struct norm1_impl : norm1_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {};
0407 
0408 template<typename Scalar>
0409 struct norm1_retval
0410 {
0411   typedef typename NumTraits<Scalar>::Real type;
0412 };
0413 
0414 /****************************************************************************
0415 * Implementation of hypot                                                *
0416 ****************************************************************************/
0417 
0418 template<typename Scalar> struct hypot_impl;
0419 
0420 template<typename Scalar>
0421 struct hypot_retval
0422 {
0423   typedef typename NumTraits<Scalar>::Real type;
0424 };
0425 
0426 /****************************************************************************
0427 * Implementation of cast                                                 *
0428 ****************************************************************************/
0429 
0430 template<typename OldType, typename NewType, typename EnableIf = void>
0431 struct cast_impl
0432 {
0433   EIGEN_DEVICE_FUNC
0434   static inline NewType run(const OldType& x)
0435   {
0436     return static_cast<NewType>(x);
0437   }
0438 };
0439 
0440 // Casting from S -> Complex<T> leads to an implicit conversion from S to T,
0441 // generating warnings on clang.  Here we explicitly cast the real component.
0442 template<typename OldType, typename NewType>
0443 struct cast_impl<OldType, NewType,
0444   typename internal::enable_if<
0445     !NumTraits<OldType>::IsComplex && NumTraits<NewType>::IsComplex
0446   >::type>
0447 {
0448   EIGEN_DEVICE_FUNC
0449   static inline NewType run(const OldType& x)
0450   {
0451     typedef typename NumTraits<NewType>::Real NewReal;
0452     return static_cast<NewType>(static_cast<NewReal>(x));
0453   }
0454 };
0455 
0456 // here, for once, we're plainly returning NewType: we don't want cast to do weird things.
0457 
0458 template<typename OldType, typename NewType>
0459 EIGEN_DEVICE_FUNC
0460 inline NewType cast(const OldType& x)
0461 {
0462   return cast_impl<OldType, NewType>::run(x);
0463 }
0464 
0465 /****************************************************************************
0466 * Implementation of round                                                   *
0467 ****************************************************************************/
0468 
0469 template<typename Scalar>
0470 struct round_impl
0471 {
0472   EIGEN_DEVICE_FUNC
0473   static inline Scalar run(const Scalar& x)
0474   {
0475     EIGEN_STATIC_ASSERT((!NumTraits<Scalar>::IsComplex), NUMERIC_TYPE_MUST_BE_REAL)
0476 #if EIGEN_HAS_CXX11_MATH
0477     EIGEN_USING_STD(round);
0478 #endif
0479     return Scalar(round(x));
0480   }
0481 };
0482 
0483 #if !EIGEN_HAS_CXX11_MATH
0484 #if EIGEN_HAS_C99_MATH
0485 // Use ::roundf for float.
0486 template<>
0487 struct round_impl<float> {
0488   EIGEN_DEVICE_FUNC
0489   static inline float run(const float& x)
0490   {
0491     return ::roundf(x);
0492   }
0493 };
0494 #else
0495 template<typename Scalar>
0496 struct round_using_floor_ceil_impl
0497 {
0498   EIGEN_DEVICE_FUNC
0499   static inline Scalar run(const Scalar& x)
0500   {
0501     EIGEN_STATIC_ASSERT((!NumTraits<Scalar>::IsComplex), NUMERIC_TYPE_MUST_BE_REAL)
0502     // Without C99 round/roundf, resort to floor/ceil.
0503     EIGEN_USING_STD(floor);
0504     EIGEN_USING_STD(ceil);
0505     // If not enough precision to resolve a decimal at all, return the input.
0506     // Otherwise, adding 0.5 can trigger an increment by 1.
0507     const Scalar limit = Scalar(1ull << (NumTraits<Scalar>::digits() - 1));
0508     if (x >= limit || x <= -limit) {
0509       return x;
0510     }
0511     return (x > Scalar(0)) ? Scalar(floor(x + Scalar(0.5))) : Scalar(ceil(x - Scalar(0.5)));
0512   }
0513 };
0514 
0515 template<>
0516 struct round_impl<float> : round_using_floor_ceil_impl<float> {};
0517 
0518 template<>
0519 struct round_impl<double> : round_using_floor_ceil_impl<double> {};
0520 #endif // EIGEN_HAS_C99_MATH
0521 #endif // !EIGEN_HAS_CXX11_MATH
0522 
0523 template<typename Scalar>
0524 struct round_retval
0525 {
0526   typedef Scalar type;
0527 };
0528 
0529 /****************************************************************************
0530 * Implementation of rint                                                    *
0531 ****************************************************************************/
0532 
0533 template<typename Scalar>
0534 struct rint_impl {
0535   EIGEN_DEVICE_FUNC
0536   static inline Scalar run(const Scalar& x)
0537   {
0538     EIGEN_STATIC_ASSERT((!NumTraits<Scalar>::IsComplex), NUMERIC_TYPE_MUST_BE_REAL)
0539 #if EIGEN_HAS_CXX11_MATH
0540       EIGEN_USING_STD(rint);
0541 #endif
0542     return rint(x);
0543   }
0544 };
0545 
0546 #if !EIGEN_HAS_CXX11_MATH
0547 template<>
0548 struct rint_impl<double> {
0549   EIGEN_DEVICE_FUNC
0550   static inline double run(const double& x)
0551   {
0552     return ::rint(x);
0553   }
0554 };
0555 template<>
0556 struct rint_impl<float> {
0557   EIGEN_DEVICE_FUNC
0558   static inline float run(const float& x)
0559   {
0560     return ::rintf(x);
0561   }
0562 };
0563 #endif
0564 
0565 template<typename Scalar>
0566 struct rint_retval
0567 {
0568   typedef Scalar type;
0569 };
0570 
0571 /****************************************************************************
0572 * Implementation of arg                                                     *
0573 ****************************************************************************/
0574 
0575 // Visual Studio 2017 has a bug where arg(float) returns 0 for negative inputs.
0576 // This seems to be fixed in VS 2019.
0577 #if EIGEN_HAS_CXX11_MATH && (!EIGEN_COMP_MSVC || EIGEN_COMP_MSVC >= 1920)
0578 // std::arg is only defined for types of std::complex, or integer types or float/double/long double
0579 template<typename Scalar,
0580           bool HasStdImpl = NumTraits<Scalar>::IsComplex || is_integral<Scalar>::value
0581                             || is_same<Scalar, float>::value || is_same<Scalar, double>::value
0582                             || is_same<Scalar, long double>::value >
0583 struct arg_default_impl;
0584 
0585 template<typename Scalar>
0586 struct arg_default_impl<Scalar, true> {
0587   typedef typename NumTraits<Scalar>::Real RealScalar;
0588   EIGEN_DEVICE_FUNC
0589   static inline RealScalar run(const Scalar& x)
0590   {
0591     #if defined(EIGEN_HIP_DEVICE_COMPILE)
0592     // HIP does not seem to have a native device side implementation for the math routine "arg"
0593     using std::arg;
0594     #else
0595     EIGEN_USING_STD(arg);
0596     #endif
0597     return static_cast<RealScalar>(arg(x));
0598   }
0599 };
0600 
0601 // Must be non-complex floating-point type (e.g. half/bfloat16).
0602 template<typename Scalar>
0603 struct arg_default_impl<Scalar, false> {
0604   typedef typename NumTraits<Scalar>::Real RealScalar;
0605   EIGEN_DEVICE_FUNC
0606   static inline RealScalar run(const Scalar& x)
0607   {
0608     return (x < Scalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
0609   }
0610 };
0611 #else
0612 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex>
0613 struct arg_default_impl
0614 {
0615   typedef typename NumTraits<Scalar>::Real RealScalar;
0616   EIGEN_DEVICE_FUNC
0617   static inline RealScalar run(const Scalar& x)
0618   {
0619     return (x < RealScalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0);
0620   }
0621 };
0622 
0623 template<typename Scalar>
0624 struct arg_default_impl<Scalar,true>
0625 {
0626   typedef typename NumTraits<Scalar>::Real RealScalar;
0627   EIGEN_DEVICE_FUNC
0628   static inline RealScalar run(const Scalar& x)
0629   {
0630     EIGEN_USING_STD(arg);
0631     return arg(x);
0632   }
0633 };
0634 #endif
0635 template<typename Scalar> struct arg_impl : arg_default_impl<Scalar> {};
0636 
0637 template<typename Scalar>
0638 struct arg_retval
0639 {
0640   typedef typename NumTraits<Scalar>::Real type;
0641 };
0642 
0643 /****************************************************************************
0644 * Implementation of expm1                                                   *
0645 ****************************************************************************/
0646 
0647 // This implementation is based on GSL Math's expm1.
0648 namespace std_fallback {
0649   // fallback expm1 implementation in case there is no expm1(Scalar) function in namespace of Scalar,
0650   // or that there is no suitable std::expm1 function available. Implementation
0651   // attributed to Kahan. See: http://www.plunk.org/~hatch/rightway.php.
0652   template<typename Scalar>
0653   EIGEN_DEVICE_FUNC inline Scalar expm1(const Scalar& x) {
0654     EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
0655     typedef typename NumTraits<Scalar>::Real RealScalar;
0656 
0657     EIGEN_USING_STD(exp);
0658     Scalar u = exp(x);
0659     if (numext::equal_strict(u, Scalar(1))) {
0660       return x;
0661     }
0662     Scalar um1 = u - RealScalar(1);
0663     if (numext::equal_strict(um1, Scalar(-1))) {
0664       return RealScalar(-1);
0665     }
0666 
0667     EIGEN_USING_STD(log);
0668     Scalar logu = log(u);
0669     return numext::equal_strict(u, logu) ? u : (u - RealScalar(1)) * x / logu;
0670   }
0671 }
0672 
0673 template<typename Scalar>
0674 struct expm1_impl {
0675   EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x)
0676   {
0677     EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
0678     #if EIGEN_HAS_CXX11_MATH
0679     using std::expm1;
0680     #else
0681     using std_fallback::expm1;
0682     #endif
0683     return expm1(x);
0684   }
0685 };
0686 
0687 template<typename Scalar>
0688 struct expm1_retval
0689 {
0690   typedef Scalar type;
0691 };
0692 
0693 /****************************************************************************
0694 * Implementation of log                                                     *
0695 ****************************************************************************/
0696 
0697 // Complex log defined in MathFunctionsImpl.h.
0698 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_log(const std::complex<T>& z);
0699 
0700 template<typename Scalar>
0701 struct log_impl {
0702   EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x)
0703   {
0704     EIGEN_USING_STD(log);
0705     return static_cast<Scalar>(log(x));
0706   }
0707 };
0708 
0709 template<typename Scalar>
0710 struct log_impl<std::complex<Scalar> > {
0711   EIGEN_DEVICE_FUNC static inline std::complex<Scalar> run(const std::complex<Scalar>& z)
0712   {
0713     return complex_log(z);
0714   }
0715 };
0716 
0717 /****************************************************************************
0718 * Implementation of log1p                                                   *
0719 ****************************************************************************/
0720 
0721 namespace std_fallback {
0722   // fallback log1p implementation in case there is no log1p(Scalar) function in namespace of Scalar,
0723   // or that there is no suitable std::log1p function available
0724   template<typename Scalar>
0725   EIGEN_DEVICE_FUNC inline Scalar log1p(const Scalar& x) {
0726     EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
0727     typedef typename NumTraits<Scalar>::Real RealScalar;
0728     EIGEN_USING_STD(log);
0729     Scalar x1p = RealScalar(1) + x;
0730     Scalar log_1p = log_impl<Scalar>::run(x1p);
0731     const bool is_small = numext::equal_strict(x1p, Scalar(1));
0732     const bool is_inf = numext::equal_strict(x1p, log_1p);
0733     return (is_small || is_inf) ? x : x * (log_1p / (x1p - RealScalar(1)));
0734   }
0735 }
0736 
0737 template<typename Scalar>
0738 struct log1p_impl {
0739   EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x)
0740   {
0741     EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar)
0742     #if EIGEN_HAS_CXX11_MATH
0743     using std::log1p;
0744     #else
0745     using std_fallback::log1p;
0746     #endif
0747     return log1p(x);
0748   }
0749 };
0750 
0751 // Specialization for complex types that are not supported by std::log1p.
0752 template <typename RealScalar>
0753 struct log1p_impl<std::complex<RealScalar> > {
0754   EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(
0755       const std::complex<RealScalar>& x) {
0756     EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
0757     return std_fallback::log1p(x);
0758   }
0759 };
0760 
0761 template<typename Scalar>
0762 struct log1p_retval
0763 {
0764   typedef Scalar type;
0765 };
0766 
0767 /****************************************************************************
0768 * Implementation of pow                                                  *
0769 ****************************************************************************/
0770 
0771 template<typename ScalarX,typename ScalarY, bool IsInteger = NumTraits<ScalarX>::IsInteger&&NumTraits<ScalarY>::IsInteger>
0772 struct pow_impl
0773 {
0774   //typedef Scalar retval;
0775   typedef typename ScalarBinaryOpTraits<ScalarX,ScalarY,internal::scalar_pow_op<ScalarX,ScalarY> >::ReturnType result_type;
0776   static EIGEN_DEVICE_FUNC inline result_type run(const ScalarX& x, const ScalarY& y)
0777   {
0778     EIGEN_USING_STD(pow);
0779     return pow(x, y);
0780   }
0781 };
0782 
0783 template<typename ScalarX,typename ScalarY>
0784 struct pow_impl<ScalarX,ScalarY, true>
0785 {
0786   typedef ScalarX result_type;
0787   static EIGEN_DEVICE_FUNC inline ScalarX run(ScalarX x, ScalarY y)
0788   {
0789     ScalarX res(1);
0790     eigen_assert(!NumTraits<ScalarY>::IsSigned || y >= 0);
0791     if(y & 1) res *= x;
0792     y >>= 1;
0793     while(y)
0794     {
0795       x *= x;
0796       if(y&1) res *= x;
0797       y >>= 1;
0798     }
0799     return res;
0800   }
0801 };
0802 
0803 /****************************************************************************
0804 * Implementation of random                                               *
0805 ****************************************************************************/
0806 
0807 template<typename Scalar,
0808          bool IsComplex,
0809          bool IsInteger>
0810 struct random_default_impl {};
0811 
0812 template<typename Scalar>
0813 struct random_impl : random_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {};
0814 
0815 template<typename Scalar>
0816 struct random_retval
0817 {
0818   typedef Scalar type;
0819 };
0820 
0821 template<typename Scalar> inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(const Scalar& x, const Scalar& y);
0822 template<typename Scalar> inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random();
0823 
0824 template<typename Scalar>
0825 struct random_default_impl<Scalar, false, false>
0826 {
0827   static inline Scalar run(const Scalar& x, const Scalar& y)
0828   {
0829     return x + (y-x) * Scalar(std::rand()) / Scalar(RAND_MAX);
0830   }
0831   static inline Scalar run()
0832   {
0833     return run(Scalar(NumTraits<Scalar>::IsSigned ? -1 : 0), Scalar(1));
0834   }
0835 };
0836 
0837 enum {
0838   meta_floor_log2_terminate,
0839   meta_floor_log2_move_up,
0840   meta_floor_log2_move_down,
0841   meta_floor_log2_bogus
0842 };
0843 
0844 template<unsigned int n, int lower, int upper> struct meta_floor_log2_selector
0845 {
0846   enum { middle = (lower + upper) / 2,
0847          value = (upper <= lower + 1) ? int(meta_floor_log2_terminate)
0848                : (n < (1 << middle)) ? int(meta_floor_log2_move_down)
0849                : (n==0) ? int(meta_floor_log2_bogus)
0850                : int(meta_floor_log2_move_up)
0851   };
0852 };
0853 
0854 template<unsigned int n,
0855          int lower = 0,
0856          int upper = sizeof(unsigned int) * CHAR_BIT - 1,
0857          int selector = meta_floor_log2_selector<n, lower, upper>::value>
0858 struct meta_floor_log2 {};
0859 
0860 template<unsigned int n, int lower, int upper>
0861 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_down>
0862 {
0863   enum { value = meta_floor_log2<n, lower, meta_floor_log2_selector<n, lower, upper>::middle>::value };
0864 };
0865 
0866 template<unsigned int n, int lower, int upper>
0867 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_up>
0868 {
0869   enum { value = meta_floor_log2<n, meta_floor_log2_selector<n, lower, upper>::middle, upper>::value };
0870 };
0871 
0872 template<unsigned int n, int lower, int upper>
0873 struct meta_floor_log2<n, lower, upper, meta_floor_log2_terminate>
0874 {
0875   enum { value = (n >= ((unsigned int)(1) << (lower+1))) ? lower+1 : lower };
0876 };
0877 
0878 template<unsigned int n, int lower, int upper>
0879 struct meta_floor_log2<n, lower, upper, meta_floor_log2_bogus>
0880 {
0881   // no value, error at compile time
0882 };
0883 
0884 template<typename Scalar>
0885 struct random_default_impl<Scalar, false, true>
0886 {
0887   static inline Scalar run(const Scalar& x, const Scalar& y)
0888   {
0889     if (y <= x)
0890       return x;
0891     // ScalarU is the unsigned counterpart of Scalar, possibly Scalar itself.
0892     typedef typename make_unsigned<Scalar>::type ScalarU;
0893     // ScalarX is the widest of ScalarU and unsigned int.
0894     // We'll deal only with ScalarX and unsigned int below thus avoiding signed
0895     // types and arithmetic and signed overflows (which are undefined behavior).
0896     typedef typename conditional<(ScalarU(-1) > unsigned(-1)), ScalarU, unsigned>::type ScalarX;
0897     // The following difference doesn't overflow, provided our integer types are two's
0898     // complement and have the same number of padding bits in signed and unsigned variants.
0899     // This is the case in most modern implementations of C++.
0900     ScalarX range = ScalarX(y) - ScalarX(x);
0901     ScalarX offset = 0;
0902     ScalarX divisor = 1;
0903     ScalarX multiplier = 1;
0904     const unsigned rand_max = RAND_MAX;
0905     if (range <= rand_max) divisor = (rand_max + 1) / (range + 1);
0906     else                   multiplier = 1 + range / (rand_max + 1);
0907     // Rejection sampling.
0908     do {
0909       offset = (unsigned(std::rand()) * multiplier) / divisor;
0910     } while (offset > range);
0911     return Scalar(ScalarX(x) + offset);
0912   }
0913 
0914   static inline Scalar run()
0915   {
0916 #ifdef EIGEN_MAKING_DOCS
0917     return run(Scalar(NumTraits<Scalar>::IsSigned ? -10 : 0), Scalar(10));
0918 #else
0919     enum { rand_bits = meta_floor_log2<(unsigned int)(RAND_MAX)+1>::value,
0920            scalar_bits = sizeof(Scalar) * CHAR_BIT,
0921            shift = EIGEN_PLAIN_ENUM_MAX(0, int(rand_bits) - int(scalar_bits)),
0922            offset = NumTraits<Scalar>::IsSigned ? (1 << (EIGEN_PLAIN_ENUM_MIN(rand_bits,scalar_bits)-1)) : 0
0923     };
0924     return Scalar((std::rand() >> shift) - offset);
0925 #endif
0926   }
0927 };
0928 
0929 template<typename Scalar>
0930 struct random_default_impl<Scalar, true, false>
0931 {
0932   static inline Scalar run(const Scalar& x, const Scalar& y)
0933   {
0934     return Scalar(random(x.real(), y.real()),
0935                   random(x.imag(), y.imag()));
0936   }
0937   static inline Scalar run()
0938   {
0939     typedef typename NumTraits<Scalar>::Real RealScalar;
0940     return Scalar(random<RealScalar>(), random<RealScalar>());
0941   }
0942 };
0943 
0944 template<typename Scalar>
0945 inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(const Scalar& x, const Scalar& y)
0946 {
0947   return EIGEN_MATHFUNC_IMPL(random, Scalar)::run(x, y);
0948 }
0949 
0950 template<typename Scalar>
0951 inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random()
0952 {
0953   return EIGEN_MATHFUNC_IMPL(random, Scalar)::run();
0954 }
0955 
0956 // Implementation of is* functions
0957 
0958 // std::is* do not work with fast-math and gcc, std::is* are available on MSVC 2013 and newer, as well as in clang.
0959 #if (EIGEN_HAS_CXX11_MATH && !(EIGEN_COMP_GNUC_STRICT && __FINITE_MATH_ONLY__)) || (EIGEN_COMP_MSVC>=1800) || (EIGEN_COMP_CLANG)
0960 #define EIGEN_USE_STD_FPCLASSIFY 1
0961 #else
0962 #define EIGEN_USE_STD_FPCLASSIFY 0
0963 #endif
0964 
0965 template<typename T>
0966 EIGEN_DEVICE_FUNC
0967 typename internal::enable_if<internal::is_integral<T>::value,bool>::type
0968 isnan_impl(const T&) { return false; }
0969 
0970 template<typename T>
0971 EIGEN_DEVICE_FUNC
0972 typename internal::enable_if<internal::is_integral<T>::value,bool>::type
0973 isinf_impl(const T&) { return false; }
0974 
0975 template<typename T>
0976 EIGEN_DEVICE_FUNC
0977 typename internal::enable_if<internal::is_integral<T>::value,bool>::type
0978 isfinite_impl(const T&) { return true; }
0979 
0980 template<typename T>
0981 EIGEN_DEVICE_FUNC
0982 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type
0983 isfinite_impl(const T& x)
0984 {
0985   #if defined(EIGEN_GPU_COMPILE_PHASE)
0986     return (::isfinite)(x);
0987   #elif EIGEN_USE_STD_FPCLASSIFY
0988     using std::isfinite;
0989     return isfinite EIGEN_NOT_A_MACRO (x);
0990   #else
0991     return x<=NumTraits<T>::highest() && x>=NumTraits<T>::lowest();
0992   #endif
0993 }
0994 
0995 template<typename T>
0996 EIGEN_DEVICE_FUNC
0997 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type
0998 isinf_impl(const T& x)
0999 {
1000   #if defined(EIGEN_GPU_COMPILE_PHASE)
1001     return (::isinf)(x);
1002   #elif EIGEN_USE_STD_FPCLASSIFY
1003     using std::isinf;
1004     return isinf EIGEN_NOT_A_MACRO (x);
1005   #else
1006     return x>NumTraits<T>::highest() || x<NumTraits<T>::lowest();
1007   #endif
1008 }
1009 
1010 template<typename T>
1011 EIGEN_DEVICE_FUNC
1012 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type
1013 isnan_impl(const T& x)
1014 {
1015   #if defined(EIGEN_GPU_COMPILE_PHASE)
1016     return (::isnan)(x);
1017   #elif EIGEN_USE_STD_FPCLASSIFY
1018     using std::isnan;
1019     return isnan EIGEN_NOT_A_MACRO (x);
1020   #else
1021     return x != x;
1022   #endif
1023 }
1024 
1025 #if (!EIGEN_USE_STD_FPCLASSIFY)
1026 
1027 #if EIGEN_COMP_MSVC
1028 
1029 template<typename T> EIGEN_DEVICE_FUNC bool isinf_msvc_helper(T x)
1030 {
1031   return _fpclass(x)==_FPCLASS_NINF || _fpclass(x)==_FPCLASS_PINF;
1032 }
1033 
1034 //MSVC defines a _isnan builtin function, but for double only
1035 EIGEN_DEVICE_FUNC inline bool isnan_impl(const long double& x) { return _isnan(x)!=0; }
1036 EIGEN_DEVICE_FUNC inline bool isnan_impl(const double& x)      { return _isnan(x)!=0; }
1037 EIGEN_DEVICE_FUNC inline bool isnan_impl(const float& x)       { return _isnan(x)!=0; }
1038 
1039 EIGEN_DEVICE_FUNC inline bool isinf_impl(const long double& x) { return isinf_msvc_helper(x); }
1040 EIGEN_DEVICE_FUNC inline bool isinf_impl(const double& x)      { return isinf_msvc_helper(x); }
1041 EIGEN_DEVICE_FUNC inline bool isinf_impl(const float& x)       { return isinf_msvc_helper(x); }
1042 
1043 #elif (defined __FINITE_MATH_ONLY__ && __FINITE_MATH_ONLY__ && EIGEN_COMP_GNUC)
1044 
1045 #if EIGEN_GNUC_AT_LEAST(5,0)
1046   #define EIGEN_TMP_NOOPT_ATTRIB EIGEN_DEVICE_FUNC inline __attribute__((optimize("no-finite-math-only")))
1047 #else
1048   // NOTE the inline qualifier and noinline attribute are both needed: the former is to avoid linking issue (duplicate symbol),
1049   //      while the second prevent too aggressive optimizations in fast-math mode:
1050   #define EIGEN_TMP_NOOPT_ATTRIB EIGEN_DEVICE_FUNC inline __attribute__((noinline,optimize("no-finite-math-only")))
1051 #endif
1052 
1053 template<> EIGEN_TMP_NOOPT_ATTRIB bool isnan_impl(const long double& x) { return __builtin_isnan(x); }
1054 template<> EIGEN_TMP_NOOPT_ATTRIB bool isnan_impl(const double& x)      { return __builtin_isnan(x); }
1055 template<> EIGEN_TMP_NOOPT_ATTRIB bool isnan_impl(const float& x)       { return __builtin_isnan(x); }
1056 template<> EIGEN_TMP_NOOPT_ATTRIB bool isinf_impl(const double& x)      { return __builtin_isinf(x); }
1057 template<> EIGEN_TMP_NOOPT_ATTRIB bool isinf_impl(const float& x)       { return __builtin_isinf(x); }
1058 template<> EIGEN_TMP_NOOPT_ATTRIB bool isinf_impl(const long double& x) { return __builtin_isinf(x); }
1059 
1060 #undef EIGEN_TMP_NOOPT_ATTRIB
1061 
1062 #endif
1063 
1064 #endif
1065 
1066 // The following overload are defined at the end of this file
1067 template<typename T> EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x);
1068 template<typename T> EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x);
1069 template<typename T> EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x);
1070 
1071 template<typename T> T generic_fast_tanh_float(const T& a_x);
1072 } // end namespace internal
1073 
1074 /****************************************************************************
1075 * Generic math functions                                                    *
1076 ****************************************************************************/
1077 
1078 namespace numext {
1079 
1080 #if (!defined(EIGEN_GPUCC) || defined(EIGEN_CONSTEXPR_ARE_DEVICE_FUNC))
1081 template<typename T>
1082 EIGEN_DEVICE_FUNC
1083 EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y)
1084 {
1085   EIGEN_USING_STD(min)
1086   return min EIGEN_NOT_A_MACRO (x,y);
1087 }
1088 
1089 template<typename T>
1090 EIGEN_DEVICE_FUNC
1091 EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y)
1092 {
1093   EIGEN_USING_STD(max)
1094   return max EIGEN_NOT_A_MACRO (x,y);
1095 }
1096 #else
1097 template<typename T>
1098 EIGEN_DEVICE_FUNC
1099 EIGEN_ALWAYS_INLINE T mini(const T& x, const T& y)
1100 {
1101   return y < x ? y : x;
1102 }
1103 template<>
1104 EIGEN_DEVICE_FUNC
1105 EIGEN_ALWAYS_INLINE float mini(const float& x, const float& y)
1106 {
1107   return fminf(x, y);
1108 }
1109 template<>
1110 EIGEN_DEVICE_FUNC
1111 EIGEN_ALWAYS_INLINE double mini(const double& x, const double& y)
1112 {
1113   return fmin(x, y);
1114 }
1115 template<>
1116 EIGEN_DEVICE_FUNC
1117 EIGEN_ALWAYS_INLINE long double mini(const long double& x, const long double& y)
1118 {
1119 #if defined(EIGEN_HIPCC)
1120   // no "fminl" on HIP yet
1121   return (x < y) ? x : y;
1122 #else
1123   return fminl(x, y);
1124 #endif
1125 }
1126 
1127 template<typename T>
1128 EIGEN_DEVICE_FUNC
1129 EIGEN_ALWAYS_INLINE T maxi(const T& x, const T& y)
1130 {
1131   return x < y ? y : x;
1132 }
1133 template<>
1134 EIGEN_DEVICE_FUNC
1135 EIGEN_ALWAYS_INLINE float maxi(const float& x, const float& y)
1136 {
1137   return fmaxf(x, y);
1138 }
1139 template<>
1140 EIGEN_DEVICE_FUNC
1141 EIGEN_ALWAYS_INLINE double maxi(const double& x, const double& y)
1142 {
1143   return fmax(x, y);
1144 }
1145 template<>
1146 EIGEN_DEVICE_FUNC
1147 EIGEN_ALWAYS_INLINE long double maxi(const long double& x, const long double& y)
1148 {
1149 #if defined(EIGEN_HIPCC)
1150   // no "fmaxl" on HIP yet
1151   return (x > y) ? x : y;
1152 #else
1153   return fmaxl(x, y);
1154 #endif
1155 }
1156 #endif
1157 
1158 #if defined(SYCL_DEVICE_ONLY)
1159 
1160 
1161 #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1162   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_char)   \
1163   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_short)  \
1164   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_int)    \
1165   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1166 #define SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1167   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_char)   \
1168   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_short)  \
1169   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_int)    \
1170   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_long)
1171 #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1172   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar)  \
1173   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1174   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_uint)   \
1175   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1176 #define SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1177   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uchar)  \
1178   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ushort) \
1179   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_uint)   \
1180   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_ulong)
1181 #define SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(NAME, FUNC) \
1182   SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY(NAME, FUNC) \
1183   SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY(NAME, FUNC)
1184 #define SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(NAME, FUNC) \
1185   SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY(NAME, FUNC) \
1186   SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY(NAME, FUNC)
1187 #define SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(NAME, FUNC) \
1188   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1189   SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC,cl::sycl::cl_double)
1190 #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(NAME, FUNC) \
1191   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, cl::sycl::cl_float) \
1192   SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC,cl::sycl::cl_double)
1193 #define SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(NAME, FUNC, RET_TYPE) \
1194   SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_float) \
1195   SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, cl::sycl::cl_double)
1196 
1197 #define SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1198 template<>                                               \
1199   EIGEN_DEVICE_FUNC                                      \
1200   EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE& x) { \
1201     return cl::sycl::FUNC(x);                            \
1202   }
1203 
1204 #define SYCL_SPECIALIZE_UNARY_FUNC(NAME, FUNC, TYPE) \
1205   SYCL_SPECIALIZE_GEN_UNARY_FUNC(NAME, FUNC, TYPE, TYPE)
1206 
1207 #define SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE1, ARG_TYPE2) \
1208   template<>                                                                  \
1209   EIGEN_DEVICE_FUNC                                                           \
1210   EIGEN_ALWAYS_INLINE RET_TYPE NAME(const ARG_TYPE1& x, const ARG_TYPE2& y) { \
1211     return cl::sycl::FUNC(x, y);                                              \
1212   }
1213 
1214 #define SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE) \
1215   SYCL_SPECIALIZE_GEN1_BINARY_FUNC(NAME, FUNC, RET_TYPE, ARG_TYPE, ARG_TYPE)
1216 
1217 #define SYCL_SPECIALIZE_BINARY_FUNC(NAME, FUNC, TYPE) \
1218   SYCL_SPECIALIZE_GEN2_BINARY_FUNC(NAME, FUNC, TYPE, TYPE)
1219 
1220 SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(mini, min)
1221 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(mini, fmin)
1222 SYCL_SPECIALIZE_INTEGER_TYPES_BINARY(maxi, max)
1223 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(maxi, fmax)
1224 
1225 #endif
1226 
1227 
1228 template<typename Scalar>
1229 EIGEN_DEVICE_FUNC
1230 inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x)
1231 {
1232   return EIGEN_MATHFUNC_IMPL(real, Scalar)::run(x);
1233 }
1234 
1235 template<typename Scalar>
1236 EIGEN_DEVICE_FUNC
1237 inline typename internal::add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) >::type real_ref(const Scalar& x)
1238 {
1239   return internal::real_ref_impl<Scalar>::run(x);
1240 }
1241 
1242 template<typename Scalar>
1243 EIGEN_DEVICE_FUNC
1244 inline EIGEN_MATHFUNC_RETVAL(real_ref, Scalar) real_ref(Scalar& x)
1245 {
1246   return EIGEN_MATHFUNC_IMPL(real_ref, Scalar)::run(x);
1247 }
1248 
1249 template<typename Scalar>
1250 EIGEN_DEVICE_FUNC
1251 inline EIGEN_MATHFUNC_RETVAL(imag, Scalar) imag(const Scalar& x)
1252 {
1253   return EIGEN_MATHFUNC_IMPL(imag, Scalar)::run(x);
1254 }
1255 
1256 template<typename Scalar>
1257 EIGEN_DEVICE_FUNC
1258 inline EIGEN_MATHFUNC_RETVAL(arg, Scalar) arg(const Scalar& x)
1259 {
1260   return EIGEN_MATHFUNC_IMPL(arg, Scalar)::run(x);
1261 }
1262 
1263 template<typename Scalar>
1264 EIGEN_DEVICE_FUNC
1265 inline typename internal::add_const_on_value_type< EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) >::type imag_ref(const Scalar& x)
1266 {
1267   return internal::imag_ref_impl<Scalar>::run(x);
1268 }
1269 
1270 template<typename Scalar>
1271 EIGEN_DEVICE_FUNC
1272 inline EIGEN_MATHFUNC_RETVAL(imag_ref, Scalar) imag_ref(Scalar& x)
1273 {
1274   return EIGEN_MATHFUNC_IMPL(imag_ref, Scalar)::run(x);
1275 }
1276 
1277 template<typename Scalar>
1278 EIGEN_DEVICE_FUNC
1279 inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x)
1280 {
1281   return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x);
1282 }
1283 
1284 template<typename Scalar>
1285 EIGEN_DEVICE_FUNC
1286 inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x)
1287 {
1288   return EIGEN_MATHFUNC_IMPL(abs2, Scalar)::run(x);
1289 }
1290 
1291 EIGEN_DEVICE_FUNC
1292 inline bool abs2(bool x) { return x; }
1293 
1294 template<typename T>
1295 EIGEN_DEVICE_FUNC
1296 EIGEN_ALWAYS_INLINE T absdiff(const T& x, const T& y)
1297 {
1298   return x > y ? x - y : y - x;
1299 }
1300 template<>
1301 EIGEN_DEVICE_FUNC
1302 EIGEN_ALWAYS_INLINE float absdiff(const float& x, const float& y)
1303 {
1304   return fabsf(x - y);
1305 }
1306 template<>
1307 EIGEN_DEVICE_FUNC
1308 EIGEN_ALWAYS_INLINE double absdiff(const double& x, const double& y)
1309 {
1310   return fabs(x - y);
1311 }
1312 
1313 #if !defined(EIGEN_GPUCC)
1314 // HIP and CUDA do not support long double.
1315 template<>
1316 EIGEN_DEVICE_FUNC
1317 EIGEN_ALWAYS_INLINE long double absdiff(const long double& x, const long double& y) {
1318   return fabsl(x - y);
1319 }
1320 #endif
1321 
1322 template<typename Scalar>
1323 EIGEN_DEVICE_FUNC
1324 inline EIGEN_MATHFUNC_RETVAL(norm1, Scalar) norm1(const Scalar& x)
1325 {
1326   return EIGEN_MATHFUNC_IMPL(norm1, Scalar)::run(x);
1327 }
1328 
1329 template<typename Scalar>
1330 EIGEN_DEVICE_FUNC
1331 inline EIGEN_MATHFUNC_RETVAL(hypot, Scalar) hypot(const Scalar& x, const Scalar& y)
1332 {
1333   return EIGEN_MATHFUNC_IMPL(hypot, Scalar)::run(x, y);
1334 }
1335 
1336 #if defined(SYCL_DEVICE_ONLY)
1337   SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(hypot, hypot)
1338 #endif
1339 
1340 template<typename Scalar>
1341 EIGEN_DEVICE_FUNC
1342 inline EIGEN_MATHFUNC_RETVAL(log1p, Scalar) log1p(const Scalar& x)
1343 {
1344   return EIGEN_MATHFUNC_IMPL(log1p, Scalar)::run(x);
1345 }
1346 
1347 #if defined(SYCL_DEVICE_ONLY)
1348 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log1p, log1p)
1349 #endif
1350 
1351 #if defined(EIGEN_GPUCC)
1352 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1353 float log1p(const float &x) { return ::log1pf(x); }
1354 
1355 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1356 double log1p(const double &x) { return ::log1p(x); }
1357 #endif
1358 
1359 template<typename ScalarX,typename ScalarY>
1360 EIGEN_DEVICE_FUNC
1361 inline typename internal::pow_impl<ScalarX,ScalarY>::result_type pow(const ScalarX& x, const ScalarY& y)
1362 {
1363   return internal::pow_impl<ScalarX,ScalarY>::run(x, y);
1364 }
1365 
1366 #if defined(SYCL_DEVICE_ONLY)
1367 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(pow, pow)
1368 #endif
1369 
1370 template<typename T> EIGEN_DEVICE_FUNC bool (isnan)   (const T &x) { return internal::isnan_impl(x); }
1371 template<typename T> EIGEN_DEVICE_FUNC bool (isinf)   (const T &x) { return internal::isinf_impl(x); }
1372 template<typename T> EIGEN_DEVICE_FUNC bool (isfinite)(const T &x) { return internal::isfinite_impl(x); }
1373 
1374 #if defined(SYCL_DEVICE_ONLY)
1375 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isnan, isnan, bool)
1376 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isinf, isinf, bool)
1377 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE(isfinite, isfinite, bool)
1378 #endif
1379 
1380 template<typename Scalar>
1381 EIGEN_DEVICE_FUNC
1382 inline EIGEN_MATHFUNC_RETVAL(rint, Scalar) rint(const Scalar& x)
1383 {
1384   return EIGEN_MATHFUNC_IMPL(rint, Scalar)::run(x);
1385 }
1386 
1387 template<typename Scalar>
1388 EIGEN_DEVICE_FUNC
1389 inline EIGEN_MATHFUNC_RETVAL(round, Scalar) round(const Scalar& x)
1390 {
1391   return EIGEN_MATHFUNC_IMPL(round, Scalar)::run(x);
1392 }
1393 
1394 #if defined(SYCL_DEVICE_ONLY)
1395 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(round, round)
1396 #endif
1397 
1398 template<typename T>
1399 EIGEN_DEVICE_FUNC
1400 T (floor)(const T& x)
1401 {
1402   EIGEN_USING_STD(floor)
1403   return floor(x);
1404 }
1405 
1406 #if defined(SYCL_DEVICE_ONLY)
1407 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(floor, floor)
1408 #endif
1409 
1410 #if defined(EIGEN_GPUCC)
1411 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1412 float floor(const float &x) { return ::floorf(x); }
1413 
1414 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1415 double floor(const double &x) { return ::floor(x); }
1416 #endif
1417 
1418 template<typename T>
1419 EIGEN_DEVICE_FUNC
1420 T (ceil)(const T& x)
1421 {
1422   EIGEN_USING_STD(ceil);
1423   return ceil(x);
1424 }
1425 
1426 #if defined(SYCL_DEVICE_ONLY)
1427 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(ceil, ceil)
1428 #endif
1429 
1430 #if defined(EIGEN_GPUCC)
1431 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1432 float ceil(const float &x) { return ::ceilf(x); }
1433 
1434 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1435 double ceil(const double &x) { return ::ceil(x); }
1436 #endif
1437 
1438 
1439 /** Log base 2 for 32 bits positive integers.
1440   * Conveniently returns 0 for x==0. */
1441 inline int log2(int x)
1442 {
1443   eigen_assert(x>=0);
1444   unsigned int v(x);
1445   static const int table[32] = { 0, 9, 1, 10, 13, 21, 2, 29, 11, 14, 16, 18, 22, 25, 3, 30, 8, 12, 20, 28, 15, 17, 24, 7, 19, 27, 23, 6, 26, 5, 4, 31 };
1446   v |= v >> 1;
1447   v |= v >> 2;
1448   v |= v >> 4;
1449   v |= v >> 8;
1450   v |= v >> 16;
1451   return table[(v * 0x07C4ACDDU) >> 27];
1452 }
1453 
1454 /** \returns the square root of \a x.
1455   *
1456   * It is essentially equivalent to
1457   * \code using std::sqrt; return sqrt(x); \endcode
1458   * but slightly faster for float/double and some compilers (e.g., gcc), thanks to
1459   * specializations when SSE is enabled.
1460   *
1461   * It's usage is justified in performance critical functions, like norm/normalize.
1462   */
1463 template<typename Scalar>
1464 EIGEN_DEVICE_FUNC
1465 EIGEN_ALWAYS_INLINE EIGEN_MATHFUNC_RETVAL(sqrt, Scalar) sqrt(const Scalar& x)
1466 {
1467   return EIGEN_MATHFUNC_IMPL(sqrt, Scalar)::run(x);
1468 }
1469 
1470 // Boolean specialization, avoids implicit float to bool conversion (-Wimplicit-conversion-floating-point-to-bool).
1471 template<>
1472 EIGEN_DEFINE_FUNCTION_ALLOWING_MULTIPLE_DEFINITIONS EIGEN_DEVICE_FUNC
1473 bool sqrt<bool>(const bool &x) { return x; }
1474 
1475 #if defined(SYCL_DEVICE_ONLY)
1476 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sqrt, sqrt)
1477 #endif
1478 
1479 /** \returns the reciprocal square root of \a x. **/
1480 template<typename T>
1481 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1482 T rsqrt(const T& x)
1483 {
1484   return internal::rsqrt_impl<T>::run(x);
1485 }
1486 
1487 template<typename T>
1488 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1489 T log(const T &x) {
1490   return internal::log_impl<T>::run(x);
1491 }
1492 
1493 #if defined(SYCL_DEVICE_ONLY)
1494 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(log, log)
1495 #endif
1496 
1497 
1498 #if defined(EIGEN_GPUCC)
1499 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1500 float log(const float &x) { return ::logf(x); }
1501 
1502 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1503 double log(const double &x) { return ::log(x); }
1504 #endif
1505 
1506 template<typename T>
1507 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1508 typename internal::enable_if<NumTraits<T>::IsSigned || NumTraits<T>::IsComplex,typename NumTraits<T>::Real>::type
1509 abs(const T &x) {
1510   EIGEN_USING_STD(abs);
1511   return abs(x);
1512 }
1513 
1514 template<typename T>
1515 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1516 typename internal::enable_if<!(NumTraits<T>::IsSigned || NumTraits<T>::IsComplex),typename NumTraits<T>::Real>::type
1517 abs(const T &x) {
1518   return x;
1519 }
1520 
1521 #if defined(SYCL_DEVICE_ONLY)
1522 SYCL_SPECIALIZE_INTEGER_TYPES_UNARY(abs, abs)
1523 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(abs, fabs)
1524 #endif
1525 
1526 #if defined(EIGEN_GPUCC)
1527 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1528 float abs(const float &x) { return ::fabsf(x); }
1529 
1530 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1531 double abs(const double &x) { return ::fabs(x); }
1532 
1533 template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1534 float abs(const std::complex<float>& x) {
1535   return ::hypotf(x.real(), x.imag());
1536 }
1537 
1538 template <> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1539 double abs(const std::complex<double>& x) {
1540   return ::hypot(x.real(), x.imag());
1541 }
1542 #endif
1543 
1544 template<typename T>
1545 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1546 T exp(const T &x) {
1547   EIGEN_USING_STD(exp);
1548   return exp(x);
1549 }
1550 
1551 #if defined(SYCL_DEVICE_ONLY)
1552 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(exp, exp)
1553 #endif
1554 
1555 #if defined(EIGEN_GPUCC)
1556 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1557 float exp(const float &x) { return ::expf(x); }
1558 
1559 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1560 double exp(const double &x) { return ::exp(x); }
1561 
1562 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1563 std::complex<float> exp(const std::complex<float>& x) {
1564   float com = ::expf(x.real());
1565   float res_real = com * ::cosf(x.imag());
1566   float res_imag = com * ::sinf(x.imag());
1567   return std::complex<float>(res_real, res_imag);
1568 }
1569 
1570 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1571 std::complex<double> exp(const std::complex<double>& x) {
1572   double com = ::exp(x.real());
1573   double res_real = com * ::cos(x.imag());
1574   double res_imag = com * ::sin(x.imag());
1575   return std::complex<double>(res_real, res_imag);
1576 }
1577 #endif
1578 
1579 template<typename Scalar>
1580 EIGEN_DEVICE_FUNC
1581 inline EIGEN_MATHFUNC_RETVAL(expm1, Scalar) expm1(const Scalar& x)
1582 {
1583   return EIGEN_MATHFUNC_IMPL(expm1, Scalar)::run(x);
1584 }
1585 
1586 #if defined(SYCL_DEVICE_ONLY)
1587 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(expm1, expm1)
1588 #endif
1589 
1590 #if defined(EIGEN_GPUCC)
1591 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1592 float expm1(const float &x) { return ::expm1f(x); }
1593 
1594 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1595 double expm1(const double &x) { return ::expm1(x); }
1596 #endif
1597 
1598 template<typename T>
1599 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1600 T cos(const T &x) {
1601   EIGEN_USING_STD(cos);
1602   return cos(x);
1603 }
1604 
1605 #if defined(SYCL_DEVICE_ONLY)
1606 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cos,cos)
1607 #endif
1608 
1609 #if defined(EIGEN_GPUCC)
1610 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1611 float cos(const float &x) { return ::cosf(x); }
1612 
1613 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1614 double cos(const double &x) { return ::cos(x); }
1615 #endif
1616 
1617 template<typename T>
1618 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1619 T sin(const T &x) {
1620   EIGEN_USING_STD(sin);
1621   return sin(x);
1622 }
1623 
1624 #if defined(SYCL_DEVICE_ONLY)
1625 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sin, sin)
1626 #endif
1627 
1628 #if defined(EIGEN_GPUCC)
1629 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1630 float sin(const float &x) { return ::sinf(x); }
1631 
1632 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1633 double sin(const double &x) { return ::sin(x); }
1634 #endif
1635 
1636 template<typename T>
1637 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1638 T tan(const T &x) {
1639   EIGEN_USING_STD(tan);
1640   return tan(x);
1641 }
1642 
1643 #if defined(SYCL_DEVICE_ONLY)
1644 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tan, tan)
1645 #endif
1646 
1647 #if defined(EIGEN_GPUCC)
1648 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1649 float tan(const float &x) { return ::tanf(x); }
1650 
1651 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1652 double tan(const double &x) { return ::tan(x); }
1653 #endif
1654 
1655 template<typename T>
1656 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1657 T acos(const T &x) {
1658   EIGEN_USING_STD(acos);
1659   return acos(x);
1660 }
1661 
1662 #if EIGEN_HAS_CXX11_MATH
1663 template<typename T>
1664 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1665 T acosh(const T &x) {
1666   EIGEN_USING_STD(acosh);
1667   return static_cast<T>(acosh(x));
1668 }
1669 #endif
1670 
1671 #if defined(SYCL_DEVICE_ONLY)
1672 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acos, acos)
1673 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(acosh, acosh)
1674 #endif
1675 
1676 #if defined(EIGEN_GPUCC)
1677 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1678 float acos(const float &x) { return ::acosf(x); }
1679 
1680 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1681 double acos(const double &x) { return ::acos(x); }
1682 #endif
1683 
1684 template<typename T>
1685 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1686 T asin(const T &x) {
1687   EIGEN_USING_STD(asin);
1688   return asin(x);
1689 }
1690 
1691 #if EIGEN_HAS_CXX11_MATH
1692 template<typename T>
1693 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1694 T asinh(const T &x) {
1695   EIGEN_USING_STD(asinh);
1696   return static_cast<T>(asinh(x));
1697 }
1698 #endif
1699 
1700 #if defined(SYCL_DEVICE_ONLY)
1701 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asin, asin)
1702 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(asinh, asinh)
1703 #endif
1704 
1705 #if defined(EIGEN_GPUCC)
1706 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1707 float asin(const float &x) { return ::asinf(x); }
1708 
1709 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1710 double asin(const double &x) { return ::asin(x); }
1711 #endif
1712 
1713 template<typename T>
1714 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1715 T atan(const T &x) {
1716   EIGEN_USING_STD(atan);
1717   return static_cast<T>(atan(x));
1718 }
1719 
1720 #if EIGEN_HAS_CXX11_MATH
1721 template<typename T>
1722 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1723 T atanh(const T &x) {
1724   EIGEN_USING_STD(atanh);
1725   return static_cast<T>(atanh(x));
1726 }
1727 #endif
1728 
1729 #if defined(SYCL_DEVICE_ONLY)
1730 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atan, atan)
1731 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(atanh, atanh)
1732 #endif
1733 
1734 #if defined(EIGEN_GPUCC)
1735 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1736 float atan(const float &x) { return ::atanf(x); }
1737 
1738 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1739 double atan(const double &x) { return ::atan(x); }
1740 #endif
1741 
1742 
1743 template<typename T>
1744 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1745 T cosh(const T &x) {
1746   EIGEN_USING_STD(cosh);
1747   return static_cast<T>(cosh(x));
1748 }
1749 
1750 #if defined(SYCL_DEVICE_ONLY)
1751 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(cosh, cosh)
1752 #endif
1753 
1754 #if defined(EIGEN_GPUCC)
1755 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1756 float cosh(const float &x) { return ::coshf(x); }
1757 
1758 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1759 double cosh(const double &x) { return ::cosh(x); }
1760 #endif
1761 
1762 template<typename T>
1763 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1764 T sinh(const T &x) {
1765   EIGEN_USING_STD(sinh);
1766   return static_cast<T>(sinh(x));
1767 }
1768 
1769 #if defined(SYCL_DEVICE_ONLY)
1770 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(sinh, sinh)
1771 #endif
1772 
1773 #if defined(EIGEN_GPUCC)
1774 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1775 float sinh(const float &x) { return ::sinhf(x); }
1776 
1777 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1778 double sinh(const double &x) { return ::sinh(x); }
1779 #endif
1780 
1781 template<typename T>
1782 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1783 T tanh(const T &x) {
1784   EIGEN_USING_STD(tanh);
1785   return tanh(x);
1786 }
1787 
1788 #if (!defined(EIGEN_GPUCC)) && EIGEN_FAST_MATH && !defined(SYCL_DEVICE_ONLY)
1789 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1790 float tanh(float x) { return internal::generic_fast_tanh_float(x); }
1791 #endif
1792 
1793 #if defined(SYCL_DEVICE_ONLY)
1794 SYCL_SPECIALIZE_FLOATING_TYPES_UNARY(tanh, tanh)
1795 #endif
1796 
1797 #if defined(EIGEN_GPUCC)
1798 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1799 float tanh(const float &x) { return ::tanhf(x); }
1800 
1801 template<> EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1802 double tanh(const double &x) { return ::tanh(x); }
1803 #endif
1804 
1805 template <typename T>
1806 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1807 T fmod(const T& a, const T& b) {
1808   EIGEN_USING_STD(fmod);
1809   return fmod(a, b);
1810 }
1811 
1812 #if defined(SYCL_DEVICE_ONLY)
1813 SYCL_SPECIALIZE_FLOATING_TYPES_BINARY(fmod, fmod)
1814 #endif
1815 
1816 #if defined(EIGEN_GPUCC)
1817 template <>
1818 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1819 float fmod(const float& a, const float& b) {
1820   return ::fmodf(a, b);
1821 }
1822 
1823 template <>
1824 EIGEN_DEVICE_FUNC EIGEN_ALWAYS_INLINE
1825 double fmod(const double& a, const double& b) {
1826   return ::fmod(a, b);
1827 }
1828 #endif
1829 
1830 #if defined(SYCL_DEVICE_ONLY)
1831 #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_BINARY
1832 #undef SYCL_SPECIALIZE_SIGNED_INTEGER_TYPES_UNARY
1833 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_BINARY
1834 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1835 #undef SYCL_SPECIALIZE_INTEGER_TYPES_BINARY
1836 #undef SYCL_SPECIALIZE_UNSIGNED_INTEGER_TYPES_UNARY
1837 #undef SYCL_SPECIALIZE_FLOATING_TYPES_BINARY
1838 #undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY
1839 #undef SYCL_SPECIALIZE_FLOATING_TYPES_UNARY_FUNC_RET_TYPE
1840 #undef SYCL_SPECIALIZE_GEN_UNARY_FUNC
1841 #undef SYCL_SPECIALIZE_UNARY_FUNC
1842 #undef SYCL_SPECIALIZE_GEN1_BINARY_FUNC
1843 #undef SYCL_SPECIALIZE_GEN2_BINARY_FUNC
1844 #undef SYCL_SPECIALIZE_BINARY_FUNC
1845 #endif
1846 
1847 } // end namespace numext
1848 
1849 namespace internal {
1850 
1851 template<typename T>
1852 EIGEN_DEVICE_FUNC bool isfinite_impl(const std::complex<T>& x)
1853 {
1854   return (numext::isfinite)(numext::real(x)) && (numext::isfinite)(numext::imag(x));
1855 }
1856 
1857 template<typename T>
1858 EIGEN_DEVICE_FUNC bool isnan_impl(const std::complex<T>& x)
1859 {
1860   return (numext::isnan)(numext::real(x)) || (numext::isnan)(numext::imag(x));
1861 }
1862 
1863 template<typename T>
1864 EIGEN_DEVICE_FUNC bool isinf_impl(const std::complex<T>& x)
1865 {
1866   return ((numext::isinf)(numext::real(x)) || (numext::isinf)(numext::imag(x))) && (!(numext::isnan)(x));
1867 }
1868 
1869 /****************************************************************************
1870 * Implementation of fuzzy comparisons                                       *
1871 ****************************************************************************/
1872 
1873 template<typename Scalar,
1874          bool IsComplex,
1875          bool IsInteger>
1876 struct scalar_fuzzy_default_impl {};
1877 
1878 template<typename Scalar>
1879 struct scalar_fuzzy_default_impl<Scalar, false, false>
1880 {
1881   typedef typename NumTraits<Scalar>::Real RealScalar;
1882   template<typename OtherScalar> EIGEN_DEVICE_FUNC
1883   static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y, const RealScalar& prec)
1884   {
1885     return numext::abs(x) <= numext::abs(y) * prec;
1886   }
1887   EIGEN_DEVICE_FUNC
1888   static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec)
1889   {
1890     return numext::abs(x - y) <= numext::mini(numext::abs(x), numext::abs(y)) * prec;
1891   }
1892   EIGEN_DEVICE_FUNC
1893   static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar& prec)
1894   {
1895     return x <= y || isApprox(x, y, prec);
1896   }
1897 };
1898 
1899 template<typename Scalar>
1900 struct scalar_fuzzy_default_impl<Scalar, false, true>
1901 {
1902   typedef typename NumTraits<Scalar>::Real RealScalar;
1903   template<typename OtherScalar> EIGEN_DEVICE_FUNC
1904   static inline bool isMuchSmallerThan(const Scalar& x, const Scalar&, const RealScalar&)
1905   {
1906     return x == Scalar(0);
1907   }
1908   EIGEN_DEVICE_FUNC
1909   static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar&)
1910   {
1911     return x == y;
1912   }
1913   EIGEN_DEVICE_FUNC
1914   static inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y, const RealScalar&)
1915   {
1916     return x <= y;
1917   }
1918 };
1919 
1920 template<typename Scalar>
1921 struct scalar_fuzzy_default_impl<Scalar, true, false>
1922 {
1923   typedef typename NumTraits<Scalar>::Real RealScalar;
1924   template<typename OtherScalar> EIGEN_DEVICE_FUNC
1925   static inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y, const RealScalar& prec)
1926   {
1927     return numext::abs2(x) <= numext::abs2(y) * prec * prec;
1928   }
1929   EIGEN_DEVICE_FUNC
1930   static inline bool isApprox(const Scalar& x, const Scalar& y, const RealScalar& prec)
1931   {
1932     return numext::abs2(x - y) <= numext::mini(numext::abs2(x), numext::abs2(y)) * prec * prec;
1933   }
1934 };
1935 
1936 template<typename Scalar>
1937 struct scalar_fuzzy_impl : scalar_fuzzy_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {};
1938 
1939 template<typename Scalar, typename OtherScalar> EIGEN_DEVICE_FUNC
1940 inline bool isMuchSmallerThan(const Scalar& x, const OtherScalar& y,
1941                               const typename NumTraits<Scalar>::Real &precision = NumTraits<Scalar>::dummy_precision())
1942 {
1943   return scalar_fuzzy_impl<Scalar>::template isMuchSmallerThan<OtherScalar>(x, y, precision);
1944 }
1945 
1946 template<typename Scalar> EIGEN_DEVICE_FUNC
1947 inline bool isApprox(const Scalar& x, const Scalar& y,
1948                      const typename NumTraits<Scalar>::Real &precision = NumTraits<Scalar>::dummy_precision())
1949 {
1950   return scalar_fuzzy_impl<Scalar>::isApprox(x, y, precision);
1951 }
1952 
1953 template<typename Scalar> EIGEN_DEVICE_FUNC
1954 inline bool isApproxOrLessThan(const Scalar& x, const Scalar& y,
1955                                const typename NumTraits<Scalar>::Real &precision = NumTraits<Scalar>::dummy_precision())
1956 {
1957   return scalar_fuzzy_impl<Scalar>::isApproxOrLessThan(x, y, precision);
1958 }
1959 
1960 /******************************************
1961 ***  The special case of the  bool type ***
1962 ******************************************/
1963 
1964 template<> struct random_impl<bool>
1965 {
1966   static inline bool run()
1967   {
1968     return random<int>(0,1)==0 ? false : true;
1969   }
1970 
1971   static inline bool run(const bool& a, const bool& b)
1972   {
1973     return random<int>(a, b)==0 ? false : true;
1974   }
1975 };
1976 
1977 template<> struct scalar_fuzzy_impl<bool>
1978 {
1979   typedef bool RealScalar;
1980 
1981   template<typename OtherScalar> EIGEN_DEVICE_FUNC
1982   static inline bool isMuchSmallerThan(const bool& x, const bool&, const bool&)
1983   {
1984     return !x;
1985   }
1986 
1987   EIGEN_DEVICE_FUNC
1988   static inline bool isApprox(bool x, bool y, bool)
1989   {
1990     return x == y;
1991   }
1992 
1993   EIGEN_DEVICE_FUNC
1994   static inline bool isApproxOrLessThan(const bool& x, const bool& y, const bool&)
1995   {
1996     return (!x) || y;
1997   }
1998 
1999 };
2000 
2001 } // end namespace internal
2002 
2003 // Default implementations that rely on other numext implementations
2004 namespace internal {
2005 
2006 // Specialization for complex types that are not supported by std::expm1.
2007 template <typename RealScalar>
2008 struct expm1_impl<std::complex<RealScalar> > {
2009   EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run(
2010       const std::complex<RealScalar>& x) {
2011     EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar)
2012     RealScalar xr = x.real();
2013     RealScalar xi = x.imag();
2014     // expm1(z) = exp(z) - 1
2015     //          = exp(x +  i * y) - 1
2016     //          = exp(x) * (cos(y) + i * sin(y)) - 1
2017     //          = exp(x) * cos(y) - 1 + i * exp(x) * sin(y)
2018     // Imag(expm1(z)) = exp(x) * sin(y)
2019     // Real(expm1(z)) = exp(x) * cos(y) - 1
2020     //          = exp(x) * cos(y) - 1.
2021     //          = expm1(x) + exp(x) * (cos(y) - 1)
2022     //          = expm1(x) + exp(x) * (2 * sin(y / 2) ** 2)
2023     RealScalar erm1 = numext::expm1<RealScalar>(xr);
2024     RealScalar er = erm1 + RealScalar(1.);
2025     RealScalar sin2 = numext::sin(xi / RealScalar(2.));
2026     sin2 = sin2 * sin2;
2027     RealScalar s = numext::sin(xi);
2028     RealScalar real_part = erm1 - RealScalar(2.) * er * sin2;
2029     return std::complex<RealScalar>(real_part, er * s);
2030   }
2031 };
2032 
2033 template<typename T>
2034 struct rsqrt_impl {
2035   EIGEN_DEVICE_FUNC
2036   static EIGEN_ALWAYS_INLINE T run(const T& x) {
2037     return T(1)/numext::sqrt(x);
2038   }
2039 };
2040 
2041 #if defined(EIGEN_GPU_COMPILE_PHASE)
2042 template<typename T>
2043 struct conj_impl<std::complex<T>, true>
2044 {
2045   EIGEN_DEVICE_FUNC
2046   static inline std::complex<T> run(const std::complex<T>& x)
2047   {
2048     return std::complex<T>(numext::real(x), -numext::imag(x));
2049   }
2050 };
2051 #endif
2052 
2053 } // end namespace internal
2054 
2055 } // end namespace Eigen
2056 
2057 #endif // EIGEN_MATHFUNCTIONS_H
[ Source navigation ] [ Diff markup ] [ Identifier search ] [ general search ]

This page was automatically generated by the 2.3.7 LXR engine.

The LXR team
Valid CSS 2.1! Valid HTML 4.01!