MathFunctions.h (60784B)
1 // This file is part of Eigen, a lightweight C++ template library 2 // for linear algebra. 3 // 4 // Copyright (C) 2006-2010 Benoit Jacob <jacob.benoit.1@gmail.com> 5 // Copyright (c) 2021, NVIDIA CORPORATION. All rights reserved. 6 // 7 // This Source Code Form is subject to the terms of the Mozilla 8 // Public License v. 2.0. If a copy of the MPL was not distributed 9 // with this file, You can obtain one at http://mozilla.org/MPL/2.0/. 10 11 #ifndef EIGEN_MATHFUNCTIONS_H 12 #define EIGEN_MATHFUNCTIONS_H 13 14 // TODO this should better be moved to NumTraits 15 // Source: WolframAlpha 16 #define EIGEN_PI 3.141592653589793238462643383279502884197169399375105820974944592307816406L 17 #define EIGEN_LOG2E 1.442695040888963407359924681001892137426645954152985934135449406931109219L 18 #define EIGEN_LN2 0.693147180559945309417232121458176568075500134360255254120680009493393621L 19 20 namespace Eigen { 21 22 // On WINCE, std::abs is defined for int only, so let's defined our own overloads: 23 // This issue has been confirmed with MSVC 2008 only, but the issue might exist for more recent versions too. 24 #if EIGEN_OS_WINCE && EIGEN_COMP_MSVC && EIGEN_COMP_MSVC<=1500 25 long abs(long x) { return (labs(x)); } 26 double abs(double x) { return (fabs(x)); } 27 float abs(float x) { return (fabsf(x)); } 28 long double abs(long double x) { return (fabsl(x)); } 29 #endif 30 31 namespace internal { 32 33 /** \internal \class global_math_functions_filtering_base 34 * 35 * What it does: 36 * Defines a typedef 'type' as follows: 37 * - if type T has a member typedef Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl, then 38 * global_math_functions_filtering_base<T>::type is a typedef for it. 39 * - otherwise, global_math_functions_filtering_base<T>::type is a typedef for T. 40 * 41 * How it's used: 42 * To allow to defined the global math functions (like sin...) in certain cases, like the Array expressions. 43 * When you do sin(array1+array2), the object array1+array2 has a complicated expression type, all what you want to know 44 * is that it inherits ArrayBase. So we implement a partial specialization of sin_impl for ArrayBase<Derived>. 45 * So we must make sure to use sin_impl<ArrayBase<Derived> > and not sin_impl<Derived>, otherwise our partial specialization 46 * won't be used. How does sin know that? That's exactly what global_math_functions_filtering_base tells it. 47 * 48 * How it's implemented: 49 * SFINAE in the style of enable_if. Highly susceptible of breaking compilers. With GCC, it sure does work, but if you replace 50 * the typename dummy by an integer template parameter, it doesn't work anymore! 51 */ 52 53 template<typename T, typename dummy = void> 54 struct global_math_functions_filtering_base 55 { 56 typedef T type; 57 }; 58 59 template<typename T> struct always_void { typedef void type; }; 60 61 template<typename T> 62 struct global_math_functions_filtering_base 63 <T, 64 typename always_void<typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl>::type 65 > 66 { 67 typedef typename T::Eigen_BaseClassForSpecializationOfGlobalMathFuncImpl type; 68 }; 69 70 #define EIGEN_MATHFUNC_IMPL(func, scalar) Eigen::internal::func##_impl<typename Eigen::internal::global_math_functions_filtering_base<scalar>::type> 71 #define EIGEN_MATHFUNC_RETVAL(func, scalar) typename Eigen::internal::func##_retval<typename Eigen::internal::global_math_functions_filtering_base<scalar>::type>::type 72 73 /**************************************************************************** 74 * Implementation of real * 75 ****************************************************************************/ 76 77 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex> 78 struct real_default_impl 79 { 80 typedef typename NumTraits<Scalar>::Real RealScalar; 81 EIGEN_DEVICE_FUNC 82 static inline RealScalar run(const Scalar& x) 83 { 84 return x; 85 } 86 }; 87 88 template<typename Scalar> 89 struct real_default_impl<Scalar,true> 90 { 91 typedef typename NumTraits<Scalar>::Real RealScalar; 92 EIGEN_DEVICE_FUNC 93 static inline RealScalar run(const Scalar& x) 94 { 95 using std::real; 96 return real(x); 97 } 98 }; 99 100 template<typename Scalar> struct real_impl : real_default_impl<Scalar> {}; 101 102 #if defined(EIGEN_GPU_COMPILE_PHASE) 103 template<typename T> 104 struct real_impl<std::complex<T> > 105 { 106 typedef T RealScalar; 107 EIGEN_DEVICE_FUNC 108 static inline T run(const std::complex<T>& x) 109 { 110 return x.real(); 111 } 112 }; 113 #endif 114 115 template<typename Scalar> 116 struct real_retval 117 { 118 typedef typename NumTraits<Scalar>::Real type; 119 }; 120 121 /**************************************************************************** 122 * Implementation of imag * 123 ****************************************************************************/ 124 125 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex> 126 struct imag_default_impl 127 { 128 typedef typename NumTraits<Scalar>::Real RealScalar; 129 EIGEN_DEVICE_FUNC 130 static inline RealScalar run(const Scalar&) 131 { 132 return RealScalar(0); 133 } 134 }; 135 136 template<typename Scalar> 137 struct imag_default_impl<Scalar,true> 138 { 139 typedef typename NumTraits<Scalar>::Real RealScalar; 140 EIGEN_DEVICE_FUNC 141 static inline RealScalar run(const Scalar& x) 142 { 143 using std::imag; 144 return imag(x); 145 } 146 }; 147 148 template<typename Scalar> struct imag_impl : imag_default_impl<Scalar> {}; 149 150 #if defined(EIGEN_GPU_COMPILE_PHASE) 151 template<typename T> 152 struct imag_impl<std::complex<T> > 153 { 154 typedef T RealScalar; 155 EIGEN_DEVICE_FUNC 156 static inline T run(const std::complex<T>& x) 157 { 158 return x.imag(); 159 } 160 }; 161 #endif 162 163 template<typename Scalar> 164 struct imag_retval 165 { 166 typedef typename NumTraits<Scalar>::Real type; 167 }; 168 169 /**************************************************************************** 170 * Implementation of real_ref * 171 ****************************************************************************/ 172 173 template<typename Scalar> 174 struct real_ref_impl 175 { 176 typedef typename NumTraits<Scalar>::Real RealScalar; 177 EIGEN_DEVICE_FUNC 178 static inline RealScalar& run(Scalar& x) 179 { 180 return reinterpret_cast<RealScalar*>(&x)[0]; 181 } 182 EIGEN_DEVICE_FUNC 183 static inline const RealScalar& run(const Scalar& x) 184 { 185 return reinterpret_cast<const RealScalar*>(&x)[0]; 186 } 187 }; 188 189 template<typename Scalar> 190 struct real_ref_retval 191 { 192 typedef typename NumTraits<Scalar>::Real & type; 193 }; 194 195 /**************************************************************************** 196 * Implementation of imag_ref * 197 ****************************************************************************/ 198 199 template<typename Scalar, bool IsComplex> 200 struct imag_ref_default_impl 201 { 202 typedef typename NumTraits<Scalar>::Real RealScalar; 203 EIGEN_DEVICE_FUNC 204 static inline RealScalar& run(Scalar& x) 205 { 206 return reinterpret_cast<RealScalar*>(&x)[1]; 207 } 208 EIGEN_DEVICE_FUNC 209 static inline const RealScalar& run(const Scalar& x) 210 { 211 return reinterpret_cast<RealScalar*>(&x)[1]; 212 } 213 }; 214 215 template<typename Scalar> 216 struct imag_ref_default_impl<Scalar, false> 217 { 218 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR 219 static inline Scalar run(Scalar&) 220 { 221 return Scalar(0); 222 } 223 EIGEN_DEVICE_FUNC EIGEN_CONSTEXPR 224 static inline const Scalar run(const Scalar&) 225 { 226 return Scalar(0); 227 } 228 }; 229 230 template<typename Scalar> 231 struct imag_ref_impl : imag_ref_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {}; 232 233 template<typename Scalar> 234 struct imag_ref_retval 235 { 236 typedef typename NumTraits<Scalar>::Real & type; 237 }; 238 239 /**************************************************************************** 240 * Implementation of conj * 241 ****************************************************************************/ 242 243 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex> 244 struct conj_default_impl 245 { 246 EIGEN_DEVICE_FUNC 247 static inline Scalar run(const Scalar& x) 248 { 249 return x; 250 } 251 }; 252 253 template<typename Scalar> 254 struct conj_default_impl<Scalar,true> 255 { 256 EIGEN_DEVICE_FUNC 257 static inline Scalar run(const Scalar& x) 258 { 259 using std::conj; 260 return conj(x); 261 } 262 }; 263 264 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex> 265 struct conj_impl : conj_default_impl<Scalar, IsComplex> {}; 266 267 template<typename Scalar> 268 struct conj_retval 269 { 270 typedef Scalar type; 271 }; 272 273 /**************************************************************************** 274 * Implementation of abs2 * 275 ****************************************************************************/ 276 277 template<typename Scalar,bool IsComplex> 278 struct abs2_impl_default 279 { 280 typedef typename NumTraits<Scalar>::Real RealScalar; 281 EIGEN_DEVICE_FUNC 282 static inline RealScalar run(const Scalar& x) 283 { 284 return x*x; 285 } 286 }; 287 288 template<typename Scalar> 289 struct abs2_impl_default<Scalar, true> // IsComplex 290 { 291 typedef typename NumTraits<Scalar>::Real RealScalar; 292 EIGEN_DEVICE_FUNC 293 static inline RealScalar run(const Scalar& x) 294 { 295 return x.real()*x.real() + x.imag()*x.imag(); 296 } 297 }; 298 299 template<typename Scalar> 300 struct abs2_impl 301 { 302 typedef typename NumTraits<Scalar>::Real RealScalar; 303 EIGEN_DEVICE_FUNC 304 static inline RealScalar run(const Scalar& x) 305 { 306 return abs2_impl_default<Scalar,NumTraits<Scalar>::IsComplex>::run(x); 307 } 308 }; 309 310 template<typename Scalar> 311 struct abs2_retval 312 { 313 typedef typename NumTraits<Scalar>::Real type; 314 }; 315 316 /**************************************************************************** 317 * Implementation of sqrt/rsqrt * 318 ****************************************************************************/ 319 320 template<typename Scalar> 321 struct sqrt_impl 322 { 323 EIGEN_DEVICE_FUNC 324 static EIGEN_ALWAYS_INLINE Scalar run(const Scalar& x) 325 { 326 EIGEN_USING_STD(sqrt); 327 return sqrt(x); 328 } 329 }; 330 331 // Complex sqrt defined in MathFunctionsImpl.h. 332 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_sqrt(const std::complex<T>& a_x); 333 334 // Custom implementation is faster than `std::sqrt`, works on 335 // GPU, and correctly handles special cases (unlike MSVC). 336 template<typename T> 337 struct sqrt_impl<std::complex<T> > 338 { 339 EIGEN_DEVICE_FUNC 340 static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x) 341 { 342 return complex_sqrt<T>(x); 343 } 344 }; 345 346 template<typename Scalar> 347 struct sqrt_retval 348 { 349 typedef Scalar type; 350 }; 351 352 // Default implementation relies on numext::sqrt, at bottom of file. 353 template<typename T> 354 struct rsqrt_impl; 355 356 // Complex rsqrt defined in MathFunctionsImpl.h. 357 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_rsqrt(const std::complex<T>& a_x); 358 359 template<typename T> 360 struct rsqrt_impl<std::complex<T> > 361 { 362 EIGEN_DEVICE_FUNC 363 static EIGEN_ALWAYS_INLINE std::complex<T> run(const std::complex<T>& x) 364 { 365 return complex_rsqrt<T>(x); 366 } 367 }; 368 369 template<typename Scalar> 370 struct rsqrt_retval 371 { 372 typedef Scalar type; 373 }; 374 375 /**************************************************************************** 376 * Implementation of norm1 * 377 ****************************************************************************/ 378 379 template<typename Scalar, bool IsComplex> 380 struct norm1_default_impl; 381 382 template<typename Scalar> 383 struct norm1_default_impl<Scalar,true> 384 { 385 typedef typename NumTraits<Scalar>::Real RealScalar; 386 EIGEN_DEVICE_FUNC 387 static inline RealScalar run(const Scalar& x) 388 { 389 EIGEN_USING_STD(abs); 390 return abs(x.real()) + abs(x.imag()); 391 } 392 }; 393 394 template<typename Scalar> 395 struct norm1_default_impl<Scalar, false> 396 { 397 EIGEN_DEVICE_FUNC 398 static inline Scalar run(const Scalar& x) 399 { 400 EIGEN_USING_STD(abs); 401 return abs(x); 402 } 403 }; 404 405 template<typename Scalar> 406 struct norm1_impl : norm1_default_impl<Scalar, NumTraits<Scalar>::IsComplex> {}; 407 408 template<typename Scalar> 409 struct norm1_retval 410 { 411 typedef typename NumTraits<Scalar>::Real type; 412 }; 413 414 /**************************************************************************** 415 * Implementation of hypot * 416 ****************************************************************************/ 417 418 template<typename Scalar> struct hypot_impl; 419 420 template<typename Scalar> 421 struct hypot_retval 422 { 423 typedef typename NumTraits<Scalar>::Real type; 424 }; 425 426 /**************************************************************************** 427 * Implementation of cast * 428 ****************************************************************************/ 429 430 template<typename OldType, typename NewType, typename EnableIf = void> 431 struct cast_impl 432 { 433 EIGEN_DEVICE_FUNC 434 static inline NewType run(const OldType& x) 435 { 436 return static_cast<NewType>(x); 437 } 438 }; 439 440 // Casting from S -> Complex<T> leads to an implicit conversion from S to T, 441 // generating warnings on clang. Here we explicitly cast the real component. 442 template<typename OldType, typename NewType> 443 struct cast_impl<OldType, NewType, 444 typename internal::enable_if< 445 !NumTraits<OldType>::IsComplex && NumTraits<NewType>::IsComplex 446 >::type> 447 { 448 EIGEN_DEVICE_FUNC 449 static inline NewType run(const OldType& x) 450 { 451 typedef typename NumTraits<NewType>::Real NewReal; 452 return static_cast<NewType>(static_cast<NewReal>(x)); 453 } 454 }; 455 456 // here, for once, we're plainly returning NewType: we don't want cast to do weird things. 457 458 template<typename OldType, typename NewType> 459 EIGEN_DEVICE_FUNC 460 inline NewType cast(const OldType& x) 461 { 462 return cast_impl<OldType, NewType>::run(x); 463 } 464 465 /**************************************************************************** 466 * Implementation of round * 467 ****************************************************************************/ 468 469 template<typename Scalar> 470 struct round_impl 471 { 472 EIGEN_DEVICE_FUNC 473 static inline Scalar run(const Scalar& x) 474 { 475 EIGEN_STATIC_ASSERT((!NumTraits<Scalar>::IsComplex), NUMERIC_TYPE_MUST_BE_REAL) 476 #if EIGEN_HAS_CXX11_MATH 477 EIGEN_USING_STD(round); 478 #endif 479 return Scalar(round(x)); 480 } 481 }; 482 483 #if !EIGEN_HAS_CXX11_MATH 484 #if EIGEN_HAS_C99_MATH 485 // Use ::roundf for float. 486 template<> 487 struct round_impl<float> { 488 EIGEN_DEVICE_FUNC 489 static inline float run(const float& x) 490 { 491 return ::roundf(x); 492 } 493 }; 494 #else 495 template<typename Scalar> 496 struct round_using_floor_ceil_impl 497 { 498 EIGEN_DEVICE_FUNC 499 static inline Scalar run(const Scalar& x) 500 { 501 EIGEN_STATIC_ASSERT((!NumTraits<Scalar>::IsComplex), NUMERIC_TYPE_MUST_BE_REAL) 502 // Without C99 round/roundf, resort to floor/ceil. 503 EIGEN_USING_STD(floor); 504 EIGEN_USING_STD(ceil); 505 // If not enough precision to resolve a decimal at all, return the input. 506 // Otherwise, adding 0.5 can trigger an increment by 1. 507 const Scalar limit = Scalar(1ull << (NumTraits<Scalar>::digits() - 1)); 508 if (x >= limit || x <= -limit) { 509 return x; 510 } 511 return (x > Scalar(0)) ? Scalar(floor(x + Scalar(0.5))) : Scalar(ceil(x - Scalar(0.5))); 512 } 513 }; 514 515 template<> 516 struct round_impl<float> : round_using_floor_ceil_impl<float> {}; 517 518 template<> 519 struct round_impl<double> : round_using_floor_ceil_impl<double> {}; 520 #endif // EIGEN_HAS_C99_MATH 521 #endif // !EIGEN_HAS_CXX11_MATH 522 523 template<typename Scalar> 524 struct round_retval 525 { 526 typedef Scalar type; 527 }; 528 529 /**************************************************************************** 530 * Implementation of rint * 531 ****************************************************************************/ 532 533 template<typename Scalar> 534 struct rint_impl { 535 EIGEN_DEVICE_FUNC 536 static inline Scalar run(const Scalar& x) 537 { 538 EIGEN_STATIC_ASSERT((!NumTraits<Scalar>::IsComplex), NUMERIC_TYPE_MUST_BE_REAL) 539 #if EIGEN_HAS_CXX11_MATH 540 EIGEN_USING_STD(rint); 541 #endif 542 return rint(x); 543 } 544 }; 545 546 #if !EIGEN_HAS_CXX11_MATH 547 template<> 548 struct rint_impl<double> { 549 EIGEN_DEVICE_FUNC 550 static inline double run(const double& x) 551 { 552 return ::rint(x); 553 } 554 }; 555 template<> 556 struct rint_impl<float> { 557 EIGEN_DEVICE_FUNC 558 static inline float run(const float& x) 559 { 560 return ::rintf(x); 561 } 562 }; 563 #endif 564 565 template<typename Scalar> 566 struct rint_retval 567 { 568 typedef Scalar type; 569 }; 570 571 /**************************************************************************** 572 * Implementation of arg * 573 ****************************************************************************/ 574 575 // Visual Studio 2017 has a bug where arg(float) returns 0 for negative inputs. 576 // This seems to be fixed in VS 2019. 577 #if EIGEN_HAS_CXX11_MATH && (!EIGEN_COMP_MSVC || EIGEN_COMP_MSVC >= 1920) 578 // std::arg is only defined for types of std::complex, or integer types or float/double/long double 579 template<typename Scalar, 580 bool HasStdImpl = NumTraits<Scalar>::IsComplex || is_integral<Scalar>::value 581 || is_same<Scalar, float>::value || is_same<Scalar, double>::value 582 || is_same<Scalar, long double>::value > 583 struct arg_default_impl; 584 585 template<typename Scalar> 586 struct arg_default_impl<Scalar, true> { 587 typedef typename NumTraits<Scalar>::Real RealScalar; 588 EIGEN_DEVICE_FUNC 589 static inline RealScalar run(const Scalar& x) 590 { 591 #if defined(EIGEN_HIP_DEVICE_COMPILE) 592 // HIP does not seem to have a native device side implementation for the math routine "arg" 593 using std::arg; 594 #else 595 EIGEN_USING_STD(arg); 596 #endif 597 return static_cast<RealScalar>(arg(x)); 598 } 599 }; 600 601 // Must be non-complex floating-point type (e.g. half/bfloat16). 602 template<typename Scalar> 603 struct arg_default_impl<Scalar, false> { 604 typedef typename NumTraits<Scalar>::Real RealScalar; 605 EIGEN_DEVICE_FUNC 606 static inline RealScalar run(const Scalar& x) 607 { 608 return (x < Scalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0); 609 } 610 }; 611 #else 612 template<typename Scalar, bool IsComplex = NumTraits<Scalar>::IsComplex> 613 struct arg_default_impl 614 { 615 typedef typename NumTraits<Scalar>::Real RealScalar; 616 EIGEN_DEVICE_FUNC 617 static inline RealScalar run(const Scalar& x) 618 { 619 return (x < RealScalar(0)) ? RealScalar(EIGEN_PI) : RealScalar(0); 620 } 621 }; 622 623 template<typename Scalar> 624 struct arg_default_impl<Scalar,true> 625 { 626 typedef typename NumTraits<Scalar>::Real RealScalar; 627 EIGEN_DEVICE_FUNC 628 static inline RealScalar run(const Scalar& x) 629 { 630 EIGEN_USING_STD(arg); 631 return arg(x); 632 } 633 }; 634 #endif 635 template<typename Scalar> struct arg_impl : arg_default_impl<Scalar> {}; 636 637 template<typename Scalar> 638 struct arg_retval 639 { 640 typedef typename NumTraits<Scalar>::Real type; 641 }; 642 643 /**************************************************************************** 644 * Implementation of expm1 * 645 ****************************************************************************/ 646 647 // This implementation is based on GSL Math's expm1. 648 namespace std_fallback { 649 // fallback expm1 implementation in case there is no expm1(Scalar) function in namespace of Scalar, 650 // or that there is no suitable std::expm1 function available. Implementation 651 // attributed to Kahan. See: http://www.plunk.org/~hatch/rightway.php. 652 template<typename Scalar> 653 EIGEN_DEVICE_FUNC inline Scalar expm1(const Scalar& x) { 654 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) 655 typedef typename NumTraits<Scalar>::Real RealScalar; 656 657 EIGEN_USING_STD(exp); 658 Scalar u = exp(x); 659 if (numext::equal_strict(u, Scalar(1))) { 660 return x; 661 } 662 Scalar um1 = u - RealScalar(1); 663 if (numext::equal_strict(um1, Scalar(-1))) { 664 return RealScalar(-1); 665 } 666 667 EIGEN_USING_STD(log); 668 Scalar logu = log(u); 669 return numext::equal_strict(u, logu) ? u : (u - RealScalar(1)) * x / logu; 670 } 671 } 672 673 template<typename Scalar> 674 struct expm1_impl { 675 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) 676 { 677 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) 678 #if EIGEN_HAS_CXX11_MATH 679 using std::expm1; 680 #else 681 using std_fallback::expm1; 682 #endif 683 return expm1(x); 684 } 685 }; 686 687 template<typename Scalar> 688 struct expm1_retval 689 { 690 typedef Scalar type; 691 }; 692 693 /**************************************************************************** 694 * Implementation of log * 695 ****************************************************************************/ 696 697 // Complex log defined in MathFunctionsImpl.h. 698 template<typename T> EIGEN_DEVICE_FUNC std::complex<T> complex_log(const std::complex<T>& z); 699 700 template<typename Scalar> 701 struct log_impl { 702 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) 703 { 704 EIGEN_USING_STD(log); 705 return static_cast<Scalar>(log(x)); 706 } 707 }; 708 709 template<typename Scalar> 710 struct log_impl<std::complex<Scalar> > { 711 EIGEN_DEVICE_FUNC static inline std::complex<Scalar> run(const std::complex<Scalar>& z) 712 { 713 return complex_log(z); 714 } 715 }; 716 717 /**************************************************************************** 718 * Implementation of log1p * 719 ****************************************************************************/ 720 721 namespace std_fallback { 722 // fallback log1p implementation in case there is no log1p(Scalar) function in namespace of Scalar, 723 // or that there is no suitable std::log1p function available 724 template<typename Scalar> 725 EIGEN_DEVICE_FUNC inline Scalar log1p(const Scalar& x) { 726 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) 727 typedef typename NumTraits<Scalar>::Real RealScalar; 728 EIGEN_USING_STD(log); 729 Scalar x1p = RealScalar(1) + x; 730 Scalar log_1p = log_impl<Scalar>::run(x1p); 731 const bool is_small = numext::equal_strict(x1p, Scalar(1)); 732 const bool is_inf = numext::equal_strict(x1p, log_1p); 733 return (is_small || is_inf) ? x : x * (log_1p / (x1p - RealScalar(1))); 734 } 735 } 736 737 template<typename Scalar> 738 struct log1p_impl { 739 EIGEN_DEVICE_FUNC static inline Scalar run(const Scalar& x) 740 { 741 EIGEN_STATIC_ASSERT_NON_INTEGER(Scalar) 742 #if EIGEN_HAS_CXX11_MATH 743 using std::log1p; 744 #else 745 using std_fallback::log1p; 746 #endif 747 return log1p(x); 748 } 749 }; 750 751 // Specialization for complex types that are not supported by std::log1p. 752 template <typename RealScalar> 753 struct log1p_impl<std::complex<RealScalar> > { 754 EIGEN_DEVICE_FUNC static inline std::complex<RealScalar> run( 755 const std::complex<RealScalar>& x) { 756 EIGEN_STATIC_ASSERT_NON_INTEGER(RealScalar) 757 return std_fallback::log1p(x); 758 } 759 }; 760 761 template<typename Scalar> 762 struct log1p_retval 763 { 764 typedef Scalar type; 765 }; 766 767 /**************************************************************************** 768 * Implementation of pow * 769 ****************************************************************************/ 770 771 template<typename ScalarX,typename ScalarY, bool IsInteger = NumTraits<ScalarX>::IsInteger&&NumTraits<ScalarY>::IsInteger> 772 struct pow_impl 773 { 774 //typedef Scalar retval; 775 typedef typename ScalarBinaryOpTraits<ScalarX,ScalarY,internal::scalar_pow_op<ScalarX,ScalarY> >::ReturnType result_type; 776 static EIGEN_DEVICE_FUNC inline result_type run(const ScalarX& x, const ScalarY& y) 777 { 778 EIGEN_USING_STD(pow); 779 return pow(x, y); 780 } 781 }; 782 783 template<typename ScalarX,typename ScalarY> 784 struct pow_impl<ScalarX,ScalarY, true> 785 { 786 typedef ScalarX result_type; 787 static EIGEN_DEVICE_FUNC inline ScalarX run(ScalarX x, ScalarY y) 788 { 789 ScalarX res(1); 790 eigen_assert(!NumTraits<ScalarY>::IsSigned || y >= 0); 791 if(y & 1) res *= x; 792 y >>= 1; 793 while(y) 794 { 795 x *= x; 796 if(y&1) res *= x; 797 y >>= 1; 798 } 799 return res; 800 } 801 }; 802 803 /**************************************************************************** 804 * Implementation of random * 805 ****************************************************************************/ 806 807 template<typename Scalar, 808 bool IsComplex, 809 bool IsInteger> 810 struct random_default_impl {}; 811 812 template<typename Scalar> 813 struct random_impl : random_default_impl<Scalar, NumTraits<Scalar>::IsComplex, NumTraits<Scalar>::IsInteger> {}; 814 815 template<typename Scalar> 816 struct random_retval 817 { 818 typedef Scalar type; 819 }; 820 821 template<typename Scalar> inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(const Scalar& x, const Scalar& y); 822 template<typename Scalar> inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(); 823 824 template<typename Scalar> 825 struct random_default_impl<Scalar, false, false> 826 { 827 static inline Scalar run(const Scalar& x, const Scalar& y) 828 { 829 return x + (y-x) * Scalar(std::rand()) / Scalar(RAND_MAX); 830 } 831 static inline Scalar run() 832 { 833 return run(Scalar(NumTraits<Scalar>::IsSigned ? -1 : 0), Scalar(1)); 834 } 835 }; 836 837 enum { 838 meta_floor_log2_terminate, 839 meta_floor_log2_move_up, 840 meta_floor_log2_move_down, 841 meta_floor_log2_bogus 842 }; 843 844 template<unsigned int n, int lower, int upper> struct meta_floor_log2_selector 845 { 846 enum { middle = (lower + upper) / 2, 847 value = (upper <= lower + 1) ? int(meta_floor_log2_terminate) 848 : (n < (1 << middle)) ? int(meta_floor_log2_move_down) 849 : (n==0) ? int(meta_floor_log2_bogus) 850 : int(meta_floor_log2_move_up) 851 }; 852 }; 853 854 template<unsigned int n, 855 int lower = 0, 856 int upper = sizeof(unsigned int) * CHAR_BIT - 1, 857 int selector = meta_floor_log2_selector<n, lower, upper>::value> 858 struct meta_floor_log2 {}; 859 860 template<unsigned int n, int lower, int upper> 861 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_down> 862 { 863 enum { value = meta_floor_log2<n, lower, meta_floor_log2_selector<n, lower, upper>::middle>::value }; 864 }; 865 866 template<unsigned int n, int lower, int upper> 867 struct meta_floor_log2<n, lower, upper, meta_floor_log2_move_up> 868 { 869 enum { value = meta_floor_log2<n, meta_floor_log2_selector<n, lower, upper>::middle, upper>::value }; 870 }; 871 872 template<unsigned int n, int lower, int upper> 873 struct meta_floor_log2<n, lower, upper, meta_floor_log2_terminate> 874 { 875 enum { value = (n >= ((unsigned int)(1) << (lower+1))) ? lower+1 : lower }; 876 }; 877 878 template<unsigned int n, int lower, int upper> 879 struct meta_floor_log2<n, lower, upper, meta_floor_log2_bogus> 880 { 881 // no value, error at compile time 882 }; 883 884 template<typename Scalar> 885 struct random_default_impl<Scalar, false, true> 886 { 887 static inline Scalar run(const Scalar& x, const Scalar& y) 888 { 889 if (y <= x) 890 return x; 891 // ScalarU is the unsigned counterpart of Scalar, possibly Scalar itself. 892 typedef typename make_unsigned<Scalar>::type ScalarU; 893 // ScalarX is the widest of ScalarU and unsigned int. 894 // We'll deal only with ScalarX and unsigned int below thus avoiding signed 895 // types and arithmetic and signed overflows (which are undefined behavior). 896 typedef typename conditional<(ScalarU(-1) > unsigned(-1)), ScalarU, unsigned>::type ScalarX; 897 // The following difference doesn't overflow, provided our integer types are two's 898 // complement and have the same number of padding bits in signed and unsigned variants. 899 // This is the case in most modern implementations of C++. 900 ScalarX range = ScalarX(y) - ScalarX(x); 901 ScalarX offset = 0; 902 ScalarX divisor = 1; 903 ScalarX multiplier = 1; 904 const unsigned rand_max = RAND_MAX; 905 if (range <= rand_max) divisor = (rand_max + 1) / (range + 1); 906 else multiplier = 1 + range / (rand_max + 1); 907 // Rejection sampling. 908 do { 909 offset = (unsigned(std::rand()) * multiplier) / divisor; 910 } while (offset > range); 911 return Scalar(ScalarX(x) + offset); 912 } 913 914 static inline Scalar run() 915 { 916 #ifdef EIGEN_MAKING_DOCS 917 return run(Scalar(NumTraits<Scalar>::IsSigned ? -10 : 0), Scalar(10)); 918 #else 919 enum { rand_bits = meta_floor_log2<(unsigned int)(RAND_MAX)+1>::value, 920 scalar_bits = sizeof(Scalar) * CHAR_BIT, 921 shift = EIGEN_PLAIN_ENUM_MAX(0, int(rand_bits) - int(scalar_bits)), 922 offset = NumTraits<Scalar>::IsSigned ? (1 << (EIGEN_PLAIN_ENUM_MIN(rand_bits,scalar_bits)-1)) : 0 923 }; 924 return Scalar((std::rand() >> shift) - offset); 925 #endif 926 } 927 }; 928 929 template<typename Scalar> 930 struct random_default_impl<Scalar, true, false> 931 { 932 static inline Scalar run(const Scalar& x, const Scalar& y) 933 { 934 return Scalar(random(x.real(), y.real()), 935 random(x.imag(), y.imag())); 936 } 937 static inline Scalar run() 938 { 939 typedef typename NumTraits<Scalar>::Real RealScalar; 940 return Scalar(random<RealScalar>(), random<RealScalar>()); 941 } 942 }; 943 944 template<typename Scalar> 945 inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(const Scalar& x, const Scalar& y) 946 { 947 return EIGEN_MATHFUNC_IMPL(random, Scalar)::run(x, y); 948 } 949 950 template<typename Scalar> 951 inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random() 952 { 953 return EIGEN_MATHFUNC_IMPL(random, Scalar)::run(); 954 } 955 956 // Implementation of is* functions 957 958 // std::is* do not work with fast-math and gcc, std::is* are available on MSVC 2013 and newer, as well as in clang. 959 #if (EIGEN_HAS_CXX11_MATH && !(EIGEN_COMP_GNUC_STRICT && __FINITE_MATH_ONLY__)) || (EIGEN_COMP_MSVC>=1800) || (EIGEN_COMP_CLANG) 960 #define EIGEN_USE_STD_FPCLASSIFY 1 961 #else 962 #define EIGEN_USE_STD_FPCLASSIFY 0 963 #endif 964 965 template<typename T> 966 EIGEN_DEVICE_FUNC 967 typename internal::enable_if<internal::is_integral<T>::value,bool>::type 968 isnan_impl(const T&) { return false; } 969 970 template<typename T> 971 EIGEN_DEVICE_FUNC 972 typename internal::enable_if<internal::is_integral<T>::value,bool>::type 973 isinf_impl(const T&) { return false; } 974 975 template<typename T> 976 EIGEN_DEVICE_FUNC 977 typename internal::enable_if<internal::is_integral<T>::value,bool>::type 978 isfinite_impl(const T&) { return true; } 979 980 template<typename T> 981 EIGEN_DEVICE_FUNC 982 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type 983 isfinite_impl(const T& x) 984 { 985 #if defined(EIGEN_GPU_COMPILE_PHASE) 986 return (::isfinite)(x); 987 #elif EIGEN_USE_STD_FPCLASSIFY 988 using std::isfinite; 989 return isfinite EIGEN_NOT_A_MACRO (x); 990 #else 991 return x<=NumTraits<T>::highest() && x>=NumTraits<T>::lowest(); 992 #endif 993 } 994 995 template<typename T> 996 EIGEN_DEVICE_FUNC 997 typename internal::enable_if<(!internal::is_integral<T>::value)&&(!NumTraits<T>::IsComplex),bool>::type 998 isinf_impl(const T& x) 999 { 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