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>
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))
38 #if defined(__cplusplus)
39 #if defined __OPENMP_AMDGCN__
51 #if !defined(__HIPCC_RTC__)
55 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
59 return __builtin_fpclassify(FP_NAN, FP_INFINITE, FP_NORMAL, FP_SUBNORMAL,
68 #if defined(__OPENMP_AMDGCN__)
75 #pragma omp begin declare variant match( \
76 implementation = {extension(disable_implicit_base)})
84 #pragma omp begin declare variant match(implementation = {vendor(llvm)})
93 #pragma omp end declare variant
103 #if defined(__OPENMP_AMDGCN__)
104 #pragma omp end declare variant
108 return __builtin_isgreater(__x,
__y);
111 return __builtin_isgreater(__x,
__y);
114 return __builtin_isgreaterequal(__x,
__y);
117 return __builtin_isgreaterequal(__x,
__y);
120 return __builtin_isless(__x,
__y);
123 return __builtin_isless(__x,
__y);
126 return __builtin_islessequal(__x,
__y);
129 return __builtin_islessequal(__x,
__y);
132 return __builtin_islessgreater(__x,
__y);
135 return __builtin_islessgreater(__x,
__y);
138 return __builtin_isnormal(__x);
141 return __builtin_isnormal(__x);
144 return __builtin_isunordered(__x,
__y);
147 return __builtin_isunordered(__x,
__y);
174 return __builtin_fmaf16(__x,
__y, __z);
177 return __ocml_pown_f16(__base, __iexp);
180 #ifndef __OPENMP_AMDGCN__
185 #pragma push_macro("__DEF_FUN1")
186 #pragma push_macro("__DEF_FUN2")
187 #pragma push_macro("__DEF_FUN2_FI")
190 #define __DEF_FUN1(__retty, __func) \
191 __DEVICE__ __CONSTEXPR__ __retty __func(float __x) { return __func##f(__x); }
194 #define __DEF_FUN2(__retty, __func) \
195 __DEVICE__ __CONSTEXPR__ __retty __func(float __x, float __y) { \
196 return __func##f(__x, __y); \
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)
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)
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)
239 __DEF_FUN1(
long,
lrint)
243 __DEF_FUN2(
float,
pow)
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)
254 __DEF_FUN1(
float,
trunc)
256 #pragma pop_macro("__DEF_FUN1")
257 #pragma pop_macro("__DEF_FUN2")
258 #pragma pop_macro("__DEF_FUN2_FI")
264 #pragma push_macro("__HIP_OVERLOAD1")
265 #pragma push_macro("__HIP_OVERLOAD2")
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> {
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> {
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> {
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> {
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; };
381 #if __cplusplus >= 201103L
383 template <
class _Tp>
struct __numeric_type {
384 static void __test(...);
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);
396 static double __test(
long double);
398 template <
typename _U>
399 static auto __test_impl(
int) -> decltype(__test(declval<_U>()));
401 template <
typename _U>
static void __test_impl(...);
403 typedef decltype(__test_impl<_Tp>(0))
type;
404 static const
bool value = !is_same<
type,
void>::value;
407 template <> struct __numeric_type<
void> {
static const bool value =
true; };
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 {
414 static const bool value =
false;
417 template <
class _A1,
class _A2,
class _A3>
418 class __promote_imp<_A1, _A2, _A3,
true> {
425 typedef decltype(__type1() + __type2() + __type3())
type;
426 static const
bool value =
true;
429 template <class _A1, class _A2> class __promote_imp<_A1, _A2,
void,
true> {
435 typedef decltype(__type1() + __type2())
type;
436 static const
bool value =
true;
439 template <class _A1> class __promote_imp<_A1,
void,
void,
true> {
442 static const bool value =
true;
445 template <
class _A1,
class _A2 =
void,
class _A3 =
void>
446 class __promote :
public __promote_imp<_A1, _A2, _A3> {};
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 \
458 return ::__fn((double)__x); \
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); \
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, \
481 __fn(__T1 __x, __T2 __y) { \
482 return __fn((double)__x, (double)__y); \
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)
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__)
512 __HIP_OVERLOAD2(
double,
hypot)
513 __HIP_OVERLOAD1(
int,
ilogb)
517 __HIP_OVERLOAD1(
bool,
isinf)
518 __HIP_OVERLOAD2(
bool,
isless)
521 __HIP_OVERLOAD1(
bool,
isnan)
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)
536 __HIP_OVERLOAD2(
double,
pow)
538 __HIP_OVERLOAD1(
double,
rint)
539 __HIP_OVERLOAD1(
double,
round)
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)
550 __HIP_OVERLOAD2(
double,
max)
551 __HIP_OVERLOAD2(
double,
min)
554 #if __cplusplus >= 201103L
555 template <
typename __T1,
typename __T2,
typename __T3>
557 __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&
558 __hip::is_arithmetic<__T3>::value,
560 fma(__T1 __x, __T2
__y, __T3 __z) {
562 return ::fma((__result_type)__x, (__result_type)
__y, (__result_type)__z);
565 template <
typename __T1,
typename __T2,
typename __T3>
567 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
568 __hip::is_arithmetic<__T2>::value &&
569 __hip::is_arithmetic<__T3>::value,
571 fma(__T1 __x, __T2
__y, __T3 __z) {
576 template <
typename __T>
578 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>
::type
579 frexp(__T __x,
int *__exp) {
583 template <
typename __T>
585 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>
::type
586 ldexp(__T __x,
int __exp) {
590 template <
typename __T>
592 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>
::type
593 modf(__T __x,
double *__exp) {
597 #if __cplusplus >= 201103L
598 template <
typename __T1,
typename __T2>
600 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
601 __hip::is_arithmetic<__T2>::value,
608 template <
typename __T1,
typename __T2>
610 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
611 __hip::is_arithmetic<__T2>::value,
618 template <
typename __T>
620 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>
::type
621 scalbln(__T __x,
long int __exp) {
625 template <
typename __T>
627 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>
::type
628 scalbn(__T __x,
int __exp) {
632 #pragma pop_macro("__HIP_OVERLOAD1")
633 #pragma pop_macro("__HIP_OVERLOAD2")
642 #ifndef __OPENMP_AMDGCN__
644 #if !defined(__HIPCC_RTC__)
645 #ifdef _LIBCPP_BEGIN_NAMESPACE_STD
646 _LIBCPP_BEGIN_NAMESPACE_STD
649 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
650 _GLIBCXX_BEGIN_NAMESPACE_VERSION
792 #ifdef _LIBCPP_END_NAMESPACE_STD
793 _LIBCPP_END_NAMESPACE_STD
795 #ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
796 _GLIBCXX_END_NAMESPACE_VERSION
803 #if !defined(__HIPCC_RTC__)
804 #if defined(_MSC_VER)
812 #if defined(__cplusplus)
837 #if defined(__cplusplus)
844 #pragma pop_macro("__DEVICE__")
845 #pragma pop_macro("__CONSTEXPR__")
__DEVICE__ bool isunordered(float __x, float __y)
Test if arguments are unordered.
__DEVICE__ bool isgreater(float __x, float __y)
Returns the component-wise compare of x > y.
__DEVICE__ bool islessgreater(float __x, float __y)
Returns the component-wise compare of (x < y) || (x > y) .
__DEVICE__ bool isnan(float __x)
Test for a NaN.
__DEVICE__ int fpclassify(float __x)
__DEVICE__ bool isfinite(float __x)
Test for finite value.
__DEVICE__ bool signbit(float __x)
Test for sign bit.
__DEVICE__ bool isinf(float __x)
Test for infinity value (+ve or -ve) .
__DEVICE__ float modf(float __x, float *__iptr)
__DEVICE__ bool islessequal(float __x, float __y)
Returns the component-wise compare of x <= y.
__DEVICE__ long long abs(long long __n)
__DEVICE__ bool isless(float __x, float __y)
Returns the component-wise compare of x < y.
__DEVICE__ bool isnormal(float __x)
Test for a normal value.
__DEVICE__ bool isgreaterequal(float __x, float __y)
Returns the component-wise compare of x >= y.
__DEVICE__ int __isinff(float __a)
__DEVICE__ int __isnan(double __a)
__DEVICE__ int __isinf(double __a)
__DEVICE__ int __signbitf(float __a)
__DEVICE__ int __finite(double __a)
__DEVICE__ int __finitef(float __a)
__DEVICE__ int __isnanf(float __a)
__DEVICE__ double powi(double __a, int __b)
__DEVICE__ float fabsf(float __a)
__DEVICE__ float fmodf(float __a, float __b)
__DEVICE__ float remainderf(float __a, float __b)
__DEVICE__ float exp2f(float __a)
__DEVICE__ float acosf(float __a)
__DEVICE__ float fmaf(float __a, float __b, float __c)
__DEVICE__ float cbrtf(float __a)
__DEVICE__ float remquof(float __a, float __b, int *__c)
__DEVICE__ float tanf(float __a)
__DEVICE__ long labs(long __a)
__DEVICE__ float nextafterf(float __a, float __b)
__DEVICE__ float fmaxf(float __a, float __b)
__DEVICE__ long long llabs(long long __a)
__DEVICE__ float fminf(float __a, float __b)
__DEVICE__ float log2f(float __a)
__DEVICE__ float copysignf(float __a, float __b)
__DEVICE__ float truncf(float __a)
__DEVICE__ float fdimf(float __a, float __b)
__DEVICE__ long lrintf(float __a)
__DEVICE__ long long llrintf(float __a)
__DEVICE__ float cosf(float __a)
__DEVICE__ float sinf(float __a)
__DEVICE__ float logf(float __a)
__DEVICE__ int min(int __a, int __b)
__DEVICE__ float erff(float __a)
__DEVICE__ float floorf(float __f)
__DEVICE__ float ceilf(float __a)
__DEVICE__ float nearbyintf(float __a)
__DEVICE__ float atanf(float __a)
__DEVICE__ float atanhf(float __a)
__DEVICE__ float tanhf(float __a)
__DEVICE__ float rintf(float __a)
__DEVICE__ float atan2f(float __a, float __b)
__DEVICE__ float sinhf(float __a)
__DEVICE__ float acoshf(float __a)
__DEVICE__ float log10f(float __a)
__DEVICE__ float ldexpf(float __a, int __b)
__DEVICE__ float modff(float __a, float *__b)
__DEVICE__ float logbf(float __a)
__DEVICE__ float powif(float __a, int __b)
__DEVICE__ float coshf(float __a)
__DEVICE__ float asinhf(float __a)
__DEVICE__ float roundf(float __a)
__DEVICE__ long lroundf(float __a)
__DEVICE__ float scalbnf(float __a, int __b)
__DEVICE__ float erfcf(float __a)
__DEVICE__ int ilogbf(float __a)
__DEVICE__ float powf(float __a, float __b)
__DEVICE__ float frexpf(float __a, int *__b)
__DEVICE__ float sqrtf(float __a)
__DEVICE__ float expf(float __a)
__DEVICE__ float expm1f(float __a)
__DEVICE__ float scalblnf(float __a, long __b)
__DEVICE__ int max(int __a, int __b)
__DEVICE__ float tgammaf(float __a)
__DEVICE__ float log1pf(float __a)
__DEVICE__ float lgammaf(float __a)
__DEVICE__ long long llroundf(float __a)
__DEVICE__ float hypotf(float __a, float __b)
__DEVICE__ float asinf(float __a)
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
__DEVICE__ __RETURN_TYPE __signbit(double __x)
static __inline__ uint32_t uint32_t __y
const internal::VariadicAllOfMatcher< Type > type
Matches Types in the clang AST.
#define scalbln(__x, __y)
#define copysign(__x, __y)
#define remquo(__x, __y, __z)
#define nextafter(__x, __y)
#define remainder(__x, __y)
#define fma(__x, __y, __z)