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> {
300template <>
struct is_integral<unsigned
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> {
344template <>
struct is_arithmetic<unsigned
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> {
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 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;
407template <>
struct __numeric_type<void> {
static const bool value =
true; };
409template <
class _A1,
class _A2 = void,
class _A3 = void,
410 bool = __numeric_type<_A1>::value &&__numeric_type<_A2>::value
411 &&__numeric_type<_A3>::value>
414 static const bool value =
false;
417template <
class _A1,
class _A2,
class _A3>
418class __promote_imp<_A1, _A2, _A3,
true> {
420 typedef typename __promote_imp<_A1>::type __type1;
421 typedef typename __promote_imp<_A2>::type __type2;
422 typedef typename __promote_imp<_A3>::type __type3;
425 typedef decltype(__type1() + __type2() + __type3()) type;
426 static const bool value =
true;
429template <
class _A1,
class _A2>
class __promote_imp<_A1, _A2, void,
true> {
431 typedef typename __promote_imp<_A1>::type __type1;
432 typedef typename __promote_imp<_A2>::type __type2;
435 typedef decltype(__type1() + __type2()) type;
436 static const bool value =
true;
439template <
class _A1>
class __promote_imp<_A1, void, void,
true> {
441 typedef typename __numeric_type<_A1>::type
type;
442 static const bool value =
true;
445template <
class _A1,
class _A2 =
void,
class _A3 =
void>
446class __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__ \
468 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \
469 __hip::is_arithmetic<__T2>::value, \
471 __fn(__T1 __x, __T2 __y) { \
472 typedef typename __hip::__promote<__T1, __T2>::type __arg_type; \
473 return __fn((__arg_type)__x, (__arg_type)__y); \
476#define __HIP_OVERLOAD2(__retty, __fn) \
477 template <typename __T1, typename __T2> \
478 __DEVICE__ __CONSTEXPR__ \
479 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value && \
480 __hip::is_arithmetic<__T2>::value, \
482 __fn(__T1 __x, __T2 __y) { \
483 return __fn((double)__x, (double)__y); \
487__HIP_OVERLOAD1(
double,
acos)
488__HIP_OVERLOAD1(
double,
acosh)
489__HIP_OVERLOAD1(
double,
asin)
490__HIP_OVERLOAD1(
double,
asinh)
491__HIP_OVERLOAD1(
double,
atan)
492__HIP_OVERLOAD2(
double,
atan2)
493__HIP_OVERLOAD1(
double,
atanh)
494__HIP_OVERLOAD1(
double,
cbrt)
495__HIP_OVERLOAD1(
double,
ceil)
497__HIP_OVERLOAD1(
double,
cos)
498__HIP_OVERLOAD1(
double,
cosh)
499__HIP_OVERLOAD1(
double,
erf)
500__HIP_OVERLOAD1(
double,
erfc)
501__HIP_OVERLOAD1(
double,
exp)
502__HIP_OVERLOAD1(
double,
exp2)
503__HIP_OVERLOAD1(
double,
expm1)
504__HIP_OVERLOAD1(
double,
fabs)
505__HIP_OVERLOAD2(
double,
fdim)
506__HIP_OVERLOAD1(
double,
floor)
507__HIP_OVERLOAD2(
double,
fmax)
508__HIP_OVERLOAD2(
double,
fmin)
509__HIP_OVERLOAD2(
double,
fmod)
510#if !defined(__HIPCC_RTC__)
513__HIP_OVERLOAD2(
double,
hypot)
514__HIP_OVERLOAD1(
int,
ilogb)
518__HIP_OVERLOAD1(
bool,
isinf)
519__HIP_OVERLOAD2(
bool,
isless)
522__HIP_OVERLOAD1(
bool,
isnan)
525__HIP_OVERLOAD1(
double,
lgamma)
526__HIP_OVERLOAD1(
double,
log)
527__HIP_OVERLOAD1(
double,
log10)
528__HIP_OVERLOAD1(
double,
log1p)
529__HIP_OVERLOAD1(
double,
log2)
530__HIP_OVERLOAD1(
double,
logb)
531__HIP_OVERLOAD1(
long long,
llrint)
532__HIP_OVERLOAD1(
long long,
llround)
533__HIP_OVERLOAD1(
long,
lrint)
534__HIP_OVERLOAD1(
long,
lround)
537__HIP_OVERLOAD2(
double,
pow)
539__HIP_OVERLOAD1(
double,
rint)
540__HIP_OVERLOAD1(
double,
round)
542__HIP_OVERLOAD1(
double,
sin)
543__HIP_OVERLOAD1(
double,
sinh)
544__HIP_OVERLOAD1(
double,
sqrt)
545__HIP_OVERLOAD1(
double,
tan)
546__HIP_OVERLOAD1(
double,
tanh)
547__HIP_OVERLOAD1(
double,
tgamma)
548__HIP_OVERLOAD1(
double,
trunc)
551__HIP_OVERLOAD2(
double,
max)
552__HIP_OVERLOAD2(
double,
min)
555#if __cplusplus >= 201103L
556template <
typename __T1,
typename __T2,
typename __T3>
558 __hip::is_arithmetic<__T1>::value && __hip::is_arithmetic<__T2>::value &&
559 __hip::is_arithmetic<__T3>::value,
560 typename __hip::__promote<__T1, __T2, __T3>::type>::type
561fma(__T1 __x, __T2
__y, __T3 __z) {
562 typedef typename __hip::__promote<__T1, __T2, __T3>::type __result_type;
563 return ::fma((__result_type)__x, (__result_type)
__y, (__result_type)__z);
566template <
typename __T1,
typename __T2,
typename __T3>
568 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
569 __hip::is_arithmetic<__T2>::value &&
570 __hip::is_arithmetic<__T3>::value,
572 fma(__T1 __x, __T2
__y, __T3 __z) {
573 return ::fma((
double)__x, (
double)
__y, (
double)__z);
577template <
typename __T>
579 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>::type
580 frexp(__T __x,
int *__exp) {
581 return ::frexp((
double)__x, __exp);
584template <
typename __T>
586 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>::type
587 ldexp(__T __x,
int __exp) {
588 return ::ldexp((
double)__x, __exp);
591template <
typename __T>
593 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>::type
594 modf(__T __x,
double *__exp) {
595 return ::modf((
double)__x, __exp);
598#if __cplusplus >= 201103L
599template <
typename __T1,
typename __T2>
601 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
602 __hip::is_arithmetic<__T2>::value,
603 typename __hip::__promote<__T1, __T2>::type>::type
605 typedef typename __hip::__promote<__T1, __T2>::type __result_type;
606 return ::remquo((__result_type)__x, (__result_type)
__y, __quo);
609template <
typename __T1,
typename __T2>
611 typename __hip_enable_if<__hip::is_arithmetic<__T1>::value &&
612 __hip::is_arithmetic<__T2>::value,
615 return ::remquo((
double)__x, (
double)
__y, __quo);
619template <
typename __T>
621 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>::type
622 scalbln(__T __x,
long int __exp) {
623 return ::scalbln((
double)__x, __exp);
626template <
typename __T>
628 typename __hip_enable_if<__hip::is_integral<__T>::value,
double>::type
629 scalbn(__T __x,
int __exp) {
630 return ::scalbn((
double)__x, __exp);
633#pragma pop_macro("__HIP_OVERLOAD1")
634#pragma pop_macro("__HIP_OVERLOAD2")
643#ifndef __OPENMP_AMDGCN__
645#if !defined(__HIPCC_RTC__)
646#ifdef _LIBCPP_BEGIN_NAMESPACE_STD
647_LIBCPP_BEGIN_NAMESPACE_STD
650#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
651_GLIBCXX_BEGIN_NAMESPACE_VERSION
793#ifdef _LIBCPP_END_NAMESPACE_STD
794_LIBCPP_END_NAMESPACE_STD
796#ifdef _GLIBCXX_BEGIN_NAMESPACE_VERSION
797_GLIBCXX_END_NAMESPACE_VERSION
804#if !defined(__HIPCC_RTC__)
813#if defined(__cplusplus)
838#if defined(__cplusplus)
845#pragma pop_macro("__DEVICE__")
846#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.
__DEVICE__ float fabsf(float __a)
__DEVICE__ float fmodf(float __a, float __b)
__DEVICE__ float remainderf(float __a, float __b)
__DEVICE__ bool isunordered(float __x, float __y)
__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__ float nextafterf(float __a, float __b)
__DEVICE__ bool isgreater(float __x, float __y)
__DEVICE__ float fmaxf(float __a, float __b)
__DEVICE__ bool islessgreater(float __x, float __y)
__DEVICE__ bool isnan(float __x)
__DEVICE__ int fpclassify(float __x)
__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__ bool isfinite(float __x)
__DEVICE__ float sinf(float __a)
__DEVICE__ float logf(float __a)
__DEVICE__ float erff(float __a)
__DEVICE__ bool signbit(float __x)
__DEVICE__ bool isinf(float __x)
__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 modf(float __x, float *__iptr)
__DEVICE__ float modff(float __a, float *__b)
__DEVICE__ bool islessequal(float __x, float __y)
__DEVICE__ float logbf(float __a)
__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__ bool isless(float __x, float __y)
__DEVICE__ float powf(float __a, float __b)
__DEVICE__ bool isnormal(float __x)
__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__ float tgammaf(float __a)
__DEVICE__ float log1pf(float __a)
__DEVICE__ float lgammaf(float __a)
__DEVICE__ long long llroundf(float __a)
__DEVICE__ bool isgreaterequal(float __x, float __y)
__DEVICE__ float hypotf(float __a, float __b)
__DEVICE__ float asinf(float __a)
#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)