Kokkos Core Kernels Package  Version of the Day
 All Classes Namespaces Files Functions Variables Typedefs Enumerations Enumerator Friends Groups Pages
Kokkos_MathematicalFunctions.hpp
1 //@HEADER
2 // ************************************************************************
3 //
4 // Kokkos v. 4.0
5 // Copyright (2022) National Technology & Engineering
6 // Solutions of Sandia, LLC (NTESS).
7 //
8 // Under the terms of Contract DE-NA0003525 with NTESS,
9 // the U.S. Government retains certain rights in this software.
10 //
11 // Part of Kokkos, under the Apache License v2.0 with LLVM Exceptions.
12 // See https://kokkos.org/LICENSE for license information.
13 // SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
14 //
15 //@HEADER
16 
17 #ifndef KOKKOS_MATHEMATICAL_FUNCTIONS_HPP
18 #define KOKKOS_MATHEMATICAL_FUNCTIONS_HPP
19 #ifndef KOKKOS_IMPL_PUBLIC_INCLUDE
20 #define KOKKOS_IMPL_PUBLIC_INCLUDE
21 #define KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_MATHFUNCTIONS
22 #endif
23 
24 #include <Kokkos_Macros.hpp>
25 #include <cmath>
26 #include <cstdlib>
27 #include <type_traits>
28 
29 #ifdef KOKKOS_ENABLE_SYCL
30 // FIXME_SYCL
31 #if __has_include(<sycl/sycl.hpp>)
32 #include <sycl/sycl.hpp>
33 #else
34 #include <CL/sycl.hpp>
35 #endif
36 #endif
37 
38 namespace Kokkos {
39 
40 namespace Impl {
41 template <class T, bool = std::is_integral_v<T>>
42 struct promote {
43  using type = double;
44 };
45 template <class T>
46 struct promote<T, false> {};
47 template <>
48 struct promote<long double> {
49  using type = long double;
50 };
51 template <>
52 struct promote<double> {
53  using type = double;
54 };
55 template <>
56 struct promote<float> {
57  using type = float;
58 };
59 template <class T>
60 using promote_t = typename promote<T>::type;
61 template <class T, class U,
62  bool = std::is_arithmetic_v<T>&& std::is_arithmetic_v<U>>
63 struct promote_2 {
64  using type = decltype(promote_t<T>() + promote_t<U>());
65 };
66 template <class T, class U>
67 struct promote_2<T, U, false> {};
68 template <class T, class U>
69 using promote_2_t = typename promote_2<T, U>::type;
70 template <class T, class U, class V,
71  bool = std::is_arithmetic_v<T>&& std::is_arithmetic_v<U>&&
72  std::is_arithmetic_v<V>>
73 struct promote_3 {
74  using type = decltype(promote_t<T>() + promote_t<U>() + promote_t<V>());
75 };
76 template <class T, class U, class V>
77 struct promote_3<T, U, V, false> {};
78 template <class T, class U, class V>
79 using promote_3_t = typename promote_3<T, U, V>::type;
80 } // namespace Impl
81 
82 // NOTE long double overloads are not available on the device
83 
84 #if defined(KOKKOS_ENABLE_SYCL)
85 #define KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE sycl
86 #else
87 #if (defined(KOKKOS_COMPILER_NVCC) || defined(KOKKOS_COMPILER_NVHPC)) && \
88  defined(__GNUC__) && (__GNUC__ < 6) && !defined(__clang__)
89 #define KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE
90 #else
91 #define KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE std
92 #endif
93 #endif
94 
95 #define KOKKOS_IMPL_MATH_UNARY_FUNCTION(FUNC) \
96  KOKKOS_INLINE_FUNCTION float FUNC(float x) { \
97  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
98  return FUNC(x); \
99  } \
100  KOKKOS_INLINE_FUNCTION double FUNC(double x) { \
101  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
102  return FUNC(x); \
103  } \
104  inline long double FUNC(long double x) { \
105  using std::FUNC; \
106  return FUNC(x); \
107  } \
108  KOKKOS_INLINE_FUNCTION float FUNC##f(float x) { \
109  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
110  return FUNC(x); \
111  } \
112  inline long double FUNC##l(long double x) { \
113  using std::FUNC; \
114  return FUNC(x); \
115  } \
116  template <class T> \
117  KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, double> FUNC( \
118  T x) { \
119  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
120  return FUNC(static_cast<double>(x)); \
121  }
122 
123 // isinf, isnan, and isinfinite do not work on Windows with CUDA with std::
124 // getting warnings about calling host function in device function then
125 // runtime test fails
126 #if defined(_WIN32) && defined(KOKKOS_ENABLE_CUDA)
127 #define KOKKOS_IMPL_MATH_UNARY_PREDICATE(FUNC) \
128  KOKKOS_INLINE_FUNCTION bool FUNC(float x) { return ::FUNC(x); } \
129  KOKKOS_INLINE_FUNCTION bool FUNC(double x) { return ::FUNC(x); } \
130  inline bool FUNC(long double x) { \
131  using std::FUNC; \
132  return FUNC(x); \
133  } \
134  template <class T> \
135  KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, bool> FUNC( \
136  T x) { \
137  return ::FUNC(static_cast<double>(x)); \
138  }
139 #else
140 #define KOKKOS_IMPL_MATH_UNARY_PREDICATE(FUNC) \
141  KOKKOS_INLINE_FUNCTION bool FUNC(float x) { \
142  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
143  return FUNC(x); \
144  } \
145  KOKKOS_INLINE_FUNCTION bool FUNC(double x) { \
146  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
147  return FUNC(x); \
148  } \
149  inline bool FUNC(long double x) { \
150  using std::FUNC; \
151  return FUNC(x); \
152  } \
153  template <class T> \
154  KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, bool> FUNC( \
155  T x) { \
156  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
157  return FUNC(static_cast<double>(x)); \
158  }
159 #endif
160 
161 #define KOKKOS_IMPL_MATH_BINARY_FUNCTION(FUNC) \
162  KOKKOS_INLINE_FUNCTION float FUNC(float x, float y) { \
163  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
164  return FUNC(x, y); \
165  } \
166  KOKKOS_INLINE_FUNCTION double FUNC(double x, double y) { \
167  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
168  return FUNC(x, y); \
169  } \
170  inline long double FUNC(long double x, long double y) { \
171  using std::FUNC; \
172  return FUNC(x, y); \
173  } \
174  KOKKOS_INLINE_FUNCTION float FUNC##f(float x, float y) { \
175  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
176  return FUNC(x, y); \
177  } \
178  inline long double FUNC##l(long double x, long double y) { \
179  using std::FUNC; \
180  return FUNC(x, y); \
181  } \
182  template <class T1, class T2> \
183  KOKKOS_INLINE_FUNCTION \
184  std::enable_if_t<std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> && \
185  !std::is_same_v<T1, long double> && \
186  !std::is_same_v<T2, long double>, \
187  Kokkos::Impl::promote_2_t<T1, T2>> \
188  FUNC(T1 x, T2 y) { \
189  using Promoted = Kokkos::Impl::promote_2_t<T1, T2>; \
190  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
191  return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y)); \
192  } \
193  template <class T1, class T2> \
194  inline std::enable_if_t<std::is_arithmetic_v<T1> && \
195  std::is_arithmetic_v<T2> && \
196  (std::is_same_v<T1, long double> || \
197  std::is_same_v<T2, long double>), \
198  long double> \
199  FUNC(T1 x, T2 y) { \
200  using Promoted = Kokkos::Impl::promote_2_t<T1, T2>; \
201  static_assert(std::is_same_v<Promoted, long double>); \
202  using std::FUNC; \
203  return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y)); \
204  }
205 
206 #define KOKKOS_IMPL_MATH_TERNARY_FUNCTION(FUNC) \
207  KOKKOS_INLINE_FUNCTION float FUNC(float x, float y, float z) { \
208  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
209  return FUNC(x, y, z); \
210  } \
211  KOKKOS_INLINE_FUNCTION double FUNC(double x, double y, double z) { \
212  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
213  return FUNC(x, y, z); \
214  } \
215  inline long double FUNC(long double x, long double y, long double z) { \
216  using std::FUNC; \
217  return FUNC(x, y, z); \
218  } \
219  KOKKOS_INLINE_FUNCTION float FUNC##f(float x, float y, float z) { \
220  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
221  return FUNC(x, y, z); \
222  } \
223  inline long double FUNC##l(long double x, long double y, long double z) { \
224  using std::FUNC; \
225  return FUNC(x, y, z); \
226  } \
227  template <class T1, class T2, class T3> \
228  KOKKOS_INLINE_FUNCTION std::enable_if_t< \
229  std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> && \
230  std::is_arithmetic_v<T3> && !std::is_same_v<T1, long double> && \
231  !std::is_same_v<T2, long double> && \
232  !std::is_same_v<T3, long double>, \
233  Kokkos::Impl::promote_3_t<T1, T2, T3>> \
234  FUNC(T1 x, T2 y, T3 z) { \
235  using Promoted = Kokkos::Impl::promote_3_t<T1, T2, T3>; \
236  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
237  return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y), \
238  static_cast<Promoted>(z)); \
239  } \
240  template <class T1, class T2, class T3> \
241  inline std::enable_if_t<std::is_arithmetic_v<T1> && \
242  std::is_arithmetic_v<T2> && \
243  std::is_arithmetic_v<T3> && \
244  (std::is_same_v<T1, long double> || \
245  std::is_same_v<T2, long double> || \
246  std::is_same_v<T3, long double>), \
247  long double> \
248  FUNC(T1 x, T2 y, T3 z) { \
249  using Promoted = Kokkos::Impl::promote_3_t<T1, T2, T3>; \
250  static_assert(std::is_same_v<Promoted, long double>); \
251  using std::FUNC; \
252  return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y), \
253  static_cast<Promoted>(z)); \
254  }
255 
256 // Basic operations
257 KOKKOS_INLINE_FUNCTION int abs(int n) {
258  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
259  return abs(n);
260 }
261 KOKKOS_INLINE_FUNCTION long abs(long n) {
262 // FIXME_NVHPC ptxas fatal : unresolved extern function 'labs'
263 #if defined(KOKKOS_COMPILER_NVHPC) && KOKKOS_COMPILER_NVHPC < 230700
264  return n > 0 ? n : -n;
265 #else
266  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
267  return abs(n);
268 #endif
269 }
270 KOKKOS_INLINE_FUNCTION long long abs(long long n) {
271 // FIXME_NVHPC ptxas fatal : unresolved extern function 'labs'
272 #if defined(KOKKOS_COMPILER_NVHPC) && KOKKOS_COMPILER_NVHPC < 230700
273  return n > 0 ? n : -n;
274 #else
275  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
276  return abs(n);
277 #endif
278 }
279 KOKKOS_INLINE_FUNCTION float abs(float x) {
280 #ifdef KOKKOS_ENABLE_SYCL
281  return sycl::fabs(x); // sycl::abs is only provided for integral types
282 #else
283  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
284  return abs(x);
285 #endif
286 }
287 KOKKOS_INLINE_FUNCTION double abs(double x) {
288 #ifdef KOKKOS_ENABLE_SYCL
289  return sycl::fabs(x); // sycl::abs is only provided for integral types
290 #else
291  using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
292  return abs(x);
293 #endif
294 }
295 inline long double abs(long double x) {
296  using std::abs;
297  return abs(x);
298 }
299 KOKKOS_IMPL_MATH_UNARY_FUNCTION(fabs)
300 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmod)
301 KOKKOS_IMPL_MATH_BINARY_FUNCTION(remainder)
302 // remquo
303 KOKKOS_IMPL_MATH_TERNARY_FUNCTION(fma)
304 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmax)
305 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmin)
306 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fdim)
307 #ifndef KOKKOS_ENABLE_SYCL
308 KOKKOS_INLINE_FUNCTION float nanf(char const* arg) { return ::nanf(arg); }
309 KOKKOS_INLINE_FUNCTION double nan(char const* arg) { return ::nan(arg); }
310 #else
311 // FIXME_SYCL
312 // sycl::nan does not follow the C/C++ standard library and takes an unsigned
313 // integer as argument. The current implementation does not attempt to convert
314 // the character string arg into the quiet NaN value.
315 KOKKOS_INLINE_FUNCTION float nanf(char const*) { return sycl::nan(0u); }
316 KOKKOS_INLINE_FUNCTION double nan(char const*) { return sycl::nan(0ul); }
317 #endif
318 inline long double nanl(char const* arg) { return ::nanl(arg); }
319 // Exponential functions
320 KOKKOS_IMPL_MATH_UNARY_FUNCTION(exp)
321 // FIXME_NVHPC nvc++ has issues with exp2
322 #if defined(KOKKOS_COMPILER_NVHPC) && KOKKOS_COMPILER_NVHPC < 230700
323 KOKKOS_INLINE_FUNCTION float exp2(float val) {
324  constexpr float ln2 = 0.693147180559945309417232121458176568L;
325  return exp(ln2 * val);
326 }
327 KOKKOS_INLINE_FUNCTION double exp2(double val) {
328  constexpr double ln2 = 0.693147180559945309417232121458176568L;
329  return exp(ln2 * val);
330 }
331 inline long double exp2(long double val) {
332  constexpr long double ln2 = 0.693147180559945309417232121458176568L;
333  return exp(ln2 * val);
334 }
335 template <class T>
336 KOKKOS_INLINE_FUNCTION double exp2(T val) {
337  constexpr double ln2 = 0.693147180559945309417232121458176568L;
338  return exp(ln2 * static_cast<double>(val));
339 }
340 #else
341 KOKKOS_IMPL_MATH_UNARY_FUNCTION(exp2)
342 #endif
343 KOKKOS_IMPL_MATH_UNARY_FUNCTION(expm1)
344 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log)
345 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log10)
346 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log2)
347 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log1p)
348 // Power functions
349 KOKKOS_IMPL_MATH_BINARY_FUNCTION(pow)
350 KOKKOS_IMPL_MATH_UNARY_FUNCTION(sqrt)
351 KOKKOS_IMPL_MATH_UNARY_FUNCTION(cbrt)
352 KOKKOS_IMPL_MATH_BINARY_FUNCTION(hypot)
353 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP) || \
354  defined(KOKKOS_ENABLE_SYCL)
355 KOKKOS_INLINE_FUNCTION float hypot(float x, float y, float z) {
356  return sqrt(x * x + y * y + z * z);
357 }
358 KOKKOS_INLINE_FUNCTION double hypot(double x, double y, double z) {
359  return sqrt(x * x + y * y + z * z);
360 }
361 inline long double hypot(long double x, long double y, long double z) {
362  return sqrt(x * x + y * y + z * z);
363 }
364 KOKKOS_INLINE_FUNCTION float hypotf(float x, float y, float z) {
365  return sqrt(x * x + y * y + z * z);
366 }
367 inline long double hypotl(long double x, long double y, long double z) {
368  return sqrt(x * x + y * y + z * z);
369 }
370 template <
371  class T1, class T2, class T3,
372  class Promoted = std::enable_if_t<
373  std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> &&
374  std::is_arithmetic_v<T3> && !std::is_same_v<T1, long double> &&
375  !std::is_same_v<T2, long double> &&
376  !std::is_same_v<T3, long double>,
377  Impl::promote_3_t<T1, T2, T3>>>
378 KOKKOS_INLINE_FUNCTION Promoted hypot(T1 x, T2 y, T3 z) {
379  return hypot(static_cast<Promoted>(x), static_cast<Promoted>(y),
380  static_cast<Promoted>(z));
381 }
382 template <
383  class T1, class T2, class T3,
384  class = std::enable_if_t<
385  std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> &&
386  std::is_arithmetic_v<T3> &&
387  (std::is_same_v<T1, long double> || std::is_same_v<T2, long double> ||
388  std::is_same_v<T3, long double>)>>
389 inline long double hypot(T1 x, T2 y, T3 z) {
390  return hypot(static_cast<long double>(x), static_cast<long double>(y),
391  static_cast<long double>(z));
392 }
393 #else
394 KOKKOS_IMPL_MATH_TERNARY_FUNCTION(hypot)
395 #endif
396 // Trigonometric functions
397 KOKKOS_IMPL_MATH_UNARY_FUNCTION(sin)
398 KOKKOS_IMPL_MATH_UNARY_FUNCTION(cos)
399 KOKKOS_IMPL_MATH_UNARY_FUNCTION(tan)
400 KOKKOS_IMPL_MATH_UNARY_FUNCTION(asin)
401 KOKKOS_IMPL_MATH_UNARY_FUNCTION(acos)
402 KOKKOS_IMPL_MATH_UNARY_FUNCTION(atan)
403 KOKKOS_IMPL_MATH_BINARY_FUNCTION(atan2)
404 // Hyperbolic functions
405 KOKKOS_IMPL_MATH_UNARY_FUNCTION(sinh)
406 KOKKOS_IMPL_MATH_UNARY_FUNCTION(cosh)
407 KOKKOS_IMPL_MATH_UNARY_FUNCTION(tanh)
408 KOKKOS_IMPL_MATH_UNARY_FUNCTION(asinh)
409 KOKKOS_IMPL_MATH_UNARY_FUNCTION(acosh)
410 KOKKOS_IMPL_MATH_UNARY_FUNCTION(atanh)
411 // Error and gamma functions
412 KOKKOS_IMPL_MATH_UNARY_FUNCTION(erf)
413 KOKKOS_IMPL_MATH_UNARY_FUNCTION(erfc)
414 KOKKOS_IMPL_MATH_UNARY_FUNCTION(tgamma)
415 KOKKOS_IMPL_MATH_UNARY_FUNCTION(lgamma)
416 // Nearest integer floating point operations
417 KOKKOS_IMPL_MATH_UNARY_FUNCTION(ceil)
418 KOKKOS_IMPL_MATH_UNARY_FUNCTION(floor)
419 KOKKOS_IMPL_MATH_UNARY_FUNCTION(trunc)
420 KOKKOS_IMPL_MATH_UNARY_FUNCTION(round)
421 // lround
422 // llround
423 // FIXME_SYCL not available as of current SYCL 2020 specification (revision 4)
424 #ifndef KOKKOS_ENABLE_SYCL // FIXME_SYCL
425 KOKKOS_IMPL_MATH_UNARY_FUNCTION(nearbyint)
426 #endif
427 // rint
428 // lrint
429 // llrint
430 // Floating point manipulation functions
431 // frexp
432 // ldexp
433 // modf
434 // scalbn
435 // scalbln
436 // ilog
437 KOKKOS_IMPL_MATH_UNARY_FUNCTION(logb)
438 KOKKOS_IMPL_MATH_BINARY_FUNCTION(nextafter)
439 // nexttoward
440 KOKKOS_IMPL_MATH_BINARY_FUNCTION(copysign)
441 // Classification and comparison
442 // fpclassify
443 KOKKOS_IMPL_MATH_UNARY_PREDICATE(isfinite)
444 KOKKOS_IMPL_MATH_UNARY_PREDICATE(isinf)
445 KOKKOS_IMPL_MATH_UNARY_PREDICATE(isnan)
446 // isnormal
447 KOKKOS_IMPL_MATH_UNARY_PREDICATE(signbit)
448 // isgreater
449 // isgreaterequal
450 // isless
451 // islessequal
452 // islessgreater
453 // isunordered
454 
455 #undef KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE
456 #undef KOKKOS_IMPL_MATH_UNARY_FUNCTION
457 #undef KOKKOS_IMPL_MATH_UNARY_PREDICATE
458 #undef KOKKOS_IMPL_MATH_BINARY_FUNCTION
459 #undef KOKKOS_IMPL_MATH_TERNARY_FUNCTION
460 
461 // non-standard math functions provided by CUDA/HIP/SYCL
462 KOKKOS_INLINE_FUNCTION float rsqrt(float val) {
463 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
464  KOKKOS_IF_ON_DEVICE(return ::rsqrtf(val);)
465  KOKKOS_IF_ON_HOST(return 1.0f / Kokkos::sqrt(val);)
466 #elif defined(KOKKOS_ENABLE_SYCL)
467  KOKKOS_IF_ON_DEVICE(return sycl::rsqrt(val);)
468  KOKKOS_IF_ON_HOST(return 1.0f / Kokkos::sqrt(val);)
469 #else
470  return 1.0f / Kokkos::sqrt(val);
471 #endif
472 }
473 KOKKOS_INLINE_FUNCTION double rsqrt(double val) {
474 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP)
475  KOKKOS_IF_ON_DEVICE(return ::rsqrt(val);)
476  KOKKOS_IF_ON_HOST(return 1.0 / Kokkos::sqrt(val);)
477 #elif defined(KOKKOS_ENABLE_SYCL)
478  KOKKOS_IF_ON_DEVICE(return sycl::rsqrt(val);)
479  KOKKOS_IF_ON_HOST(return 1.0 / Kokkos::sqrt(val);)
480 #else
481  return 1.0 / Kokkos::sqrt(val);
482 #endif
483 }
484 inline long double rsqrt(long double val) { return 1.0l / Kokkos::sqrt(val); }
485 KOKKOS_INLINE_FUNCTION float rsqrtf(float x) { return Kokkos::rsqrt(x); }
486 inline long double rsqrtl(long double x) { return Kokkos::rsqrt(x); }
487 template <class T>
488 KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, double> rsqrt(
489  T x) {
490  return Kokkos::rsqrt(static_cast<double>(x));
491 }
492 
493 } // namespace Kokkos
494 
495 #ifdef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_MATHFUNCTIONS
496 #undef KOKKOS_IMPL_PUBLIC_INCLUDE
497 #undef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_MATHFUNCTIONS
498 #endif
499 #endif
Generated by   doxygen 1.8.5