xref: /llvm-project/clang/lib/Headers/__clang_hip_cmath.h (revision db1d3b23a37c7a57fa8b9e5bc94e1b22e278d361) 
1 /*===---- __clang_hip_cmath.h - HIP cmath decls -----------------------------===
2  *
3  * Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4  * See https://llvm.org/LICENSE.txt for license information.
5  * SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6  *
7  *===-----------------------------------------------------------------------===
8  */
9 
10 #ifndef __CLANG_HIP_CMATH_H__
11 #define __CLANG_HIP_CMATH_H__
12 
13 #if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
14 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
15 #endif
16 
17 #if !defined(__HIPCC_RTC__)
18 #if defined(__cplusplus)
19 #include <limits>
20 #include <type_traits>
21 #include <utility>
22 #endif
23 #include <limits.h>
24 #include <stdint.h>
25 #endif // !defined(__HIPCC_RTC__)
26 
27 #pragma push_macro("__DEVICE__")
28 #pragma push_macro("__CONSTEXPR__")
29 #ifdef __OPENMP_AMDGCN__
30 #define __DEVICE__ static __attribute__((always_inline, nothrow))
31 #define __CONSTEXPR__ constexpr
32 #else
33 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
34 #define __CONSTEXPR__
35 #endif // __OPENMP_AMDGCN__
36 
37 // Start with functions that cannot be defined by DEF macros below.
38 #if defined(__cplusplus)
39 #if defined __OPENMP_AMDGCN__
40 __DEVICE__ __CONSTEXPR__ float fabs(float __x) { return ::fabsf(__x); }
41 __DEVICE__ __CONSTEXPR__ float sin(float __x) { return ::sinf(__x); }
42 __DEVICE__ __CONSTEXPR__ float cos(float __x) { return ::cosf(__x); }
43 #endif
44 __DEVICE__ __CONSTEXPR__ double abs(double __x) { return ::fabs(__x); }
45 __DEVICE__ __CONSTEXPR__ float abs(float __x) { return ::fabsf(__x); }
46 __DEVICE__ __CONSTEXPR__ long long abs(long long __n) { return ::llabs(__n); }
47 __DEVICE__ __CONSTEXPR__ long abs(long __n) { return ::labs(__n); }
48 __DEVICE__ __CONSTEXPR__ float fma(float __x, float __y, float __z) {
49   return ::fmaf(__x, __y, __z);
50 }
51 #if !defined(__HIPCC_RTC__)
52 // The value returned by fpclassify is platform dependent, therefore it is not
53 // supported by hipRTC.
54 __DEVICE__ __CONSTEXPR__ int fpclassify(float __x) {
55   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
56                               FP_ZERO, __x);
57 }
58 __DEVICE__ __CONSTEXPR__ int fpclassify(double __x) {
59   return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
60                               FP_ZERO, __x);
61 }
62 #endif // !defined(__HIPCC_RTC__)
63 
64 __DEVICE__ __CONSTEXPR__ float frexp(float __arg, int *__exp) {
65   return ::frexpf(__arg, __exp);
66 }
67 
68 #if defined(__OPENMP_AMDGCN__)
69 // For OpenMP we work around some old system headers that have non-conforming
70 // `isinf(float)` and `isnan(float)` implementations that return an `int`. We do
71 // this by providing two versions of these functions, differing only in the
72 // return type. To avoid conflicting definitions we disable implicit base
73 // function generation. That means we will end up with two specializations, one
74 // per type, but only one has a base function defined by the system header.
75 #pragma omp begin declare variant match(                                       \
76     implementation = {extension(disable_implicit_base)})
77 
78 // FIXME: We lack an extension to customize the mangling of the variants, e.g.,
79 //        add a suffix. This means we would clash with the names of the variants
80 //        (note that we do not create implicit base functions here). To avoid
81 //        this clash we add a new trait to some of them that is always true
82 //        (this is LLVM after all ;)). It will only influence the mangled name
83 //        of the variants inside the inner region and avoid the clash.
84 #pragma omp begin declare variant match(implementation = {vendor(llvm)})
85 
86 __DEVICE__ __CONSTEXPR__ int isinf(float __x) { return ::__isinff(__x); }
87 __DEVICE__ __CONSTEXPR__ int isinf(double __x) { return ::__isinf(__x); }
88 __DEVICE__ __CONSTEXPR__ int isfinite(float __x) { return ::__finitef(__x); }
89 __DEVICE__ __CONSTEXPR__ int isfinite(double __x) { return ::__finite(__x); }
90 __DEVICE__ __CONSTEXPR__ int isnan(float __x) { return ::__isnanf(__x); }
91 __DEVICE__ __CONSTEXPR__ int isnan(double __x) { return ::__isnan(__x); }
92 
93 #pragma omp end declare variant
94 #endif // defined(__OPENMP_AMDGCN__)
95 
96 __DEVICE__ __CONSTEXPR__ bool isinf(float __x) { return ::__isinff(__x); }
97 __DEVICE__ __CONSTEXPR__ bool isinf(double __x) { return ::__isinf(__x); }
98 __DEVICE__ __CONSTEXPR__ bool isfinite(float __x) { return ::__finitef(__x); }
99 __DEVICE__ __CONSTEXPR__ bool isfinite(double __x) { return ::__finite(__x); }
100 __DEVICE__ __CONSTEXPR__ bool isnan(float __x) { return ::__isnanf(__x); }
101 __DEVICE__ __CONSTEXPR__ bool isnan(double __x) { return ::__isnan(__x); }
102 
103 #if defined(__OPENMP_AMDGCN__)
104 #pragma omp end declare variant
105 #endif // defined(__OPENMP_AMDGCN__)
106 
107 __DEVICE__ __CONSTEXPR__ bool isgreater(float __x, float __y) {
108   return __builtin_isgreater(__x, __y);
109 }
110 __DEVICE__ __CONSTEXPR__ bool isgreater(double __x, double __y) {
111   return __builtin_isgreater(__x, __y);
112 }
113 __DEVICE__ __CONSTEXPR__ bool isgreaterequal(float __x, float __y) {
114   return __builtin_isgreaterequal(__x, __y);
115 }
116 __DEVICE__ __CONSTEXPR__ bool isgreaterequal(double __x, double __y) {
117   return __builtin_isgreaterequal(__x, __y);
118 }
119 __DEVICE__ __CONSTEXPR__ bool isless(float __x, float __y) {
120   return __builtin_isless(__x, __y);
121 }
122 __DEVICE__ __CONSTEXPR__ bool isless(double __x, double __y) {
123   return __builtin_isless(__x, __y);
124 }
125 __DEVICE__ __CONSTEXPR__ bool islessequal(float __x, float __y) {
126   return __builtin_islessequal(__x, __y);
127 }
128 __DEVICE__ __CONSTEXPR__ bool islessequal(double __x, double __y) {
129   return __builtin_islessequal(__x, __y);
130 }
131 __DEVICE__ __CONSTEXPR__ bool islessgreater(float __x, float __y) {
132   return __builtin_islessgreater(__x, __y);
133 }
134 __DEVICE__ __CONSTEXPR__ bool islessgreater(double __x, double __y) {
135   return __builtin_islessgreater(__x, __y);
136 }
137 __DEVICE__ __CONSTEXPR__ bool isnormal(float __x) {
138   return __builtin_isnormal(__x);
139 }
140 __DEVICE__ __CONSTEXPR__ bool isnormal(double __x) {
141   return __builtin_isnormal(__x);
142 }
143 __DEVICE__ __CONSTEXPR__ bool isunordered(float __x, float __y) {
144   return __builtin_isunordered(__x, __y);
145 }
146 __DEVICE__ __CONSTEXPR__ bool isunordered(double __x, double __y) {
147   return __builtin_isunordered(__x, __y);
148 }
149 __DEVICE__ __CONSTEXPR__ float modf(float __x, float *__iptr) {
150   return ::modff(__x, __iptr);
151 }
152 __DEVICE__ __CONSTEXPR__ float pow(float __base, int __iexp) {
153   return ::powif(__base, __iexp);
154 }
155 __DEVICE__ __CONSTEXPR__ double pow(double __base, int __iexp) {
156   return ::powi(__base, __iexp);
157 }
158 __DEVICE__ __CONSTEXPR__ float remquo(float __x, float __y, int *__quo) {
159   return ::remquof(__x, __y, __quo);
160 }
161 __DEVICE__ __CONSTEXPR__ float scalbln(float __x, long int __n) {
162   return ::scalblnf(__x, __n);
163 }
164 __DEVICE__ __CONSTEXPR__ bool signbit(float __x) { return ::__signbitf(__x); }
165 __DEVICE__ __CONSTEXPR__ bool signbit(double __x) { return ::__signbit(__x); }
166 
167 // Notably missing above is nexttoward.  We omit it because
168 // ocml doesn't provide an implementation, and we don't want to be in the
169 // business of implementing tricky libm functions in this header.
170 
171 // Other functions.
172 __DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y,
173                                       _Float16 __z) {
174   return __builtin_fmaf16(__x, __y, __z);
175 }
176 __DEVICE__ __CONSTEXPR__ _Float16 pow(_Float16 __base, int __iexp) {
177   return __ocml_pown_f16(__base, __iexp);
178 }
179 
180 #ifndef __OPENMP_AMDGCN__
181 // BEGIN DEF_FUN and HIP_OVERLOAD
182 
183 // BEGIN DEF_FUN
184 
185 #pragma push_macro("__DEF_FUN1")
186 #pragma push_macro("__DEF_FUN2")
187 #pragma push_macro("__DEF_FUN2_FI")
188 
189 // Define cmath functions with float argument and returns __retty.
190 #define __DEF_FUN1(__retty, __func)                                            \
191   __DEVICE__ __CONSTEXPR__ __retty __func(float __x) { return __func##f(__x); }
192 
193 // Define cmath functions with two float arguments and returns __retty.
194 #define __DEF_FUN2(__retty, __func)                                            \
195   __DEVICE__ __CONSTEXPR__ __retty __func(float __x, float __y) {              \
196     return __func##f(__x, __y);                                                \
197   }
198 
199 // Define cmath functions with a float and an int argument and returns __retty.
200 #define __DEF_FUN2_FI(__retty, __func)                                         \
201   __DEVICE__ __CONSTEXPR__ __retty __func(float __x, int __y) {                \
202     return __func##f(__x, __y);                                                \
203   }
204 
205 __DEF_FUN1(float, acos)
206 __DEF_FUN1(float, acosh)
207 __DEF_FUN1(float, asin)
208 __DEF_FUN1(float, asinh)
209 __DEF_FUN1(float, atan)
210 __DEF_FUN2(float, atan2)
211 __DEF_FUN1(float, atanh)
212 __DEF_FUN1(float, cbrt)
213 __DEF_FUN1(float, ceil)
214 __DEF_FUN2(float, copysign)
215 __DEF_FUN1(float, cos)
216 __DEF_FUN1(float, cosh)
217 __DEF_FUN1(float, erf)
218 __DEF_FUN1(float, erfc)
219 __DEF_FUN1(float, exp)
220 __DEF_FUN1(float, exp2)
221 __DEF_FUN1(float, expm1)
222 __DEF_FUN1(float, fabs)
223 __DEF_FUN2(float, fdim)
224 __DEF_FUN1(float, floor)
225 __DEF_FUN2(float, fmax)
226 __DEF_FUN2(float, fmin)
227 __DEF_FUN2(float, fmod)
228 __DEF_FUN2(float, hypot)
229 __DEF_FUN1(int, ilogb)
230 __DEF_FUN2_FI(float, ldexp)
231 __DEF_FUN1(float, lgamma)
232 __DEF_FUN1(float, log)
233 __DEF_FUN1(float, log10)
234 __DEF_FUN1(float, log1p)
235 __DEF_FUN1(float, log2)
236 __DEF_FUN1(float, logb)
237 __DEF_FUN1(long long, llrint)
238 __DEF_FUN1(long long, llround)
239 __DEF_FUN1(long, lrint)
240 __DEF_FUN1(long, lround)
241 __DEF_FUN1(float, nearbyint)
242 __DEF_FUN2(float, nextafter)
243 __DEF_FUN2(float, pow)
244 __DEF_FUN2(float, remainder)
245 __DEF_FUN1(float, rint)
246 __DEF_FUN1(float, round)
247 __DEF_FUN2_FI(float, scalbn)
248 __DEF_FUN1(float, sin)
249 __DEF_FUN1(float, sinh)
250 __DEF_FUN1(float, sqrt)
251 __DEF_FUN1(float, tan)
252 __DEF_FUN1(float, tanh)
253 __DEF_FUN1(float, tgamma)
254 __DEF_FUN1(float, trunc)
255 
256 #pragma pop_macro("__DEF_FUN1")
257 #pragma pop_macro("__DEF_FUN2")
258 #pragma pop_macro("__DEF_FUN2_FI")
259 
260 // END DEF_FUN
261 
262 // BEGIN HIP_OVERLOAD
263 
264 #pragma push_macro("__HIP_OVERLOAD1")
265 #pragma push_macro("__HIP_OVERLOAD2")
266 
267 // __hip_enable_if::type is a type function which returns __T if __B is true.
268 template <bool __B, class __T = void> struct __hip_enable_if {};
269 
270 template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
271 
272 namespace __hip {
273 template <class _Tp> struct is_integral {
274   enum { value = 0 };
275 };
276 template <> struct is_integral<bool> {
277   enum { value = 1 };
278 };
279 template <> struct is_integral<char> {
280   enum { value = 1 };
281 };
282 template <> struct is_integral<signed char> {
283   enum { value = 1 };
284 };
285 template <> struct is_integral<unsigned char> {
286   enum { value = 1 };
287 };
288 template <> struct is_integral<wchar_t> {
289   enum { value = 1 };
290 };
291 template <> struct is_integral<short> {
292   enum { value = 1 };
293 };
294 template <> struct is_integral<unsigned short> {
295   enum { value = 1 };
296 };
297 template <> struct is_integral<int> {
298   enum { value = 1 };
299 };
300 template <> struct is_integral<unsigned int> {
301   enum { value = 1 };
302 };
303 template <> struct is_integral<long> {
304   enum { value = 1 };
305 };
306 template <> struct is_integral<unsigned long> {
307   enum { value = 1 };
308 };
309 template <> struct is_integral<long long> {
310   enum { value = 1 };
311 };
312 template <> struct is_integral<unsigned long long> {
313   enum { value = 1 };
314 };
315 
316 // ToDo: specializes is_arithmetic<_Float16>
317 template <class _Tp> struct is_arithmetic {
318   enum { value = 0 };
319 };
320 template <> struct is_arithmetic<bool> {
321   enum { value = 1 };
322 };
323 template <> struct is_arithmetic<char> {
324   enum { value = 1 };
325 };
326 template <> struct is_arithmetic<signed char> {
327   enum { value = 1 };
328 };
329 template <> struct is_arithmetic<unsigned char> {
330   enum { value = 1 };
331 };
332 template <> struct is_arithmetic<wchar_t> {
333   enum { value = 1 };
334 };
335 template <> struct is_arithmetic<short> {
336   enum { value = 1 };
337 };
338 template <> struct is_arithmetic<unsigned short> {
339   enum { value = 1 };
340 };
341 template <> struct is_arithmetic<int> {
342   enum { value = 1 };
343 };
344 template <> struct is_arithmetic<unsigned int> {
345   enum { value = 1 };
346 };
347 template <> struct is_arithmetic<long> {
348   enum { value = 1 };
349 };
350 template <> struct is_arithmetic<unsigned long> {
351   enum { value = 1 };
352 };
353 template <> struct is_arithmetic<long long> {
354   enum { value = 1 };
355 };
356 template <> struct is_arithmetic<unsigned long long> {
357   enum { value = 1 };
358 };
359 template <> struct is_arithmetic<float> {
360   enum { value = 1 };
361 };
362 template <> struct is_arithmetic<double> {
363   enum { value = 1 };
364 };
365 
366 struct true_type {
367   static const __constant__ bool value = true;
368 };
369 struct false_type {
370   static const __constant__ bool value = false;
371 };
372 
373 template <typename __T, typename __U> struct is_same : public false_type {};
374 template <typename __T> struct is_same<__T, __T> : public true_type {};
375 
376 template <typename __T> struct add_rvalue_reference { typedef __T &&type; };
377 
378 template <typename __T> typename add_rvalue_reference<__T>::type declval();
379 
380 // decltype is only available in C++11 and above.
381 #if __cplusplus >= 201103L
382 // __hip_promote
383 template <class _Tp> struct __numeric_type {
384   static void __test(...);
385   static _Float16 __test(_Float16);
386   static float __test(float);
387   static double __test(char);
388   static double __test(int);
389   static double __test(unsigned);
390   static double __test(long);
391   static double __test(unsigned long);
392   static double __test(long long);
393   static double __test(unsigned long long);
394   static double __test(double);
395   // No support for long double, use double instead.
396   static double __test(long double);
397 
398   template <typename _U>
399   static auto __test_impl(int) -> decltype(__test(declval<_U>()));
400 
401   template <typename _U> static void __test_impl(...);
402 
403   typedef decltype(__test_impl<_Tp>(0)) type;
404   static const bool value = !is_same<type, void>::value;
405 };
406 
407 template <> struct __numeric_type<void> { static const bool value = true; };
408 
409 template <class _A1, class _A2 = void, class _A3 = void,
410           bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
411               &&__numeric_type<_A3>::value>
412 class __promote_imp {
413 public:
414   static const bool value = false;
415 };
416 
417 template <class _A1, class _A2, class _A3>
418 class __promote_imp<_A1, _A2, _A3, true> {
419 private:
420   typedef typename __promote_imp<_A1>::type __type1;
421   typedef typename __promote_imp<_A2>::type __type2;
422   typedef typename __promote_imp<_A3>::type __type3;
423 
424 public:
425   typedef decltype(__type1() + __type2() + __type3()) type;
426   static const bool value = true;
427 };
428 
429 template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
430 private:
431   typedef typename __promote_imp<_A1>::type __type1;
432   typedef typename __promote_imp<_A2>::type __type2;
433 
434 public:
435   typedef decltype(__type1() + __type2()) type;
436   static const bool value = true;
437 };
438 
439 template <class _A1> class __promote_imp<_A1, void, void, true> {
440 public:
441   typedef typename __numeric_type<_A1>::type type;
442   static const bool value = true;
443 };
444 
445 template <class _A1, class _A2 = void, class _A3 = void>
446 class __promote : public __promote_imp<_A1, _A2, _A3> {};
447 #endif //__cplusplus >= 201103L
448 } // namespace __hip
449 
450 // __HIP_OVERLOAD1 is used to resolve function calls with integer argument to
451 // avoid compilation error due to ambibuity. e.g. floor(5) is resolved with
452 // floor(double).
453 #define __HIP_OVERLOAD1(__retty, __fn)                                         \
454   template <typename __T>                                                      \
455   __DEVICE__ __CONSTEXPR__                                                     \
456       typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type  \
457       __fn(__T __x) {                                                          \
458     return ::__fn((double)__x);                                                \
459   }
460 
461 // __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double
462 // or integer argument to avoid compilation error due to ambibuity. e.g.
463 // max(5.0f, 6.0) is resolved with max(double, double).
464 #if __cplusplus >= 201103L
465 #define __HIP_OVERLOAD2(__retty, __fn)                                         \
466   template <typename __T1, typename __T2>                                      \
467   __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<                           \
468       __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value,  \
469       typename __hip::__promote<__T1, __T2>::type>::type                       \
470   __fn(__T1 __x, __T2 __y) {                                                   \
471     typedef typename __hip::__promote<__T1, __T2>::type __result_type;         \
472     return __fn((__result_type)__x, (__result_type)__y);                       \
473   }
474 #else
475 #define __HIP_OVERLOAD2(__retty, __fn)                                         \
476   template <typename __T1, typename __T2>                                      \
477   __DEVICE__ __CONSTEXPR__                                                     \
478       typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&            \
479                                    __hip::is_arithmetic<__T2>::value,          \
480                                __retty>::type                                  \
481       __fn(__T1 __x, __T2 __y) {                                               \
482     return __fn((double)__x, (double)__y);                                     \
483   }
484 #endif
485 
486 __HIP_OVERLOAD1(double, acos)
487 __HIP_OVERLOAD1(double, acosh)
488 __HIP_OVERLOAD1(double, asin)
489 __HIP_OVERLOAD1(double, asinh)
490 __HIP_OVERLOAD1(double, atan)
491 __HIP_OVERLOAD2(double, atan2)
492 __HIP_OVERLOAD1(double, atanh)
493 __HIP_OVERLOAD1(double, cbrt)
494 __HIP_OVERLOAD1(double, ceil)
495 __HIP_OVERLOAD2(double, copysign)
496 __HIP_OVERLOAD1(double, cos)
497 __HIP_OVERLOAD1(double, cosh)
498 __HIP_OVERLOAD1(double, erf)
499 __HIP_OVERLOAD1(double, erfc)
500 __HIP_OVERLOAD1(double, exp)
501 __HIP_OVERLOAD1(double, exp2)
502 __HIP_OVERLOAD1(double, expm1)
503 __HIP_OVERLOAD1(double, fabs)
504 __HIP_OVERLOAD2(double, fdim)
505 __HIP_OVERLOAD1(double, floor)
506 __HIP_OVERLOAD2(double, fmax)
507 __HIP_OVERLOAD2(double, fmin)
508 __HIP_OVERLOAD2(double, fmod)
509 #if !defined(__HIPCC_RTC__)
510 __HIP_OVERLOAD1(int, fpclassify)
511 #endif // !defined(__HIPCC_RTC__)
512 __HIP_OVERLOAD2(double, hypot)
513 __HIP_OVERLOAD1(int, ilogb)
514 __HIP_OVERLOAD1(bool, isfinite)
515 __HIP_OVERLOAD2(bool, isgreater)
516 __HIP_OVERLOAD2(bool, isgreaterequal)
517 __HIP_OVERLOAD1(bool, isinf)
518 __HIP_OVERLOAD2(bool, isless)
519 __HIP_OVERLOAD2(bool, islessequal)
520 __HIP_OVERLOAD2(bool, islessgreater)
521 __HIP_OVERLOAD1(bool, isnan)
522 __HIP_OVERLOAD1(bool, isnormal)
523 __HIP_OVERLOAD2(bool, isunordered)
524 __HIP_OVERLOAD1(double, lgamma)
525 __HIP_OVERLOAD1(double, log)
526 __HIP_OVERLOAD1(double, log10)
527 __HIP_OVERLOAD1(double, log1p)
528 __HIP_OVERLOAD1(double, log2)
529 __HIP_OVERLOAD1(double, logb)
530 __HIP_OVERLOAD1(long long, llrint)
531 __HIP_OVERLOAD1(long long, llround)
532 __HIP_OVERLOAD1(long, lrint)
533 __HIP_OVERLOAD1(long, lround)
534 __HIP_OVERLOAD1(double, nearbyint)
535 __HIP_OVERLOAD2(double, nextafter)
536 __HIP_OVERLOAD2(double, pow)
537 __HIP_OVERLOAD2(double, remainder)
538 __HIP_OVERLOAD1(double, rint)
539 __HIP_OVERLOAD1(double, round)
540 __HIP_OVERLOAD1(bool, signbit)
541 __HIP_OVERLOAD1(double, sin)
542 __HIP_OVERLOAD1(double, sinh)
543 __HIP_OVERLOAD1(double, sqrt)
544 __HIP_OVERLOAD1(double, tan)
545 __HIP_OVERLOAD1(double, tanh)
546 __HIP_OVERLOAD1(double, tgamma)
547 __HIP_OVERLOAD1(double, trunc)
548 
549 // Overload these but don't add them to std, they are not part of cmath.
550 __HIP_OVERLOAD2(double, max)
551 __HIP_OVERLOAD2(double, min)
552 
553 // Additional Overloads that don't quite match HIP_OVERLOAD.
554 #if __cplusplus >= 201103L
555 template <typename __T1, typename __T2, typename __T3>
556 __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<
557     __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&
558         __hip::is_arithmetic<__T3>::value,
559     typename __hip::__promote<__T1, __T2, __T3>::type>::type
560 fma(__T1 __x, __T2 __y, __T3 __z) {
561   typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;
562   return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z);
563 }
564 #else
565 template <typename __T1, typename __T2, typename __T3>
566 __DEVICE__ __CONSTEXPR__
567     typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
568                                  __hip::is_arithmetic<__T2>::value &&
569                                  __hip::is_arithmetic<__T3>::value,
570                              double>::type
571     fma(__T1 __x, __T2 __y, __T3 __z) {
572   return ::fma((double)__x, (double)__y, (double)__z);
573 }
574 #endif
575 
576 template <typename __T>
577 __DEVICE__ __CONSTEXPR__
578     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
579     frexp(__T __x, int *__exp) {
580   return ::frexp((double)__x, __exp);
581 }
582 
583 template <typename __T>
584 __DEVICE__ __CONSTEXPR__
585     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
586     ldexp(__T __x, int __exp) {
587   return ::ldexp((double)__x, __exp);
588 }
589 
590 template <typename __T>
591 __DEVICE__ __CONSTEXPR__
592     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
593     modf(__T __x, double *__exp) {
594   return ::modf((double)__x, __exp);
595 }
596 
597 #if __cplusplus >= 201103L
598 template <typename __T1, typename __T2>
599 __DEVICE__ __CONSTEXPR__
600     typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
601                                  __hip::is_arithmetic<__T2>::value,
602                              typename __hip::__promote<__T1, __T2>::type>::type
603     remquo(__T1 __x, __T2 __y, int *__quo) {
604   typedef typename __hip::__promote<__T1, __T2>::type __result_type;
605   return ::remquo((__result_type)__x, (__result_type)__y, __quo);
606 }
607 #else
608 template <typename __T1, typename __T2>
609 __DEVICE__ __CONSTEXPR__
610     typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
611                                  __hip::is_arithmetic<__T2>::value,
612                              double>::type
613     remquo(__T1 __x, __T2 __y, int *__quo) {
614   return ::remquo((double)__x, (double)__y, __quo);
615 }
616 #endif
617 
618 template <typename __T>
619 __DEVICE__ __CONSTEXPR__
620     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
621     scalbln(__T __x, long int __exp) {
622   return ::scalbln((double)__x, __exp);
623 }
624 
625 template <typename __T>
626 __DEVICE__ __CONSTEXPR__
627     typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
628     scalbn(__T __x, int __exp) {
629   return ::scalbn((double)__x, __exp);
630 }
631 
632 #pragma pop_macro("__HIP_OVERLOAD1")
633 #pragma pop_macro("__HIP_OVERLOAD2")
634 
635 // END HIP_OVERLOAD
636 
637 // END DEF_FUN and HIP_OVERLOAD
638 
639 #endif // ifndef __OPENMP_AMDGCN__
640 #endif // defined(__cplusplus)
641 
642 #ifndef __OPENMP_AMDGCN__
643 // Define these overloads inside the namespace our standard library uses.
644 #if !defined(__HIPCC_RTC__)
645 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
646 _LIBCPP_BEGIN_NAMESPACE_STD
647 #else
648 namespace std {
649 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
650 _GLIBCXX_BEGIN_NAMESPACE_VERSION
651 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
652 #endif // _LIBCPP_BEGIN_NAMESPACE_STD
653 
654 // Pull the new overloads we defined above into namespace std.
655 // using ::abs; - This may be considered for C++.
656 using ::acos;
657 using ::acosh;
658 using ::asin;
659 using ::asinh;
660 using ::atan;
661 using ::atan2;
662 using ::atanh;
663 using ::cbrt;
664 using ::ceil;
665 using ::copysign;
666 using ::cos;
667 using ::cosh;
668 using ::erf;
669 using ::erfc;
670 using ::exp;
671 using ::exp2;
672 using ::expm1;
673 using ::fabs;
674 using ::fdim;
675 using ::floor;
676 using ::fma;
677 using ::fmax;
678 using ::fmin;
679 using ::fmod;
680 using ::fpclassify;
681 using ::frexp;
682 using ::hypot;
683 using ::ilogb;
684 using ::isfinite;
685 using ::isgreater;
686 using ::isgreaterequal;
687 using ::isless;
688 using ::islessequal;
689 using ::islessgreater;
690 using ::isnormal;
691 using ::isunordered;
692 using ::ldexp;
693 using ::lgamma;
694 using ::llrint;
695 using ::llround;
696 using ::log;
697 using ::log10;
698 using ::log1p;
699 using ::log2;
700 using ::logb;
701 using ::lrint;
702 using ::lround;
703 using ::modf;
704 // using ::nan; - This may be considered for C++.
705 // using ::nanf; - This may be considered for C++.
706 // using ::nanl; - This is not yet defined.
707 using ::nearbyint;
708 using ::nextafter;
709 // using ::nexttoward; - Omit this since we do not have a definition.
710 using ::pow;
711 using ::remainder;
712 using ::remquo;
713 using ::rint;
714 using ::round;
715 using ::scalbln;
716 using ::scalbn;
717 using ::signbit;
718 using ::sin;
719 using ::sinh;
720 using ::sqrt;
721 using ::tan;
722 using ::tanh;
723 using ::tgamma;
724 using ::trunc;
725 
726 // Well this is fun: We need to pull these symbols in for libc++, but we can't
727 // pull them in with libstdc++, because its ::isinf and ::isnan are different
728 // than its std::isinf and std::isnan.
729 #ifndef __GLIBCXX__
730 using ::isinf;
731 using ::isnan;
732 #endif
733 
734 // Finally, pull the "foobarf" functions that HIP defines into std.
735 using ::acosf;
736 using ::acoshf;
737 using ::asinf;
738 using ::asinhf;
739 using ::atan2f;
740 using ::atanf;
741 using ::atanhf;
742 using ::cbrtf;
743 using ::ceilf;
744 using ::copysignf;
745 using ::cosf;
746 using ::coshf;
747 using ::erfcf;
748 using ::erff;
749 using ::exp2f;
750 using ::expf;
751 using ::expm1f;
752 using ::fabsf;
753 using ::fdimf;
754 using ::floorf;
755 using ::fmaf;
756 using ::fmaxf;
757 using ::fminf;
758 using ::fmodf;
759 using ::frexpf;
760 using ::hypotf;
761 using ::ilogbf;
762 using ::ldexpf;
763 using ::lgammaf;
764 using ::llrintf;
765 using ::llroundf;
766 using ::log10f;
767 using ::log1pf;
768 using ::log2f;
769 using ::logbf;
770 using ::logf;
771 using ::lrintf;
772 using ::lroundf;
773 using ::modff;
774 using ::nearbyintf;
775 using ::nextafterf;
776 // using ::nexttowardf; - Omit this since we do not have a definition.
777 using ::powf;
778 using ::remainderf;
779 using ::remquof;
780 using ::rintf;
781 using ::roundf;
782 using ::scalblnf;
783 using ::scalbnf;
784 using ::sinf;
785 using ::sinhf;
786 using ::sqrtf;
787 using ::tanf;
788 using ::tanhf;
789 using ::tgammaf;
790 using ::truncf;
791 
792 #ifdef _LIBCPP_END_NAMESPACE_STD
793 _LIBCPP_END_NAMESPACE_STD
794 #else
795 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
796 _GLIBCXX_END_NAMESPACE_VERSION
797 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
798 } // namespace std
799 #endif // _LIBCPP_END_NAMESPACE_STD
800 #endif // !defined(__HIPCC_RTC__)
801 
802 // Define device-side math functions from <ymath.h> on MSVC.
803 #if !defined(__HIPCC_RTC__)
804 #if defined(_MSC_VER)
805 
806 // Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers.
807 // But, from VS2019, it's only included in `<complex>`. Need to include
808 // `<ymath.h>` here to ensure C functions declared there won't be markded as
809 // `__host__` and `__device__` through `<complex>` wrapper.
810 #include <ymath.h>
811 
812 #if defined(__cplusplus)
813 extern "C" {
814 #endif // defined(__cplusplus)
815 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Cosh(double x,
816                                                                     double y) {
817   return cosh(x) * y;
818 }
819 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FCosh(float x,
820                                                                     float y) {
821   return coshf(x) * y;
822 }
823 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _Dtest(double *p) {
824   return fpclassify(*p);
825 }
826 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _FDtest(float *p) {
827   return fpclassify(*p);
828 }
829 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Sinh(double x,
830                                                                     double y) {
831   return sinh(x) * y;
832 }
833 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FSinh(float x,
834                                                                     float y) {
835   return sinhf(x) * y;
836 }
837 #if defined(__cplusplus)
838 }
839 #endif // defined(__cplusplus)
840 #endif // defined(_MSC_VER)
841 #endif // !defined(__HIPCC_RTC__)
842 #endif // ifndef __OPENMP_AMDGCN__
843 
844 #pragma pop_macro("__DEVICE__")
845 #pragma pop_macro("__CONSTEXPR__")
846 
847 #endif // __CLANG_HIP_CMATH_H__
848