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 24 #include <Kokkos_Macros.hpp> 27 #include <type_traits> 29 #ifdef KOKKOS_ENABLE_SYCL 31 #if __has_include(<sycl/sycl.hpp>) 32 #include <sycl/sycl.hpp> 34 #include <CL/sycl.hpp> 41 template <
class T,
bool = std::is_
integral_v<T>>
46 struct promote<T, false> {};
48 struct promote<long double> {
49 using type =
long double;
52 struct promote<double> {
56 struct promote<float> {
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>>
64 using type = decltype(promote_t<T>() + promote_t<U>());
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>>
74 using type = decltype(promote_t<T>() + promote_t<U>() + promote_t<V>());
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;
84 #if defined(KOKKOS_ENABLE_SYCL) 85 #define KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE sycl 87 #if (defined(KOKKOS_COMPILER_NVCC) || defined(KOKKOS_COMPILER_NVHPC)) && \ 88 defined(__GNUC__) && (__GNUC__ < 6) && !defined(__clang__) 89 #define KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE 91 #define KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE std 95 #if defined(KOKKOS_ENABLE_DEPRECATED_CODE_3) 96 #define KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \ 97 USING_DECLARATIONS_IN_EXPERIMENTAL_NAMESPACE) \ 98 USING_DECLARATIONS_IN_EXPERIMENTAL_NAMESPACE 100 #define KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \ 101 USING_DECLARATIONS_IN_EXPERIMENTAL_NAMESPACE) \ 105 #define KOKKOS_IMPL_MATH_UNARY_FUNCTION(FUNC) \ 106 KOKKOS_INLINE_FUNCTION float FUNC(float x) { \ 107 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 110 KOKKOS_INLINE_FUNCTION double FUNC(double x) { \ 111 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 114 inline long double FUNC(long double x) { \ 118 KOKKOS_INLINE_FUNCTION float FUNC##f(float x) { \ 119 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 122 inline long double FUNC##l(long double x) { \ 127 KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, double> FUNC( \ 129 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 130 return FUNC(static_cast<double>(x)); \ 132 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \ 133 namespace Experimental { \ 134 using ::Kokkos::FUNC; \ 135 using ::Kokkos::FUNC##f; \ 136 using ::Kokkos::FUNC##l; \ 142 #if defined(_WIN32) && defined(KOKKOS_ENABLE_CUDA) 143 #define KOKKOS_IMPL_MATH_UNARY_PREDICATE(FUNC) \ 144 KOKKOS_INLINE_FUNCTION bool FUNC(float x) { return ::FUNC(x); } \ 145 KOKKOS_INLINE_FUNCTION bool FUNC(double x) { return ::FUNC(x); } \ 146 inline bool FUNC(long double x) { \ 151 KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, bool> FUNC( \ 153 return ::FUNC(static_cast<double>(x)); \ 155 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \ 156 namespace Experimental { using ::Kokkos::FUNC; }) 158 #define KOKKOS_IMPL_MATH_UNARY_PREDICATE(FUNC) \ 159 KOKKOS_INLINE_FUNCTION bool FUNC(float x) { \ 160 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 163 KOKKOS_INLINE_FUNCTION bool FUNC(double x) { \ 164 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 167 inline bool FUNC(long double x) { \ 172 KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, bool> FUNC( \ 174 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 175 return FUNC(static_cast<double>(x)); \ 177 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \ 178 namespace Experimental { using ::Kokkos::FUNC; }) 181 #define KOKKOS_IMPL_MATH_BINARY_FUNCTION(FUNC) \ 182 KOKKOS_INLINE_FUNCTION float FUNC(float x, float y) { \ 183 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 186 KOKKOS_INLINE_FUNCTION double FUNC(double x, double y) { \ 187 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 190 inline long double FUNC(long double x, long double y) { \ 194 KOKKOS_INLINE_FUNCTION float FUNC##f(float x, float y) { \ 195 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 198 inline long double FUNC##l(long double x, long double y) { \ 202 template <class T1, class T2> \ 203 KOKKOS_INLINE_FUNCTION \ 204 std::enable_if_t<std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> && \ 205 !std::is_same_v<T1, long double> && \ 206 !std::is_same_v<T2, long double>, \ 207 Kokkos::Impl::promote_2_t<T1, T2>> \ 209 using Promoted = Kokkos::Impl::promote_2_t<T1, T2>; \ 210 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 211 return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y)); \ 213 template <class T1, class T2> \ 214 inline std::enable_if_t<std::is_arithmetic_v<T1> && \ 215 std::is_arithmetic_v<T2> && \ 216 (std::is_same_v<T1, long double> || \ 217 std::is_same_v<T2, long double>), \ 220 using Promoted = Kokkos::Impl::promote_2_t<T1, T2>; \ 221 static_assert(std::is_same_v<Promoted, long double>, ""); \ 223 return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y)); \ 225 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \ 226 namespace Experimental { \ 227 using ::Kokkos::FUNC; \ 228 using ::Kokkos::FUNC##f; \ 229 using ::Kokkos::FUNC##l; \ 232 #define KOKKOS_IMPL_MATH_TERNARY_FUNCTION(FUNC) \ 233 KOKKOS_INLINE_FUNCTION float FUNC(float x, float y, float z) { \ 234 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 235 return FUNC(x, y, z); \ 237 KOKKOS_INLINE_FUNCTION double FUNC(double x, double y, double z) { \ 238 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 239 return FUNC(x, y, z); \ 241 inline long double FUNC(long double x, long double y, long double z) { \ 243 return FUNC(x, y, z); \ 245 KOKKOS_INLINE_FUNCTION float FUNC##f(float x, float y, float z) { \ 246 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 247 return FUNC(x, y, z); \ 249 inline long double FUNC##l(long double x, long double y, long double z) { \ 251 return FUNC(x, y, z); \ 253 template <class T1, class T2, class T3> \ 254 KOKKOS_INLINE_FUNCTION std::enable_if_t< \ 255 std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> && \ 256 std::is_arithmetic_v<T3> && !std::is_same_v<T1, long double> && \ 257 !std::is_same_v<T2, long double> && \ 258 !std::is_same_v<T3, long double>, \ 259 Kokkos::Impl::promote_3_t<T1, T2, T3>> \ 260 FUNC(T1 x, T2 y, T3 z) { \ 261 using Promoted = Kokkos::Impl::promote_3_t<T1, T2, T3>; \ 262 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \ 263 return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y), \ 264 static_cast<Promoted>(z)); \ 266 template <class T1, class T2, class T3> \ 267 inline std::enable_if_t<std::is_arithmetic_v<T1> && \ 268 std::is_arithmetic_v<T2> && \ 269 std::is_arithmetic_v<T3> && \ 270 (std::is_same_v<T1, long double> || \ 271 std::is_same_v<T2, long double> || \ 272 std::is_same_v<T3, long double>), \ 274 FUNC(T1 x, T2 y, T3 z) { \ 275 using Promoted = Kokkos::Impl::promote_3_t<T1, T2, T3>; \ 276 static_assert(std::is_same_v<Promoted, long double>); \ 278 return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y), \ 279 static_cast<Promoted>(z)); \ 283 KOKKOS_INLINE_FUNCTION
int abs(
int n) {
284 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
287 KOKKOS_INLINE_FUNCTION
long abs(
long n) {
289 #ifdef KOKKOS_COMPILER_NVHPC 290 return n > 0 ? n : -n;
292 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
296 KOKKOS_INLINE_FUNCTION
long long abs(
long long n) {
298 #ifdef KOKKOS_COMPILER_NVHPC 299 return n > 0 ? n : -n;
301 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
305 KOKKOS_INLINE_FUNCTION
float abs(
float x) {
306 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
309 KOKKOS_INLINE_FUNCTION
double abs(
double x) {
310 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
313 inline long double abs(
long double x) {
317 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED(
319 KOKKOS_IMPL_MATH_UNARY_FUNCTION(fabs)
320 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmod)
321 KOKKOS_IMPL_MATH_BINARY_FUNCTION(remainder)
323 KOKKOS_IMPL_MATH_TERNARY_FUNCTION(fma)
324 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmax)
325 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmin)
326 KOKKOS_IMPL_MATH_BINARY_FUNCTION(fdim)
327 #ifndef KOKKOS_ENABLE_SYCL 328 KOKKOS_INLINE_FUNCTION
float nanf(
char const* arg) { return ::nanf(arg); }
329 KOKKOS_INLINE_FUNCTION
double nan(
char const* arg) { return ::nan(arg); }
335 KOKKOS_INLINE_FUNCTION
float nanf(
char const*) {
return sycl::nan(0u); }
336 KOKKOS_INLINE_FUNCTION
double nan(
char const*) {
return sycl::nan(0ul); }
338 inline long double nanl(
char const* arg) { return ::nanl(arg); }
339 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED(
342 using ::Kokkos::nanf;
343 using ::Kokkos::nanl;
346 KOKKOS_IMPL_MATH_UNARY_FUNCTION(exp)
348 #ifndef KOKKOS_COMPILER_NVHPC 349 KOKKOS_IMPL_MATH_UNARY_FUNCTION(exp2)
351 KOKKOS_INLINE_FUNCTION
float exp2(
float val) {
352 constexpr
float ln2 = 0.693147180559945309417232121458176568L;
353 return exp(ln2 * val);
355 KOKKOS_INLINE_FUNCTION
double exp2(
double val) {
356 constexpr
double ln2 = 0.693147180559945309417232121458176568L;
357 return exp(ln2 * val);
359 inline long double exp2(
long double val) {
360 constexpr
long double ln2 = 0.693147180559945309417232121458176568L;
361 return exp(ln2 * val);
364 KOKKOS_INLINE_FUNCTION
double exp2(T val) {
365 constexpr
double ln2 = 0.693147180559945309417232121458176568L;
366 return exp(ln2 * static_cast<double>(val));
369 KOKKOS_IMPL_MATH_UNARY_FUNCTION(expm1)
370 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log)
371 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log10)
372 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log2)
373 KOKKOS_IMPL_MATH_UNARY_FUNCTION(log1p)
375 KOKKOS_IMPL_MATH_BINARY_FUNCTION(pow)
376 KOKKOS_IMPL_MATH_UNARY_FUNCTION(sqrt)
377 KOKKOS_IMPL_MATH_UNARY_FUNCTION(cbrt)
378 KOKKOS_IMPL_MATH_BINARY_FUNCTION(hypot)
379 #if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP) || \ 380 defined(KOKKOS_ENABLE_SYCL) 381 KOKKOS_INLINE_FUNCTION
float hypot(
float x,
float y,
float z) {
382 return sqrt(x * x + y * y + z * z);
384 KOKKOS_INLINE_FUNCTION
double hypot(
double x,
double y,
double z) {
385 return sqrt(x * x + y * y + z * z);
387 inline long double hypot(
long double x,
long double y,
long double z) {
388 return sqrt(x * x + y * y + z * z);
390 KOKKOS_INLINE_FUNCTION
float hypotf(
float x,
float y,
float z) {
391 return sqrt(x * x + y * y + z * z);
393 inline long double hypotl(
long double x,
long double y,
long double z) {
394 return sqrt(x * x + y * y + z * z);
397 class T1,
class T2,
class T3,
398 class Promoted = std::enable_if_t<
399 std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> &&
400 std::is_arithmetic_v<T3> && !std::is_same_v<T1, long double> &&
401 !std::is_same_v<T2, long double> &&
402 !std::is_same_v<T3, long double>,
403 Impl::promote_3_t<T1, T2, T3>>>
404 KOKKOS_INLINE_FUNCTION Promoted hypot(T1 x, T2 y, T3 z) {
405 return hypot(static_cast<Promoted>(x), static_cast<Promoted>(y),
406 static_cast<Promoted>(z));
409 class T1,
class T2,
class T3,
410 class = std::enable_if_t<
411 std::is_arithmetic_v<T1> && std::is_arithmetic_v<T2> &&
412 std::is_arithmetic_v<T3> &&
413 (std::is_same_v<T1, long double> || std::is_same_v<T2, long double> ||
414 std::is_same_v<T3, long double>)>>
415 inline long double hypot(T1 x, T2 y, T3 z) {
416 return hypot(static_cast<long double>(x), static_cast<long double>(y),
417 static_cast<long double>(z));
420 KOKKOS_IMPL_MATH_TERNARY_FUNCTION(hypot)
423 KOKKOS_IMPL_MATH_UNARY_FUNCTION(sin)
424 KOKKOS_IMPL_MATH_UNARY_FUNCTION(cos)
425 KOKKOS_IMPL_MATH_UNARY_FUNCTION(tan)
426 KOKKOS_IMPL_MATH_UNARY_FUNCTION(asin)
427 KOKKOS_IMPL_MATH_UNARY_FUNCTION(acos)
428 KOKKOS_IMPL_MATH_UNARY_FUNCTION(atan)
429 KOKKOS_IMPL_MATH_BINARY_FUNCTION(atan2)
431 KOKKOS_IMPL_MATH_UNARY_FUNCTION(sinh)
432 KOKKOS_IMPL_MATH_UNARY_FUNCTION(cosh)
433 KOKKOS_IMPL_MATH_UNARY_FUNCTION(tanh)
434 KOKKOS_IMPL_MATH_UNARY_FUNCTION(asinh)
435 KOKKOS_IMPL_MATH_UNARY_FUNCTION(acosh)
436 KOKKOS_IMPL_MATH_UNARY_FUNCTION(atanh)
438 KOKKOS_IMPL_MATH_UNARY_FUNCTION(erf)
439 KOKKOS_IMPL_MATH_UNARY_FUNCTION(erfc)
440 KOKKOS_IMPL_MATH_UNARY_FUNCTION(tgamma)
441 KOKKOS_IMPL_MATH_UNARY_FUNCTION(lgamma)
443 KOKKOS_IMPL_MATH_UNARY_FUNCTION(ceil)
444 KOKKOS_IMPL_MATH_UNARY_FUNCTION(floor)
445 KOKKOS_IMPL_MATH_UNARY_FUNCTION(trunc)
446 KOKKOS_IMPL_MATH_UNARY_FUNCTION(round)
450 #ifndef KOKKOS_ENABLE_SYCL // FIXME_SYCL 451 KOKKOS_IMPL_MATH_UNARY_FUNCTION(nearbyint)
463 KOKKOS_IMPL_MATH_UNARY_FUNCTION(logb)
464 KOKKOS_IMPL_MATH_BINARY_FUNCTION(nextafter)
466 KOKKOS_IMPL_MATH_BINARY_FUNCTION(copysign)
469 KOKKOS_IMPL_MATH_UNARY_PREDICATE(isfinite)
470 KOKKOS_IMPL_MATH_UNARY_PREDICATE(isinf)
471 KOKKOS_IMPL_MATH_UNARY_PREDICATE(isnan)
473 KOKKOS_IMPL_MATH_UNARY_PREDICATE(signbit)
481 #undef KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED 482 #undef KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE 483 #undef KOKKOS_IMPL_MATH_UNARY_FUNCTION 484 #undef KOKKOS_IMPL_MATH_UNARY_PREDICATE 485 #undef KOKKOS_IMPL_MATH_BINARY_FUNCTION 486 #undef KOKKOS_IMPL_MATH_TERNARY_FUNCTION 490 #ifdef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_MATHFUNCTIONS 491 #undef KOKKOS_IMPL_PUBLIC_INCLUDE 492 #undef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_MATHFUNCTIONS