Kokkos Core Kernels Package Version of the Day
Loading...
Searching...
No Matches
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
38namespace Kokkos {
39
40namespace Impl {
41template <class T, bool = std::is_integral_v<T>>
42struct promote {
43 using type = double;
44};
45template <class T>
46struct promote<T, false> {};
47template <>
48struct promote<long double> {
49 using type = long double;
50};
51template <>
52struct promote<double> {
53 using type = double;
54};
55template <>
56struct promote<float> {
57 using type = float;
58};
59template <class T>
60using promote_t = typename promote<T>::type;
61template <class T, class U,
62 bool = std::is_arithmetic_v<T>&& std::is_arithmetic_v<U>>
63struct promote_2 {
64 using type = decltype(promote_t<T>() + promote_t<U>());
65};
66template <class T, class U>
67struct promote_2<T, U, false> {};
68template <class T, class U>
69using promote_2_t = typename promote_2<T, U>::type;
70template <class T, class U, class V,
71 bool = std::is_arithmetic_v<T>&& std::is_arithmetic_v<U>&&
72 std::is_arithmetic_v<V>>
73struct promote_3 {
74 using type = decltype(promote_t<T>() + promote_t<U>() + promote_t<V>());
75};
76template <class T, class U, class V>
77struct promote_3<T, U, V, false> {};
78template <class T, class U, class V>
79using 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#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
99#else
100#define KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \
101 USING_DECLARATIONS_IN_EXPERIMENTAL_NAMESPACE) \
102 /* nothing */
103#endif
104
105#define KOKKOS_IMPL_MATH_UNARY_FUNCTION(FUNC) \
106 KOKKOS_INLINE_FUNCTION float FUNC(float x) { \
107 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
108 return FUNC(x); \
109 } \
110 KOKKOS_INLINE_FUNCTION double FUNC(double x) { \
111 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
112 return FUNC(x); \
113 } \
114 inline long double FUNC(long double x) { \
115 using std::FUNC; \
116 return FUNC(x); \
117 } \
118 KOKKOS_INLINE_FUNCTION float FUNC##f(float x) { \
119 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
120 return FUNC(x); \
121 } \
122 inline long double FUNC##l(long double x) { \
123 using std::FUNC; \
124 return FUNC(x); \
125 } \
126 template <class T> \
127 KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, double> FUNC( \
128 T x) { \
129 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
130 return FUNC(static_cast<double>(x)); \
131 } \
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; \
137 })
138
139// isinf, isnan, and isinfinite do not work on Windows with CUDA with std::
140// getting warnings about calling host function in device function then
141// runtime test fails
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) { \
147 using std::FUNC; \
148 return FUNC(x); \
149 } \
150 template <class T> \
151 KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, bool> FUNC( \
152 T x) { \
153 return ::FUNC(static_cast<double>(x)); \
154 } \
155 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \
156 namespace Experimental { using ::Kokkos::FUNC; })
157#else
158#define KOKKOS_IMPL_MATH_UNARY_PREDICATE(FUNC) \
159 KOKKOS_INLINE_FUNCTION bool FUNC(float x) { \
160 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
161 return FUNC(x); \
162 } \
163 KOKKOS_INLINE_FUNCTION bool FUNC(double x) { \
164 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
165 return FUNC(x); \
166 } \
167 inline bool FUNC(long double x) { \
168 using std::FUNC; \
169 return FUNC(x); \
170 } \
171 template <class T> \
172 KOKKOS_INLINE_FUNCTION std::enable_if_t<std::is_integral_v<T>, bool> FUNC( \
173 T x) { \
174 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
175 return FUNC(static_cast<double>(x)); \
176 } \
177 KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED( \
178 namespace Experimental { using ::Kokkos::FUNC; })
179#endif
180
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; \
184 return FUNC(x, y); \
185 } \
186 KOKKOS_INLINE_FUNCTION double FUNC(double x, double y) { \
187 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
188 return FUNC(x, y); \
189 } \
190 inline long double FUNC(long double x, long double y) { \
191 using std::FUNC; \
192 return FUNC(x, y); \
193 } \
194 KOKKOS_INLINE_FUNCTION float FUNC##f(float x, float y) { \
195 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::FUNC; \
196 return FUNC(x, y); \
197 } \
198 inline long double FUNC##l(long double x, long double y) { \
199 using std::FUNC; \
200 return FUNC(x, y); \
201 } \
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>> \
208 FUNC(T1 x, T2 y) { \
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)); \
212 } \
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>), \
218 long double> \
219 FUNC(T1 x, T2 y) { \
220 using Promoted = Kokkos::Impl::promote_2_t<T1, T2>; \
221 static_assert(std::is_same_v<Promoted, long double>, ""); \
222 using std::FUNC; \
223 return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y)); \
224 } \
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; \
230 })
231
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); \
236 } \
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); \
240 } \
241 inline long double FUNC(long double x, long double y, long double z) { \
242 using std::FUNC; \
243 return FUNC(x, y, z); \
244 } \
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); \
248 } \
249 inline long double FUNC##l(long double x, long double y, long double z) { \
250 using std::FUNC; \
251 return FUNC(x, y, z); \
252 } \
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)); \
265 } \
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>), \
273 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>); \
277 using std::FUNC; \
278 return FUNC(static_cast<Promoted>(x), static_cast<Promoted>(y), \
279 static_cast<Promoted>(z)); \
280 }
281
282// Basic operations
283KOKKOS_INLINE_FUNCTION int abs(int n) {
284 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
285 return abs(n);
286}
287KOKKOS_INLINE_FUNCTION long abs(long n) {
288// FIXME_NVHPC ptxas fatal : unresolved extern function 'labs'
289#ifdef KOKKOS_COMPILER_NVHPC
290 return n > 0 ? n : -n;
291#else
292 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
293 return abs(n);
294#endif
295}
296KOKKOS_INLINE_FUNCTION long long abs(long long n) {
297// FIXME_NVHPC ptxas fatal : unresolved extern function 'labs'
298#ifdef KOKKOS_COMPILER_NVHPC
299 return n > 0 ? n : -n;
300#else
301 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
302 return abs(n);
303#endif
304}
305KOKKOS_INLINE_FUNCTION float abs(float x) {
306 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
307 return abs(x);
308}
309KOKKOS_INLINE_FUNCTION double abs(double x) {
310 using KOKKOS_IMPL_MATH_FUNCTIONS_NAMESPACE::abs;
311 return abs(x);
312}
313inline long double abs(long double x) {
314 using std::abs;
315 return abs(x);
316}
317KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED(
318 namespace Experimental { using ::Kokkos::abs; })
319KOKKOS_IMPL_MATH_UNARY_FUNCTION(fabs)
320KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmod)
321KOKKOS_IMPL_MATH_BINARY_FUNCTION(remainder)
322// remquo
323KOKKOS_IMPL_MATH_TERNARY_FUNCTION(fma)
324KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmax)
325KOKKOS_IMPL_MATH_BINARY_FUNCTION(fmin)
326KOKKOS_IMPL_MATH_BINARY_FUNCTION(fdim)
327#ifndef KOKKOS_ENABLE_SYCL
328KOKKOS_INLINE_FUNCTION float nanf(char const* arg) { return ::nanf(arg); }
329KOKKOS_INLINE_FUNCTION double nan(char const* arg) { return ::nan(arg); }
330#else
331// FIXME_SYCL
332// sycl::nan does not follow the C/C++ standard library and takes an unsigned
333// integer as argument. The current implementation does not attempt to convert
334// the character string arg into the quiet NaN value.
335KOKKOS_INLINE_FUNCTION float nanf(char const*) { return sycl::nan(0u); }
336KOKKOS_INLINE_FUNCTION double nan(char const*) { return sycl::nan(0ul); }
337#endif
338inline long double nanl(char const* arg) { return ::nanl(arg); }
339KOKKOS_IMPL_MATH_FUNCTIONS_DEFINED_IF_DEPRECATED_CODE_ENABLED(
340 namespace Experimental {
341 using ::Kokkos::nan;
342 using ::Kokkos::nanf;
343 using ::Kokkos::nanl;
344 })
345// Exponential functions
346KOKKOS_IMPL_MATH_UNARY_FUNCTION(exp)
347// FIXME_NVHPC nvc++ has issues with exp2
348#ifndef KOKKOS_COMPILER_NVHPC
349KOKKOS_IMPL_MATH_UNARY_FUNCTION(exp2)
350#else
351KOKKOS_INLINE_FUNCTION float exp2(float val) {
352 constexpr float ln2 = 0.693147180559945309417232121458176568L;
353 return exp(ln2 * val);
354}
355KOKKOS_INLINE_FUNCTION double exp2(double val) {
356 constexpr double ln2 = 0.693147180559945309417232121458176568L;
357 return exp(ln2 * val);
358}
359inline long double exp2(long double val) {
360 constexpr long double ln2 = 0.693147180559945309417232121458176568L;
361 return exp(ln2 * val);
362}
363template <class T>
364KOKKOS_INLINE_FUNCTION double exp2(T val) {
365 constexpr double ln2 = 0.693147180559945309417232121458176568L;
366 return exp(ln2 * static_cast<double>(val));
367}
368#endif
369KOKKOS_IMPL_MATH_UNARY_FUNCTION(expm1)
370KOKKOS_IMPL_MATH_UNARY_FUNCTION(log)
371KOKKOS_IMPL_MATH_UNARY_FUNCTION(log10)
372KOKKOS_IMPL_MATH_UNARY_FUNCTION(log2)
373KOKKOS_IMPL_MATH_UNARY_FUNCTION(log1p)
374// Power functions
375KOKKOS_IMPL_MATH_BINARY_FUNCTION(pow)
376KOKKOS_IMPL_MATH_UNARY_FUNCTION(sqrt)
377KOKKOS_IMPL_MATH_UNARY_FUNCTION(cbrt)
378KOKKOS_IMPL_MATH_BINARY_FUNCTION(hypot)
379#if defined(KOKKOS_ENABLE_CUDA) || defined(KOKKOS_ENABLE_HIP) || \
380 defined(KOKKOS_ENABLE_SYCL)
381KOKKOS_INLINE_FUNCTION float hypot(float x, float y, float z) {
382 return sqrt(x * x + y * y + z * z);
383}
384KOKKOS_INLINE_FUNCTION double hypot(double x, double y, double z) {
385 return sqrt(x * x + y * y + z * z);
386}
387inline long double hypot(long double x, long double y, long double z) {
388 return sqrt(x * x + y * y + z * z);
389}
390KOKKOS_INLINE_FUNCTION float hypotf(float x, float y, float z) {
391 return sqrt(x * x + y * y + z * z);
392}
393inline long double hypotl(long double x, long double y, long double z) {
394 return sqrt(x * x + y * y + z * z);
395}
396template <
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>>>
404KOKKOS_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));
407}
408template <
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>)>>
415inline 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));
418}
419#else
420KOKKOS_IMPL_MATH_TERNARY_FUNCTION(hypot)
421#endif
422// Trigonometric functions
423KOKKOS_IMPL_MATH_UNARY_FUNCTION(sin)
424KOKKOS_IMPL_MATH_UNARY_FUNCTION(cos)
425KOKKOS_IMPL_MATH_UNARY_FUNCTION(tan)
426KOKKOS_IMPL_MATH_UNARY_FUNCTION(asin)
427KOKKOS_IMPL_MATH_UNARY_FUNCTION(acos)
428KOKKOS_IMPL_MATH_UNARY_FUNCTION(atan)
429KOKKOS_IMPL_MATH_BINARY_FUNCTION(atan2)
430// Hyperbolic functions
431KOKKOS_IMPL_MATH_UNARY_FUNCTION(sinh)
432KOKKOS_IMPL_MATH_UNARY_FUNCTION(cosh)
433KOKKOS_IMPL_MATH_UNARY_FUNCTION(tanh)
434KOKKOS_IMPL_MATH_UNARY_FUNCTION(asinh)
435KOKKOS_IMPL_MATH_UNARY_FUNCTION(acosh)
436KOKKOS_IMPL_MATH_UNARY_FUNCTION(atanh)
437// Error and gamma functions
438KOKKOS_IMPL_MATH_UNARY_FUNCTION(erf)
439KOKKOS_IMPL_MATH_UNARY_FUNCTION(erfc)
440KOKKOS_IMPL_MATH_UNARY_FUNCTION(tgamma)
441KOKKOS_IMPL_MATH_UNARY_FUNCTION(lgamma)
442// Nearest integer floating point operations
443KOKKOS_IMPL_MATH_UNARY_FUNCTION(ceil)
444KOKKOS_IMPL_MATH_UNARY_FUNCTION(floor)
445KOKKOS_IMPL_MATH_UNARY_FUNCTION(trunc)
446KOKKOS_IMPL_MATH_UNARY_FUNCTION(round)
447// lround
448// llround
449// FIXME_SYCL not available as of current SYCL 2020 specification (revision 4)
450#ifndef KOKKOS_ENABLE_SYCL // FIXME_SYCL
451KOKKOS_IMPL_MATH_UNARY_FUNCTION(nearbyint)
452#endif
453// rint
454// lrint
455// llrint
456// Floating point manipulation functions
457// frexp
458// ldexp
459// modf
460// scalbn
461// scalbln
462// ilog
463KOKKOS_IMPL_MATH_UNARY_FUNCTION(logb)
464KOKKOS_IMPL_MATH_BINARY_FUNCTION(nextafter)
465// nexttoward
466KOKKOS_IMPL_MATH_BINARY_FUNCTION(copysign)
467// Classification and comparison
468// fpclassify
469KOKKOS_IMPL_MATH_UNARY_PREDICATE(isfinite)
470KOKKOS_IMPL_MATH_UNARY_PREDICATE(isinf)
471KOKKOS_IMPL_MATH_UNARY_PREDICATE(isnan)
472// isnormal
473KOKKOS_IMPL_MATH_UNARY_PREDICATE(signbit)
474// isgreater
475// isgreaterequal
476// isless
477// islessequal
478// islessgreater
479// isunordered
480
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
487
488} // namespace Kokkos
489
490#ifdef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_MATHFUNCTIONS
491#undef KOKKOS_IMPL_PUBLIC_INCLUDE
492#undef KOKKOS_IMPL_PUBLIC_INCLUDE_NOTDEFINED_MATHFUNCTIONS
493#endif
494#endif