Branch data Line data Source code
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 : 21336 : static inline RealScalar run(const Scalar& x)
83 : : {
84 : 21336 : 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 : 54906 : static inline Scalar run(const Scalar& x)
248 : : {
249 : 54906 : 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 : 4526 : static inline RealScalar run(const Scalar& x)
283 : : {
284 : 4526 : 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 : 4526 : static inline RealScalar run(const Scalar& x)
305 : : {
306 : 4526 : 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 : 73227 : static inline Scalar run(const Scalar& x, const Scalar& y)
828 : : {
829 : 73227 : return x + (y-x) * Scalar(std::rand()) / Scalar(RAND_MAX);
830 : : }
831 : 64227 : static inline Scalar run()
832 : : {
833 : 64227 : 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 : 9000 : inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random(const Scalar& x, const Scalar& y)
946 : : {
947 : 9000 : return EIGEN_MATHFUNC_IMPL(random, Scalar)::run(x, y);
948 : : }
949 : :
950 : : template<typename Scalar>
951 : 64227 : inline EIGEN_MATHFUNC_RETVAL(random, Scalar) random()
952 : : {
953 : 64227 : 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 : 306 : 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 : 21336 : inline EIGEN_MATHFUNC_RETVAL(real, Scalar) real(const Scalar& x)
1231 : : {
1232 : 21336 : 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 : 54906 : inline EIGEN_MATHFUNC_RETVAL(conj, Scalar) conj(const Scalar& x)
1280 : : {
1281 : 54906 : return EIGEN_MATHFUNC_IMPL(conj, Scalar)::run(x);
1282 : : }
1283 : :
1284 : : template<typename Scalar>
1285 : : EIGEN_DEVICE_FUNC
1286 : 4526 : inline EIGEN_MATHFUNC_RETVAL(abs2, Scalar) abs2(const Scalar& x)
1287 : : {
1288 : 4526 : 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
|
