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)
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__
49 return ::fmaf(__x,
__y, __z);
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,
65 return ::frexpf(__arg, __exp);
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);
150 return ::modff(__x, __iptr);
153 return ::powif(__base, __iexp);
156 return ::powi(__base, __iexp);
159 return ::remquof(__x,
__y, __quo);
162 return ::scalblnf(__x, __n);
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")
268template <
bool __B,
class __T =
void>
struct __hip_enable_if {};
270template <
class __T>
struct __hip_enable_if<
true, __T> {
typedef __T
type; };
273template <
class _Tp>
struct is_integral {
276template <>
struct is_integral<
bool> {
279template <>
struct is_integral<char> {
282template <>
struct is_integral<signed char> {
285template <>
struct is_integral<
unsigned char> {
288template <>
struct is_integral<
wchar_t> {
291template <>
struct is_integral<short> {
294template <>
struct is_integral<
unsigned short> {
297template <>
struct is_integral<
int> {
303template <>
struct is_integral<long> {
306template <>
struct is_integral<
unsigned long> {
309template <>
struct is_integral<long long> {
312template <>
struct is_integral<
unsigned long long> {
317template <
class _Tp>
struct is_arithmetic {
320template <>
struct is_arithmetic<
bool> {
323template <>
struct is_arithmetic<char> {
326template <>
struct is_arithmetic<signed char> {
329template <>
struct is_arithmetic<
unsigned char> {
332template <>
struct is_arithmetic<
wchar_t> {
335template <>
struct is_arithmetic<short> {
338template <>
struct is_arithmetic<
unsigned short> {
341template <>
struct is_arithmetic<
int> {
347template <>
struct is_arithmetic<long> {
350template <>
struct is_arithmetic<
unsigned long> {
353template <>
struct is_arithmetic<long long> {
356template <>
struct is_arithmetic<
unsigned long long> {
359template <>
struct is_arithmetic<
float> {
362template <>
struct is_arithmetic<
double> {
367 static const __constant__
bool value =
true;
370 static const __constant__
bool value =
false;
373template <
typename __T,
typename __U>
struct is_same :
public false_type {};
374template <
typename __T>
struct is_same<__T, __T> :
public true_type {};
376template <
typename __T>
struct add_rvalue_reference {
typedef __T &&
type; };
378template <
typename __T>
typename add_rvalue_reference<__T>::type declval();
381#if __cplusplus >= 201103L
383template <
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 typedef decltype(__test(declval<_Tp>())) type;
399 static const bool value = !is_same<type, void>::value;
402template <>
struct __numeric_type<void> {
static const bool value =
true; };
404template <
class _A1,
class _A2 = void,
class _A3 = void,
405 bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
406 &&__numeric_type<_A3>::value>
409 static const bool value =
false;
412template <
class _A1,
class _A2,
class _A3>
413class __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;
424template <
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;
434template <
class _A1>
class __promote_imp<_A1, void, void,
true> {
436 typedef typename __numeric_type<_A1>::type
type;
437 static const bool value =
true;
440template <
class _A1,
class _A2 =
void,
class _A3 =
void>
441class __promote :
public __promote_imp<_A1, _A2, _A3> {};
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); \
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)
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__)
507__HIP_OVERLOAD2(
double,
hypot)
508__HIP_OVERLOAD1(
int,
ilogb)
512__HIP_OVERLOAD1(
bool,
isinf)
513__HIP_OVERLOAD2(
bool,
isless)
516__HIP_OVERLOAD1(
bool,
isnan)
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)
531__HIP_OVERLOAD2(
double,
pow)
533__HIP_OVERLOAD1(
double,
rint)
534__HIP_OVERLOAD1(
double,
round)
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)
545__HIP_OVERLOAD2(
double,
max)
546__HIP_OVERLOAD2(
double,
min)
549#if __cplusplus >= 201103L
550template <
typename __T1,
typename __T2,
typename __T3>
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
555fma(__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);
560template <
typename __T1,
typename __T2,
typename __T3>
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);
571template <
typename __T>
573 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>::type
574 frexp(__T __x,
int *__exp) {
575 return ::frexp((
double)__x, __exp);
578template <
typename __T>
580 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>::type
581 ldexp(__T __x,
int __exp) {
582 return ::ldexp((
double)__x, __exp);
585template <
typename __T>
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
593template <
typename __T1,
typename __T2>
595 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
596 __hip::is_arithmetic<__T2>::value,
597 typename __hip::__promote<__T1, __T2>::type>::type
599 typedef typename __hip::__promote<__T1, __T2>::type __result_type;
600 return ::remquo((__result_type)__x, (__result_type)
__y, __quo);
603template <
typename __T1,
typename __T2>
605 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
606 __hip::is_arithmetic<__T2>::value,
609 return ::remquo((
double)__x, (
double)
__y, __quo);
613template <
typename __T>
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);
620template <
typename __T>
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")
637#ifndef __OPENMP_AMDGCN__
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
681using ::isgreaterequal;
684using ::islessgreater;
787#ifdef _LIBCPP_END_NAMESPACE_STD
788_LIBCPP_END_NAMESPACE_STD
790#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
791_GLIBCXX_END_NAMESPACE_VERSION
798#if !defined(__HIPCC_RTC__)
807#if defined(__cplusplus)
832#if defined(__cplusplus)
839#pragma pop_macro("__DEVICE__")
840#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 min(int __a, int __b)
__DEVICE__ float sinhf(float __a)
__DEVICE__ float coshf(float __a)
__DEVICE__ int max(int __a, int __b)
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
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)