1 /*===---- __clang_hip_cmath.h - HIP cmath decls -----------------------------===
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
7 *===-----------------------------------------------------------------------===
10 #ifndef __CLANG_HIP_CMATH_H__
11 #define __CLANG_HIP_CMATH_H__
13 #if !defined(__HIP__) && !defined(__OPENMP_AMDGCN__)
14 #error "This file is for HIP and OpenMP AMDGCN device compilation only."
17 #if !defined(__HIPCC_RTC__)
18 #if defined(__cplusplus)
20 #include <type_traits>
25 #endif // !defined(__HIPCC_RTC__)
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
33 #define __DEVICE__ static __device__ inline __attribute__((always_inline))
35 #endif // __OPENMP_AMDGCN__
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); }
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);
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,
58 __DEVICE__ __CONSTEXPR__ int fpclassify(double __x) {
59 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
62 #endif // !defined(__HIPCC_RTC__)
64 __DEVICE__ __CONSTEXPR__ float frexp(float __arg, int *__exp) {
65 return ::frexpf(__arg, __exp);
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)})
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)})
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); }
93 #pragma omp end declare variant
94 #endif // defined(__OPENMP_AMDGCN__)
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); }
103 #if defined(__OPENMP_AMDGCN__)
104 #pragma omp end declare variant
105 #endif // defined(__OPENMP_AMDGCN__)
107 __DEVICE__ __CONSTEXPR__ bool isgreater(float __x, float __y) {
108 return __builtin_isgreater(__x, __y);
110 __DEVICE__ __CONSTEXPR__ bool isgreater(double __x, double __y) {
111 return __builtin_isgreater(__x, __y);
113 __DEVICE__ __CONSTEXPR__ bool isgreaterequal(float __x, float __y) {
114 return __builtin_isgreaterequal(__x, __y);
116 __DEVICE__ __CONSTEXPR__ bool isgreaterequal(double __x, double __y) {
117 return __builtin_isgreaterequal(__x, __y);
119 __DEVICE__ __CONSTEXPR__ bool isless(float __x, float __y) {
120 return __builtin_isless(__x, __y);
122 __DEVICE__ __CONSTEXPR__ bool isless(double __x, double __y) {
123 return __builtin_isless(__x, __y);
125 __DEVICE__ __CONSTEXPR__ bool islessequal(float __x, float __y) {
126 return __builtin_islessequal(__x, __y);
128 __DEVICE__ __CONSTEXPR__ bool islessequal(double __x, double __y) {
129 return __builtin_islessequal(__x, __y);
131 __DEVICE__ __CONSTEXPR__ bool islessgreater(float __x, float __y) {
132 return __builtin_islessgreater(__x, __y);
134 __DEVICE__ __CONSTEXPR__ bool islessgreater(double __x, double __y) {
135 return __builtin_islessgreater(__x, __y);
137 __DEVICE__ __CONSTEXPR__ bool isnormal(float __x) {
138 return __builtin_isnormal(__x);
140 __DEVICE__ __CONSTEXPR__ bool isnormal(double __x) {
141 return __builtin_isnormal(__x);
143 __DEVICE__ __CONSTEXPR__ bool isunordered(float __x, float __y) {
144 return __builtin_isunordered(__x, __y);
146 __DEVICE__ __CONSTEXPR__ bool isunordered(double __x, double __y) {
147 return __builtin_isunordered(__x, __y);
149 __DEVICE__ __CONSTEXPR__ float modf(float __x, float *__iptr) {
150 return ::modff(__x, __iptr);
152 __DEVICE__ __CONSTEXPR__ float pow(float __base, int __iexp) {
153 return ::powif(__base, __iexp);
155 __DEVICE__ __CONSTEXPR__ double pow(double __base, int __iexp) {
156 return ::powi(__base, __iexp);
158 __DEVICE__ __CONSTEXPR__ float remquo(float __x, float __y, int *__quo) {
159 return ::remquof(__x, __y, __quo);
161 __DEVICE__ __CONSTEXPR__ float scalbln(float __x, long int __n) {
162 return ::scalblnf(__x, __n);
164 __DEVICE__ __CONSTEXPR__ bool signbit(float __x) { return ::__signbitf(__x); }
165 __DEVICE__ __CONSTEXPR__ bool signbit(double __x) { return ::__signbit(__x); }
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.
172 __DEVICE__ __CONSTEXPR__ _Float16 fma(_Float16 __x, _Float16 __y,
174 return __ocml_fma_f16(__x, __y, __z);
176 __DEVICE__ __CONSTEXPR__ _Float16 pow(_Float16 __base, int __iexp) {
177 return __ocml_pown_f16(__base, __iexp);
180 #ifndef __OPENMP_AMDGCN__
181 // BEGIN DEF_FUN and HIP_OVERLOAD
185 #pragma push_macro("__DEF_FUN1")
186 #pragma push_macro("__DEF_FUN2")
187 #pragma push_macro("__DEF_FUN2_FI")
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); }
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); \
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); \
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)
256 #pragma pop_macro("__DEF_FUN1")
257 #pragma pop_macro("__DEF_FUN2")
258 #pragma pop_macro("__DEF_FUN2_FI")
262 // BEGIN HIP_OVERLOAD
264 #pragma push_macro("__HIP_OVERLOAD1")
265 #pragma push_macro("__HIP_OVERLOAD2")
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 {};
270 template <class __T> struct __hip_enable_if<true, __T> { typedef __T type; };
273 template <class _Tp> struct is_integral {
276 template <> struct is_integral<bool> {
279 template <> struct is_integral<char> {
282 template <> struct is_integral<signed char> {
285 template <> struct is_integral<unsigned char> {
288 template <> struct is_integral<wchar_t> {
291 template <> struct is_integral<short> {
294 template <> struct is_integral<unsigned short> {
297 template <> struct is_integral<int> {
300 template <> struct is_integral<unsigned int> {
303 template <> struct is_integral<long> {
306 template <> struct is_integral<unsigned long> {
309 template <> struct is_integral<long long> {
312 template <> struct is_integral<unsigned long long> {
316 // ToDo: specializes is_arithmetic<_Float16>
317 template <class _Tp> struct is_arithmetic {
320 template <> struct is_arithmetic<bool> {
323 template <> struct is_arithmetic<char> {
326 template <> struct is_arithmetic<signed char> {
329 template <> struct is_arithmetic<unsigned char> {
332 template <> struct is_arithmetic<wchar_t> {
335 template <> struct is_arithmetic<short> {
338 template <> struct is_arithmetic<unsigned short> {
341 template <> struct is_arithmetic<int> {
344 template <> struct is_arithmetic<unsigned int> {
347 template <> struct is_arithmetic<long> {
350 template <> struct is_arithmetic<unsigned long> {
353 template <> struct is_arithmetic<long long> {
356 template <> struct is_arithmetic<unsigned long long> {
359 template <> struct is_arithmetic<float> {
362 template <> struct is_arithmetic<double> {
367 static const __constant__ bool value = true;
370 static const __constant__ bool value = false;
373 template <typename __T, typename __U> struct is_same : public false_type {};
374 template <typename __T> struct is_same<__T, __T> : public true_type {};
376 template <typename __T> struct add_rvalue_reference { typedef __T &&type; };
378 template <typename __T> typename add_rvalue_reference<__T>::type declval();
380 // decltype is only available in C++11 and above.
381 #if __cplusplus >= 201103L
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);
398 typedef decltype(__test(declval<_Tp>())) type;
399 static const bool value = !is_same<type, void>::value;
402 template <> struct __numeric_type<void> { static const bool value = true; };
404 template <class _A1, class _A2 = void, class _A3 = void,
405 bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
406 &&__numeric_type<_A3>::value>
407 class __promote_imp {
409 static const bool value = false;
412 template <class _A1, class _A2, class _A3>
413 class __promote_imp<_A1, _A2, _A3, true> {
415 typedef typename __promote_imp<_A1>::type __type1;
416 typedef typename __promote_imp<_A2>::type __type2;
417 typedef typename __promote_imp<_A3>::type __type3;
420 typedef decltype(__type1() + __type2() + __type3()) type;
421 static const bool value = true;
424 template <class _A1, class _A2> class __promote_imp<_A1, _A2, void, true> {
426 typedef typename __promote_imp<_A1>::type __type1;
427 typedef typename __promote_imp<_A2>::type __type2;
430 typedef decltype(__type1() + __type2()) type;
431 static const bool value = true;
434 template <class _A1> class __promote_imp<_A1, void, void, true> {
436 typedef typename __numeric_type<_A1>::type type;
437 static const bool value = true;
440 template <class _A1, class _A2 = void, class _A3 = void>
441 class __promote : public __promote_imp<_A1, _A2, _A3> {};
442 #endif //__cplusplus >= 201103L
445 // __HIP_OVERLOAD1 is used to resolve function calls with integer argument to
446 // avoid compilation error due to ambibuity. e.g. floor(5) is resolved with
448 #define __HIP_OVERLOAD1(__retty, __fn) \
449 template <typename __T> \
450 __DEVICE__ __CONSTEXPR__ \
451 typename __hip_enable_if<__hip::is_integral<__T>::value, __retty>::type \
453 return ::__fn((double)__x); \
456 // __HIP_OVERLOAD2 is used to resolve function calls with mixed float/double
457 // or integer argument to avoid compilation error due to ambibuity. e.g.
458 // max(5.0f, 6.0) is resolved with max(double, double).
459 #if __cplusplus >= 201103L
460 #define __HIP_OVERLOAD2(__retty, __fn) \
461 template <typename __T1, typename __T2> \
462 __DEVICE__ __CONSTEXPR__ typename __hip_enable_if< \
463 __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value, \
464 typename __hip::__promote<__T1, __T2>::type>::type \
465 __fn(__T1 __x, __T2 __y) { \
466 typedef typename __hip::__promote<__T1, __T2>::type __result_type; \
467 return __fn((__result_type)__x, (__result_type)__y); \
470 #define __HIP_OVERLOAD2(__retty, __fn) \
471 template <typename __T1, typename __T2> \
472 __DEVICE__ __CONSTEXPR__ \
473 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \
474 __hip::is_arithmetic<__T2>::value, \
476 __fn(__T1 __x, __T2 __y) { \
477 return __fn((double)__x, (double)__y); \
481 __HIP_OVERLOAD1(double, acos)
482 __HIP_OVERLOAD1(double, acosh)
483 __HIP_OVERLOAD1(double, asin)
484 __HIP_OVERLOAD1(double, asinh)
485 __HIP_OVERLOAD1(double, atan)
486 __HIP_OVERLOAD2(double, atan2)
487 __HIP_OVERLOAD1(double, atanh)
488 __HIP_OVERLOAD1(double, cbrt)
489 __HIP_OVERLOAD1(double, ceil)
490 __HIP_OVERLOAD2(double, copysign)
491 __HIP_OVERLOAD1(double, cos)
492 __HIP_OVERLOAD1(double, cosh)
493 __HIP_OVERLOAD1(double, erf)
494 __HIP_OVERLOAD1(double, erfc)
495 __HIP_OVERLOAD1(double, exp)
496 __HIP_OVERLOAD1(double, exp2)
497 __HIP_OVERLOAD1(double, expm1)
498 __HIP_OVERLOAD1(double, fabs)
499 __HIP_OVERLOAD2(double, fdim)
500 __HIP_OVERLOAD1(double, floor)
501 __HIP_OVERLOAD2(double, fmax)
502 __HIP_OVERLOAD2(double, fmin)
503 __HIP_OVERLOAD2(double, fmod)
504 #if !defined(__HIPCC_RTC__)
505 __HIP_OVERLOAD1(int, fpclassify)
506 #endif // !defined(__HIPCC_RTC__)
507 __HIP_OVERLOAD2(double, hypot)
508 __HIP_OVERLOAD1(int, ilogb)
509 __HIP_OVERLOAD1(bool, isfinite)
510 __HIP_OVERLOAD2(bool, isgreater)
511 __HIP_OVERLOAD2(bool, isgreaterequal)
512 __HIP_OVERLOAD1(bool, isinf)
513 __HIP_OVERLOAD2(bool, isless)
514 __HIP_OVERLOAD2(bool, islessequal)
515 __HIP_OVERLOAD2(bool, islessgreater)
516 __HIP_OVERLOAD1(bool, isnan)
517 __HIP_OVERLOAD1(bool, isnormal)
518 __HIP_OVERLOAD2(bool, isunordered)
519 __HIP_OVERLOAD1(double, lgamma)
520 __HIP_OVERLOAD1(double, log)
521 __HIP_OVERLOAD1(double, log10)
522 __HIP_OVERLOAD1(double, log1p)
523 __HIP_OVERLOAD1(double, log2)
524 __HIP_OVERLOAD1(double, logb)
525 __HIP_OVERLOAD1(long long, llrint)
526 __HIP_OVERLOAD1(long long, llround)
527 __HIP_OVERLOAD1(long, lrint)
528 __HIP_OVERLOAD1(long, lround)
529 __HIP_OVERLOAD1(double, nearbyint)
530 __HIP_OVERLOAD2(double, nextafter)
531 __HIP_OVERLOAD2(double, pow)
532 __HIP_OVERLOAD2(double, remainder)
533 __HIP_OVERLOAD1(double, rint)
534 __HIP_OVERLOAD1(double, round)
535 __HIP_OVERLOAD1(bool, signbit)
536 __HIP_OVERLOAD1(double, sin)
537 __HIP_OVERLOAD1(double, sinh)
538 __HIP_OVERLOAD1(double, sqrt)
539 __HIP_OVERLOAD1(double, tan)
540 __HIP_OVERLOAD1(double, tanh)
541 __HIP_OVERLOAD1(double, tgamma)
542 __HIP_OVERLOAD1(double, trunc)
544 // Overload these but don't add them to std, they are not part of cmath.
545 __HIP_OVERLOAD2(double, max)
546 __HIP_OVERLOAD2(double, min)
548 // Additional Overloads that don't quite match HIP_OVERLOAD.
549 #if __cplusplus >= 201103L
550 template <typename __T1, typename __T2, typename __T3>
551 __DEVICE__ __CONSTEXPR__ typename __hip_enable_if<
552 __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&
553 __hip::is_arithmetic<__T3>::value,
554 typename __hip::__promote<__T1, __T2, __T3>::type>::type
555 fma(__T1 __x, __T2 __y, __T3 __z) {
556 typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;
557 return ::fma((__result_type)__x, (__result_type)__y, (__result_type)__z);
560 template <typename __T1, typename __T2, typename __T3>
561 __DEVICE__ __CONSTEXPR__
562 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
563 __hip::is_arithmetic<__T2>::value &&
564 __hip::is_arithmetic<__T3>::value,
566 fma(__T1 __x, __T2 __y, __T3 __z) {
567 return ::fma((double)__x, (double)__y, (double)__z);
571 template <typename __T>
572 __DEVICE__ __CONSTEXPR__
573 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
574 frexp(__T __x, int *__exp) {
575 return ::frexp((double)__x, __exp);
578 template <typename __T>
579 __DEVICE__ __CONSTEXPR__
580 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
581 ldexp(__T __x, int __exp) {
582 return ::ldexp((double)__x, __exp);
585 template <typename __T>
586 __DEVICE__ __CONSTEXPR__
587 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
588 modf(__T __x, double *__exp) {
589 return ::modf((double)__x, __exp);
592 #if __cplusplus >= 201103L
593 template <typename __T1, typename __T2>
594 __DEVICE__ __CONSTEXPR__
595 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
596 __hip::is_arithmetic<__T2>::value,
597 typename __hip::__promote<__T1, __T2>::type>::type
598 remquo(__T1 __x, __T2 __y, int *__quo) {
599 typedef typename __hip::__promote<__T1, __T2>::type __result_type;
600 return ::remquo((__result_type)__x, (__result_type)__y, __quo);
603 template <typename __T1, typename __T2>
604 __DEVICE__ __CONSTEXPR__
605 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
606 __hip::is_arithmetic<__T2>::value,
608 remquo(__T1 __x, __T2 __y, int *__quo) {
609 return ::remquo((double)__x, (double)__y, __quo);
613 template <typename __T>
614 __DEVICE__ __CONSTEXPR__
615 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
616 scalbln(__T __x, long int __exp) {
617 return ::scalbln((double)__x, __exp);
620 template <typename __T>
621 __DEVICE__ __CONSTEXPR__
622 typename __hip_enable_if<__hip::is_integral<__T>::value, double>::type
623 scalbn(__T __x, int __exp) {
624 return ::scalbn((double)__x, __exp);
627 #pragma pop_macro("__HIP_OVERLOAD1")
628 #pragma pop_macro("__HIP_OVERLOAD2")
632 // END DEF_FUN and HIP_OVERLOAD
634 #endif // ifndef __OPENMP_AMDGCN__
635 #endif // defined(__cplusplus)
637 #ifndef __OPENMP_AMDGCN__
638 // Define these overloads inside the namespace our standard library uses.
639 #if !defined(__HIPCC_RTC__)
640 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
641 _LIBCPP_BEGIN_NAMESPACE_STD
644 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
645 _GLIBCXX_BEGIN_NAMESPACE_VERSION
646 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
647 #endif // _LIBCPP_BEGIN_NAMESPACE_STD
649 // Pull the new overloads we defined above into namespace std.
650 // using ::abs; - This may be considered for C++.
681 using ::isgreaterequal;
684 using ::islessgreater;
699 // using ::nan; - This may be considered for C++.
700 // using ::nanf; - This may be considered for C++.
701 // using ::nanl; - This is not yet defined.
704 // using ::nexttoward; - Omit this since we do not have a definition.
721 // Well this is fun: We need to pull these symbols in for libc++, but we can't
722 // pull them in with libstdc++, because its ::isinf and ::isnan are different
723 // than its std::isinf and std::isnan.
729 // Finally, pull the "foobarf" functions that HIP defines into std.
771 // using ::nexttowardf; - Omit this since we do not have a definition.
787 #ifdef _LIBCPP_END_NAMESPACE_STD
788 _LIBCPP_END_NAMESPACE_STD
790 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
791 _GLIBCXX_END_NAMESPACE_VERSION
792 #endif // _GLIBCXX_BEGIN_NAMESPACE_VERSION
794 #endif // _LIBCPP_END_NAMESPACE_STD
795 #endif // !defined(__HIPCC_RTC__)
797 // Define device-side math functions from <ymath.h> on MSVC.
798 #if !defined(__HIPCC_RTC__)
799 #if defined(_MSC_VER)
801 // Before VS2019, `<ymath.h>` is also included in `<limits>` and other headers.
802 // But, from VS2019, it's only included in `<complex>`. Need to include
803 // `<ymath.h>` here to ensure C functions declared there won't be markded as
804 // `__host__` and `__device__` through `<complex>` wrapper.
807 #if defined(__cplusplus)
809 #endif // defined(__cplusplus)
810 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Cosh(double x,
814 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FCosh(float x,
818 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _Dtest(double *p) {
819 return fpclassify(*p);
821 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) short _FDtest(float *p) {
822 return fpclassify(*p);
824 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) double _Sinh(double x,
828 __DEVICE__ __CONSTEXPR__ __attribute__((overloadable)) float _FSinh(float x,
832 #if defined(__cplusplus)
834 #endif // defined(__cplusplus)
835 #endif // defined(_MSC_VER)
836 #endif // !defined(__HIPCC_RTC__)
837 #endif // ifndef __OPENMP_AMDGCN__
839 #pragma pop_macro("__DEVICE__")
840 #pragma pop_macro("__CONSTEXPR__")
842 #endif // __CLANG_HIP_CMATH_H__