9 #ifndef __CLANG_HIP_MATH_H__ 10 #define __CLANG_HIP_MATH_H__ 13 #error "This file is for HIP and OpenMP AMDGCN device compilation only." 16 #if defined(__cplusplus) 22 #pragma push_macro("__DEVICE__") 23 #define __DEVICE__ static __device__ inline __attribute__((always_inline)) 26 #pragma push_macro("__RETURN_TYPE") 27 #if defined(__cplusplus) 28 #define __RETURN_TYPE bool 30 #define __RETURN_TYPE int 33 #if defined (__cplusplus) && __cplusplus < 201103L 36 struct __compare_result{};
38 struct __compare_result<
true> {
39 static const bool valid;
43 void __suppress_unused_warning(
bool b){};
44 template <
unsigned int S,
unsigned int T>
46 __suppress_unused_warning(__compare_result<S == T>::valid);
49 #define __static_assert_type_size_equal(A, B) \ 50 __static_assert_equal_size<A,B>() 53 #define __static_assert_type_size_equal(A,B) \ 54 static_assert((A) == (B), "") 64 if (__tmp >=
'0' && __tmp <=
'7')
65 __r = (__r * 8u) + __tmp -
'0';
81 if (__tmp >=
'0' && __tmp <=
'9')
82 __r = (__r * 10u) + __tmp -
'0';
98 if (__tmp >=
'0' && __tmp <=
'9')
99 __r = (__r * 16u) + __tmp -
'0';
100 else if (__tmp >=
'a' && __tmp <=
'f')
101 __r = (__r * 16u) + __tmp -
'a' + 10;
102 else if (__tmp >=
'A' && __tmp <=
'F')
103 __r = (__r * 16u) + __tmp -
'A' + 10;
118 if (*__tagp ==
'0') {
121 if (*__tagp ==
'x' || *__tagp ==
'X')
131 #if defined(__cplusplus) 135 return (
__x ^ __sgn) - __sgn;
140 return (
__x ^ __sgn) - __sgn;
144 long long __sgn =
__x >> (
sizeof(
long long) *
CHAR_BIT - 1);
145 return (
__x ^ __sgn) - __sgn;
235 return __ocml_fma_f32(
__x,
__y, __z);
290 for (
int __i = 1; __i < __n; ++__i) {
291 float __x2 = (2 * __i) /
__x * __x1 - __x0;
300 float ldexpf(
float __x,
int __e) {
return __ocml_ldexp_f32(
__x, __e); }
342 float nanf(
const char *__tagp) {
346 unsigned int mantissa : 22;
347 unsigned int quiet : 1;
348 unsigned int exponent : 8;
349 unsigned int sign : 1;
354 __tmp.bits.sign = 0u;
355 __tmp.bits.exponent = ~0u;
356 __tmp.bits.quiet = 1u;
367 return __ocml_nextafter_f32(
__x,
__y);
372 return __ocml_len3_f32(
__x,
__y, __z);
377 return __ocml_len4_f32(
__x,
__y, __z, __w);
395 return __ocml_sqrt_f32(__r);
409 return __ocml_remainder_f32(
__x,
__y);
430 return __ocml_rlen3_f32(
__x,
__y, __z);
435 return __ocml_rlen4_f32(
__x,
__y, __z, __w);
447 return __ocml_rsqrt_f32(__r);
458 return (__n <
INT_MAX) ? __ocml_scalbn_f32(
__x, __n)
459 : __ocml_scalb_f32(
__x, __n);
527 for (
int __i = 1; __i < __n; ++__i) {
528 float __x2 = (2 * __i) /
__x * __x1 - __x0;
547 #if defined OCML_BASIC_ROUNDED_OPERATIONS 561 #if defined OCML_BASIC_ROUNDED_OPERATIONS 578 #if defined OCML_BASIC_ROUNDED_OPERATIONS 581 return __ocml_fma_rtn_f32(
__x,
__y, __z);
585 return __ocml_fma_rte_f32(
__x,
__y, __z);
589 return __ocml_fma_rtp_f32(
__x,
__y, __z);
593 return __ocml_fma_rtz_f32(
__x,
__y, __z);
598 return __ocml_fma_f32(
__x,
__y, __z);
602 #if defined OCML_BASIC_ROUNDED_OPERATIONS 616 #if defined OCML_BASIC_ROUNDED_OPERATIONS 633 #if defined OCML_BASIC_ROUNDED_OPERATIONS 647 #if defined OCML_BASIC_ROUNDED_OPERATIONS 720 return __ocml_copysign_f64(
__x,
__y);
776 return __ocml_fma_f64(
__x,
__y, __z);
829 double __x0 =
j0(
__x);
830 double __x1 =
j1(
__x);
831 for (
int __i = 1; __i < __n; ++__i) {
832 double __x2 = (2 * __i) /
__x * __x1 - __x0;
840 double ldexp(
double __x,
int __e) {
return __ocml_ldexp_f64(
__x, __e); }
883 double nan(
const char *__tagp) {
888 uint64_t mantissa : 51;
890 uint32_t exponent : 11;
896 __tmp.bits.sign = 0u;
897 __tmp.bits.exponent = ~0u;
898 __tmp.bits.quiet = 1u;
905 __val |= 0xFFF << 51;
906 return *reinterpret_cast<double *>(&__val);
915 return __ocml_nextafter_f64(
__x,
__y);
927 return __ocml_sqrt_f64(__r);
932 return __ocml_len3_f64(
__x,
__y, __z);
937 return __ocml_len4_f64(
__x,
__y, __z, __w);
957 return __ocml_remainder_f64(
__x,
__y);
985 return __ocml_rsqrt_f64(__r);
990 return __ocml_rlen3_f64(
__x,
__y, __z);
995 return __ocml_rlen4_f64(
__x,
__y, __z, __w);
1006 return (__n <
INT_MAX) ? __ocml_scalbn_f64(
__x, __n)
1007 : __ocml_scalb_f64(
__x, __n);
1019 void sincos(
double __x,
double *__sinptr,
double *__cosptr) {
1072 double __x0 =
y0(
__x);
1073 double __x1 =
y1(
__x);
1074 for (
int __i = 1; __i < __n; ++__i) {
1075 double __x2 = (2 * __i) /
__x * __x1 - __x0;
1084 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1087 return __ocml_add_rtn_f64(
__x,
__y);
1091 return __ocml_add_rte_f64(
__x,
__y);
1095 return __ocml_add_rtp_f64(
__x,
__y);
1099 return __ocml_add_rtz_f64(
__x,
__y);
1106 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1109 return __ocml_div_rtn_f64(
__x,
__y);
1113 return __ocml_div_rte_f64(
__x,
__y);
1117 return __ocml_div_rtp_f64(
__x,
__y);
1121 return __ocml_div_rtz_f64(
__x,
__y);
1128 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1131 return __ocml_mul_rtn_f64(
__x,
__y);
1135 return __ocml_mul_rte_f64(
__x,
__y);
1139 return __ocml_mul_rtp_f64(
__x,
__y);
1143 return __ocml_mul_rtz_f64(
__x,
__y);
1150 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1164 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1178 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1181 return __ocml_sub_rtn_f64(
__x,
__y);
1185 return __ocml_sub_rte_f64(
__x,
__y);
1189 return __ocml_sub_rtp_f64(
__x,
__y);
1193 return __ocml_sub_rtz_f64(
__x,
__y);
1200 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1203 return __ocml_fma_rtn_f64(
__x,
__y, __z);
1207 return __ocml_fma_rte_f64(
__x,
__y, __z);
1211 return __ocml_fma_rtp_f64(
__x,
__y, __z);
1215 return __ocml_fma_rtz_f64(
__x,
__y, __z);
1220 return __ocml_fma_f64(
__x,
__y, __z);
1227 #if !defined(__cplusplus) && __STDC_VERSION__ >= 201112L 1228 #define isfinite(__x) _Generic((__x), float : __finitef, double : __finite)(__x) 1229 #define isinf(__x) _Generic((__x), float : __isinff, double : __isinf)(__x) 1230 #define isnan(__x) _Generic((__x), float : __isnanf, double : __isnan)(__x) 1231 #define signbit(__x) \ 1232 _Generic((__x), float : __signbitf, double : __signbit)(__x) 1233 #endif // !defined(__cplusplus) && __STDC_VERSION__ >= 201112L 1235 #if defined(__cplusplus) 1237 return (__arg1 < __arg2) ? __arg1 : __arg2;
1241 return (__arg1 > __arg2) ? __arg1 : __arg2;
1245 return (__arg1 < __arg2) ? __arg1 : __arg2;
1248 return (__arg1 > __arg2) ? __arg1 : __arg2;
1263 __host__
inline static int min(
int __arg1,
int __arg2) {
1267 __host__
inline static int max(
int __arg1,
int __arg2) {
1272 #pragma pop_macro("__DEVICE__") 1273 #pragma pop_macro("__RETURN_TYPE") 1275 #endif // __CLANG_HIP_MATH_H__ __DEVICE__ float rnormf(int __dim, const float *__a)
static __inline unsigned char unsigned int unsigned int __y
__DEVICE__ float sqrtf(float __x)
__device__ float __ocml_lgamma_f32(float)
__DEVICE__ float atan2f(float __x, float __y)
__DEVICE__ float erfcxf(float __x)
__DEVICE__ float __fmaf_ru(float __a, float __b, float __c)
__DEVICE__ float fmodf(float __x, float __y)
__DEVICE__ float __fmul_ru(float __a, float __b)
__DEVICE__ float atanhf(float __x)
__DEVICE__ double frexp(double __x, int *__nptr)
__DEVICE__ float cbrtf(float __x)
__DEVICE__ float coshf(float __x)
__DEVICE__ double trunc(double __x)
__DEVICE__ double sin(double __x)
__DEVICE__ float __fsub_rz(float __a, float __b)
__DEVICE__ float __frcp_rd(float __a)
__DEVICE__ double remainder(double __x, double __y)
__DEVICE__ double __fma_rz(double __a, double __b, double __c)
__DEVICE__ float truncf(float __x)
__DEVICE__ double __dmul_rn(double __x, double __y)
__device__ double __ocml_lgamma_f64(double)
__DEVICE__ float __sinf(float __x)
__DEVICE__ int max(int __a, int __b)
__DEVICE__ float normf(int __dim, const float *__a)
__DEVICE__ long int lrintf(float __x)
__DEVICE__ double copysign(double __x, double __y)
__DEVICE__ double expm1(double __x)
__device__ double __ocml_cospi_f64(double)
__DEVICE__ long long abs(long long __n)
__DEVICE__ float __fsub_rn(float __x, float __y)
__DEVICE__ float sinf(float __x)
__device__ double __ocml_i0_f64(double)
__DEVICE__ float rcbrtf(float __x)
__device__ float __ocml_y0_f32(float)
__DEVICE__ double ldexp(double __x, int __e)
__DEVICE__ float __fdiv_rn(float __x, float __y)
__DEVICE__ long int lround(double __x)
__DEVICE__ float y1f(float __x)
__DEVICE__ double rcbrt(double __x)
__DEVICE__ double asin(double __x)
__DEVICE__ float scalblnf(float __x, long int __n)
__DEVICE__ float lgammaf(float __x)
__DEVICE__ void __sincosf(float __x, float *__sinptr, float *__cosptr)
__DEVICE__ int ilogbf(float __x)
__DEVICE__ double __dsqrt_ru(double __a)
__DEVICE__ long labs(long __a)
__DEVICE__ float __exp10f(float __x)
__DEVICE__ float y0f(float __x)
__device__ float __ocml_tgamma_f32(float)
__DEVICE__ double fabs(double __x)
__DEVICE__ double log10(double __x)
__DEVICE__ double sqrt(double __x)
__DEVICE__ double rnorm(int __dim, const double *__a)
__DEVICE__ float acoshf(float __x)
__DEVICE__ double scalbn(double __x, int __n)
__DEVICE__ __RETURN_TYPE __signbitf(float __x)
__DEVICE__ float __log10f(float __x)
__device__ double __ocml_cos_f64(double)
__DEVICE__ long long int llrint(double __x)
__DEVICE__ float fmaf(float __x, float __y, float __z)
__DEVICE__ double nextafter(double __x, double __y)
__DEVICE__ double normcdf(double __x)
__DEVICE__ long long llabs(long long __a)
__DEVICE__ float modff(float __x, float *__iptr)
__device__ float __ocml_modf_f32(float, __attribute__((address_space(5))) float *)
__DEVICE__ float __fsub_ru(float __a, float __b)
__DEVICE__ float tanhf(float __x)
__DEVICE__ double tanh(double __x)
__DEVICE__ double norm4d(double __x, double __y, double __z, double __w)
__DEVICE__ uint64_t __make_mantissa_base16(const char *__tagp)
__DEVICE__ float expm1f(float __x)
__DEVICE__ float __fadd_rd(float __a, float __b)
__DEVICE__ double cyl_bessel_i1(double __x)
__DEVICE__ float cyl_bessel_i1f(float __x)
__DEVICE__ double __dmul_rd(double __a, double __b)
__DEVICE__ double __dsub_ru(double __a, double __b)
__DEVICE__ float erfcf(float __x)
__device__ float __ocml_i1_f32(float)
__DEVICE__ double j0(double __x)
__DEVICE__ double exp10(double __x)
__DEVICE__ double fmin(double __x, double __y)
__DEVICE__ float __fmaf_rd(float __a, float __b, float __c)
__DEVICE__ double floor(double __x)
__DEVICE__ double lgamma(double __x)
__DEVICE__ long long int llround(double __x)
__DEVICE__ double erf(double __x)
__DEVICE__ float erfcinvf(float __x)
__device__ double __ocml_tan_f64(double)
__DEVICE__ float exp2f(float __x)
__DEVICE__ float __fmaf_rz(float __a, float __b, float __c)
__DEVICE__ float acosf(float __x)
__DEVICE__ float ynf(int __n, float __x)
__DEVICE__ float norm4df(float __x, float __y, float __z, float __w)
__DEVICE__ float frexpf(float __x, int *__nptr)
__DEVICE__ float __fdiv_rd(float __a, float __b)
__DEVICE__ double rsqrt(double __x)
__DEVICE__ float __fsqrt_rd(float __a)
__DEVICE__ float __fmul_rd(float __a, float __b)
__device__ float __ocml_native_sin_f32(float)
__DEVICE__ float __fmul_rn(float __x, float __y)
__device__ float __ocml_j1_f32(float)
static __inline__ void int __a
__DEVICE__ float logbf(float __x)
__DEVICE__ double __dsub_rd(double __a, double __b)
__DEVICE__ float __fsqrt_rz(float __a)
__DEVICE__ double __dadd_rd(double __a, double __b)
__DEVICE__ float cosf(float __x)
__device__ double __ocml_sincospi_f64(double, __attribute__((address_space(5))) double *)
__DEVICE__ double log1p(double __x)
__DEVICE__ double __dsqrt_rd(double __a)
__DEVICE__ float erff(float __x)
__DEVICE__ float __fmaf_rn(float __x, float __y, float __z)
__DEVICE__ float j1f(float __x)
__DEVICE__ void sincospi(double __x, double *__sinptr, double *__cosptr)
__DEVICE__ double __fma_rn(double __x, double __y, double __z)
__DEVICE__ float roundf(float __x)
__DEVICE__ double __ddiv_rn(double __x, double __y)
__DEVICE__ double __ddiv_rd(double __a, double __b)
__DEVICE__ double tgamma(double __x)
__DEVICE__ double acosh(double __x)
__DEVICE__ float rintf(float __x)
__DEVICE__ uint64_t __make_mantissa(const char *__tagp)
__DEVICE__ double acos(double __x)
__DEVICE__ long int lroundf(float __x)
__DEVICE__ double normcdfinv(double __x)
__DEVICE__ float __fdiv_rz(float __a, float __b)
__DEVICE__ int ilogb(double __x)
__device__ float __ocml_remquo_f32(float, float, __attribute__((address_space(5))) int *)
__device__ double __ocml_remquo_f64(double, double, __attribute__((address_space(5))) int *)
__device__ double __ocml_tgamma_f64(double)
__DEVICE__ double yn(int __n, double __x)
__DEVICE__ float ceilf(float __x)
__DEVICE__ double logb(double __x)
__DEVICE__ float log2f(float __x)
__DEVICE__ __RETURN_TYPE __finite(double __x)
__DEVICE__ float log10f(float __x)
__device__ float __ocml_cos_f32(float)
__DEVICE__ void sincospif(float __x, float *__sinptr, float *__cosptr)
__device__ float __ocml_cospi_f32(float)
__DEVICE__ float __fsqrt_ru(float __a)
__DEVICE__ double cosh(double __x)
__DEVICE__ double cbrt(double __x)
__DEVICE__ double exp(double __x)
__DEVICE__ __RETURN_TYPE __finitef(float __x)
__DEVICE__ float __powf(float __x, float __y)
__DEVICE__ float __fadd_ru(float __a, float __b)
__DEVICE__ float j0f(float __x)
__DEVICE__ float rsqrtf(float __x)
__DEVICE__ double __drcp_ru(double __a)
__DEVICE__ double j1(double __x)
__device__ float __ocml_sinpi_f32(float)
__device__ float __ocml_j0_f32(float)
__DEVICE__ double exp2(double __x)
__DEVICE__ float asinhf(float __x)
__device__ double __ocml_y1_f64(double)
__DEVICE__ float copysignf(float __x, float __y)
__DEVICE__ float rnorm3df(float __x, float __y, float __z)
__device__ float __ocml_frexp_f32(float, __attribute__((address_space(5))) int *)
__DEVICE__ double __dsqrt_rn(double __x)
__DEVICE__ double norm3d(double __x, double __y, double __z)
__DEVICE__ void sincosf(float __x, float *__sinptr, float *__cosptr)
__DEVICE__ float asinf(float __x)
__DEVICE__ double rnorm3d(double __x, double __y, double __z)
__DEVICE__ double sinh(double __x)
__DEVICE__ double fmod(double __x, double __y)
__DEVICE__ float normcdfinvf(float __x)
__DEVICE__ double fdim(double __x, double __y)
__DEVICE__ double asinh(double __x)
__DEVICE__ double modf(double __x, double *__iptr)
__DEVICE__ double __dadd_rn(double __x, double __y)
__device__ float __ocml_y1_f32(float)
__DEVICE__ float cyl_bessel_i0f(float __x)
static __inline unsigned char unsigned int __x
__DEVICE__ float __frsqrt_rn(float __x)
__DEVICE__ float log1pf(float __x)
__DEVICE__ double __drcp_rn(double __x)
__DEVICE__ double __fma_ru(double __a, double __b, double __c)
__DEVICE__ float rnorm4df(float __x, float __y, float __z, float __w)
__DEVICE__ double y1(double __x)
__DEVICE__ double log(double __x)
__DEVICE__ uint64_t __make_mantissa_base8(const char *__tagp)
__DEVICE__ double fma(double __x, double __y, double __z)
__DEVICE__ float __frcp_ru(float __a)
__DEVICE__ double sinpi(double __x)
__DEVICE__ double __dadd_ru(double __a, double __b)
__DEVICE__ float scalbnf(float __x, int __n)
__DEVICE__ double cospi(double __x)
__DEVICE__ double __dadd_rz(double __a, double __b)
__device__ float __ocml_sincospi_f32(float, __attribute__((address_space(5))) float *)
__device__ double __ocml_j1_f64(double)
__DEVICE__ float powf(float __x, float __y)
#define __static_assert_type_size_equal(A, B)
__DEVICE__ double hypot(double __x, double __y)
__DEVICE__ double __ddiv_rz(double __a, double __b)
__DEVICE__ float __fdiv_ru(float __a, float __b)
__DEVICE__ double cyl_bessel_i0(double __x)
__DEVICE__ float tgammaf(float __x)
__DEVICE__ __RETURN_TYPE __isnan(double __x)
__DEVICE__ double log2(double __x)
__DEVICE__ float powif(float __x, int __y)
__DEVICE__ float __fdividef(float __x, float __y)
__DEVICE__ long long int llroundf(float __x)
__DEVICE__ double rint(double __x)
__DEVICE__ float __log2f(float __x)
__DEVICE__ float nanf(const char *__tagp)
__device__ double __ocml_i1_f64(double)
__device__ double __ocml_sin_f64(double)
__DEVICE__ float __logf(float __x)
__DEVICE__ double jn(int __n, double __x)
__DEVICE__ double scalbln(double __x, long int __n)
__DEVICE__ float nearbyintf(float __x)
__DEVICE__ float normcdff(float __x)
__device__ double __ocml_sincos_f64(double, __attribute__((address_space(5))) double *)
__DEVICE__ double __dsqrt_rz(double __a)
__DEVICE__ float __frcp_rn(float __x)
__DEVICE__ float fdividef(float __x, float __y)
__DEVICE__ double __dsub_rn(double __x, double __y)
__device__ double __ocml_frexp_f64(double, __attribute__((address_space(5))) int *)
__DEVICE__ double erfinv(double __x)
__DEVICE__ double __fma_rd(double __a, double __b, double __c)
_Float16 __2f16 __attribute__((ext_vector_type(2)))
Zeroes the upper 128 bits (bits 255:128) of all YMM registers.
__device__ float __ocml_native_cos_f32(float)
__device__ double __ocml_y0_f64(double)
__DEVICE__ float __fadd_rz(float __a, float __b)
__DEVICE__ double rnorm4d(double __x, double __y, double __z, double __w)
__DEVICE__ double powi(double __x, int __y)
__DEVICE__ float fabsf(float __x)
__DEVICE__ float remquof(float __x, float __y, int *__quo)
__device__ double __ocml_modf_f64(double, __attribute__((address_space(5))) double *)
__DEVICE__ __RETURN_TYPE __isinff(float __x)
__device__ float __ocml_tan_f32(float)
float __ovld __cnfn sign(float x)
Returns 1.0 if x > 0, -0.0 if x = -0.0, +0.0 if x = +0.0, or -1.0 if x < 0.
__DEVICE__ __RETURN_TYPE __isinf(double __x)
__DEVICE__ float erfinvf(float __x)
__DEVICE__ void sincos(double __x, double *__sinptr, double *__cosptr)
__DEVICE__ float exp10f(float __x)
__DEVICE__ double erfcinv(double __x)
__DEVICE__ float remainderf(float __x, float __y)
__DEVICE__ double nan(const char *__tagp)
__DEVICE__ double norm(int __dim, const double *__a)
__DEVICE__ float floorf(float __x)
__DEVICE__ double nearbyint(double __x)
__DEVICE__ float __fsub_rd(float __a, float __b)
__DEVICE__ float sinhf(float __x)
__DEVICE__ float hypotf(float __x, float __y)
__DEVICE__ double pow(double __x, double __y)
__DEVICE__ float __cosf(float __x)
__device__ float __ocml_sincos_f32(float, __attribute__((address_space(5))) float *)
__DEVICE__ double rhypot(double __x, double __y)
__DEVICE__ float sinpif(float __x)
__DEVICE__ __RETURN_TYPE __signbit(double __x)
__DEVICE__ float tanf(float __x)
__DEVICE__ float cospif(float __x)
__DEVICE__ double atanh(double __x)
__DEVICE__ long int lrint(double __x)
__device__ float __ocml_sin_f32(float)
__DEVICE__ double __dsub_rz(double __a, double __b)
__DEVICE__ float __fsqrt_rn(float __x)
__DEVICE__ float logf(float __x)
__DEVICE__ double remquo(double __x, double __y, int *__quo)
__DEVICE__ double erfc(double __x)
__DEVICE__ float nextafterf(float __x, float __y)
__DEVICE__ float fmaxf(float __x, float __y)
__DEVICE__ float __frcp_rz(float __a)
__DEVICE__ float ldexpf(float __x, int __e)
__DEVICE__ double ceil(double __x)
__DEVICE__ int min(int __a, int __b)
__DEVICE__ double __ddiv_ru(double __a, double __b)
__DEVICE__ double round(double __x)
__DEVICE__ float __tanf(float __x)
__DEVICE__ double cos(double __x)
__DEVICE__ float expf(float __x)
__DEVICE__ float __fadd_rn(float __x, float __y)
__DEVICE__ uint64_t __make_mantissa_base10(const char *__tagp)
__device__ float __ocml_i0_f32(float)
__device__ double __ocml_sinpi_f64(double)
__DEVICE__ double __dmul_rz(double __a, double __b)
__DEVICE__ double atan(double __x)
__DEVICE__ float jnf(int __n, float __x)
__DEVICE__ double atan2(double __x, double __y)
__DEVICE__ float __saturatef(float __x)
__DEVICE__ float atanf(float __x)
__DEVICE__ long long int llrintf(float __x)
__DEVICE__ double y0(double __x)
__DEVICE__ double tan(double __x)
__device__ double __ocml_j0_f64(double)
__DEVICE__ double __drcp_rd(double __a)
__DEVICE__ __RETURN_TYPE __isnanf(float __x)
__DEVICE__ double fmax(double __x, double __y)
__DEVICE__ double __dmul_ru(double __a, double __b)
__DEVICE__ double erfcx(double __x)
__DEVICE__ float __expf(float __x)
__DEVICE__ float norm3df(float __x, float __y, float __z)
__DEVICE__ double __drcp_rz(double __a)
__DEVICE__ float rhypotf(float __x, float __y)
__DEVICE__ float fminf(float __x, float __y)
__DEVICE__ float fdimf(float __x, float __y)
__DEVICE__ float __fmul_rz(float __a, float __b)