25 #include "hip_fp16_math_fwd.h" 36 #if !__HIP_DEVICE_COMPILE__ 47 #include "kalmar_math.h" 50 #pragma push_macro("__DEVICE__") 51 #pragma push_macro("__RETURN_TYPE") 54 #define __DEVICE__ __device__ 55 #define __RETURN_TYPE int 56 #else // to be consistent with __clang_cuda_math_forward_declares 57 #define __DEVICE__ static __device__ 58 #define __RETURN_TYPE bool 63 uint64_t __make_mantissa_base8(
const char* tagp)
69 if (tmp >=
'0' && tmp <=
'7') r = (r * 8u) + tmp -
'0';
80 uint64_t __make_mantissa_base10(
const char* tagp)
86 if (tmp >=
'0' && tmp <=
'9') r = (r * 10u) + tmp -
'0';
97 uint64_t __make_mantissa_base16(
const char* tagp)
103 if (tmp >=
'0' && tmp <=
'9') r = (r * 16u) + tmp -
'0';
104 else if (tmp >=
'a' && tmp <=
'f') r = (r * 16u) + tmp -
'a' + 10;
105 else if (tmp >=
'A' && tmp <=
'F') r = (r * 16u) + tmp -
'A' + 10;
116 uint64_t __make_mantissa(
const char* tagp)
118 if (!tagp)
return 0u;
123 if (*tagp ==
'x' || *tagp ==
'X')
return __make_mantissa_base16(tagp);
124 else return __make_mantissa_base8(tagp);
127 return __make_mantissa_base10(tagp);
131 #if (__hcc_workweek__ >= 19015) || __HIP_CLANG_ONLY__ 134 int amd_mixed_dot(short2 a, short2 b,
int c,
bool saturate) {
135 return __ockl_sdot2(a.data, b.data, c, saturate);
139 uint amd_mixed_dot(ushort2 a, ushort2 b, uint c,
bool saturate) {
140 return __ockl_udot2(a.data, b.data, c, saturate);
144 int amd_mixed_dot(char4 a, char4 b,
int c,
bool saturate) {
145 return __ockl_sdot4(a.data, b.data, c, saturate);
149 uint amd_mixed_dot(uchar4 a, uchar4 b, uint c,
bool saturate) {
150 return __ockl_udot4(a.data, b.data, c, saturate);
154 int amd_mixed_dot(
int a,
int b,
int c,
bool saturate) {
155 return __ockl_sdot8(a, b, c, saturate);
159 uint amd_mixed_dot(uint a, uint b, uint c,
bool saturate) {
160 return __ockl_udot8(a, b, c, saturate);
167 float abs(
float x) {
return __ocml_fabs_f32(x); }
170 float acosf(
float x) {
return __ocml_acos_f32(x); }
173 float acoshf(
float x) {
return __ocml_acosh_f32(x); }
176 float asinf(
float x) {
return __ocml_asin_f32(x); }
179 float asinhf(
float x) {
return __ocml_asinh_f32(x); }
182 float atan2f(
float x,
float y) {
return __ocml_atan2_f32(x, y); }
185 float atanf(
float x) {
return __ocml_atan_f32(x); }
188 float atanhf(
float x) {
return __ocml_atanh_f32(x); }
191 float cbrtf(
float x) {
return __ocml_cbrt_f32(x); }
194 float ceilf(
float x) {
return __ocml_ceil_f32(x); }
197 float copysignf(
float x,
float y) {
return __ocml_copysign_f32(x, y); }
200 float cosf(
float x) {
return __ocml_cos_f32(x); }
203 float coshf(
float x) {
return __ocml_cosh_f32(x); }
206 float cospif(
float x) {
return __ocml_cospi_f32(x); }
209 float cyl_bessel_i0f(
float x) {
return __ocml_i0_f32(x); }
212 float cyl_bessel_i1f(
float x) {
return __ocml_i1_f32(x); }
215 float erfcf(
float x) {
return __ocml_erfc_f32(x); }
218 float erfcinvf(
float x) {
return __ocml_erfcinv_f32(x); }
221 float erfcxf(
float x) {
return __ocml_erfcx_f32(x); }
224 float erff(
float x) {
return __ocml_erf_f32(x); }
227 float erfinvf(
float x) {
return __ocml_erfinv_f32(x); }
230 float exp10f(
float x) {
return __ocml_exp10_f32(x); }
233 float exp2f(
float x) {
return __ocml_exp2_f32(x); }
236 float expf(
float x) {
return __ocml_exp_f32(x); }
239 float expm1f(
float x) {
return __ocml_expm1_f32(x); }
242 float fabsf(
float x) {
return __ocml_fabs_f32(x); }
245 float fdimf(
float x,
float y) {
return __ocml_fdim_f32(x, y); }
248 float fdividef(
float x,
float y) {
return x / y; }
251 float floorf(
float x) {
return __ocml_floor_f32(x); }
254 float fmaf(
float x,
float y,
float z) {
return __ocml_fma_f32(x, y, z); }
257 float fmaxf(
float x,
float y) {
return __ocml_fmax_f32(x, y); }
260 float fminf(
float x,
float y) {
return __ocml_fmin_f32(x, y); }
263 float fmodf(
float x,
float y) {
return __ocml_fmod_f32(x, y); }
266 float frexpf(
float x,
int* nptr)
270 __ocml_frexp_f32(x, (__attribute__((address_space(5)))
int*) &tmp);
277 float hypotf(
float x,
float y) {
return __ocml_hypot_f32(x, y); }
280 int ilogbf(
float x) {
return __ocml_ilogb_f32(x); }
283 __RETURN_TYPE isfinite(
float x) {
return __ocml_isfinite_f32(x); }
286 __RETURN_TYPE isinf(
float x) {
return __ocml_isinf_f32(x); }
289 __RETURN_TYPE isnan(
float x) {
return __ocml_isnan_f32(x); }
292 float j0f(
float x) {
return __ocml_j0_f32(x); }
295 float j1f(
float x) {
return __ocml_j1_f32(x); }
298 float jnf(
int n,
float x)
302 if (n == 0)
return j0f(x);
303 if (n == 1)
return j1f(x);
307 for (
int i = 1; i < n; ++i) {
308 float x2 = (2 * i) / x * x1 - x0;
317 float ldexpf(
float x,
int e) {
return __ocml_ldexp_f32(x, e); }
320 float lgammaf(
float x) {
return __ocml_lgamma_f32(x); }
323 long long int llrintf(
float x) {
return __ocml_rint_f32(x); }
326 long long int llroundf(
float x) {
return __ocml_round_f32(x); }
329 float log10f(
float x) {
return __ocml_log10_f32(x); }
332 float log1pf(
float x) {
return __ocml_log1p_f32(x); }
335 float log2f(
float x) {
return __ocml_log2_f32(x); }
338 float logbf(
float x) {
return __ocml_logb_f32(x); }
341 float logf(
float x) {
return __ocml_log_f32(x); }
344 long int lrintf(
float x) {
return __ocml_rint_f32(x); }
347 long int lroundf(
float x) {
return __ocml_round_f32(x); }
350 float modff(
float x,
float* iptr)
354 __ocml_modf_f32(x, (__attribute__((address_space(5)))
float*) &tmp);
361 float nanf(
const char* tagp)
366 uint32_t mantissa : 22;
368 uint32_t exponent : 8;
372 static_assert(
sizeof(
float) ==
sizeof(ieee_float),
"");
376 tmp.bits.exponent = ~0u;
378 tmp.bits.mantissa = __make_mantissa(tagp);
384 float nearbyintf(
float x) {
return __ocml_nearbyint_f32(x); }
387 float nextafterf(
float x,
float y) {
return __ocml_nextafter_f32(x, y); }
390 float norm3df(
float x,
float y,
float z) {
return __ocml_len3_f32(x, y, z); }
393 float norm4df(
float x,
float y,
float z,
float w)
395 return __ocml_len4_f32(x, y, z, w);
399 float normcdff(
float x) {
return __ocml_ncdf_f32(x); }
402 float normcdfinvf(
float x) {
return __ocml_ncdfinv_f32(x); }
405 float normf(
int dim,
const float* a)
408 while (dim--) { r += a[0] * a[0]; ++a; }
410 return __ocml_sqrt_f32(r);
414 float powf(
float x,
float y) {
return __ocml_pow_f32(x, y); }
417 float rcbrtf(
float x) {
return __ocml_rcbrt_f32(x); }
420 float remainderf(
float x,
float y) {
return __ocml_remainder_f32(x, y); }
423 float remquof(
float x,
float y,
int* quo)
427 __ocml_remquo_f32(x, y, (__attribute__((address_space(5)))
int*) &tmp);
434 float rhypotf(
float x,
float y) {
return __ocml_rhypot_f32(x, y); }
437 float rintf(
float x) {
return __ocml_rint_f32(x); }
440 float rnorm3df(
float x,
float y,
float z)
442 return __ocml_rlen3_f32(x, y, z);
447 float rnorm4df(
float x,
float y,
float z,
float w)
449 return __ocml_rlen4_f32(x, y, z, w);
453 float rnormf(
int dim,
const float* a)
456 while (dim--) { r += a[0] * a[0]; ++a; }
458 return __ocml_rsqrt_f32(r);
462 float roundf(
float x) {
return __ocml_round_f32(x); }
465 float rsqrtf(
float x) {
return __ocml_rsqrt_f32(x); }
468 float scalblnf(
float x,
long int n)
470 return (n < INT_MAX) ? __ocml_scalbn_f32(x, n) : __ocml_scalb_f32(x, n);
474 float scalbnf(
float x,
int n) {
return __ocml_scalbn_f32(x, n); }
477 __RETURN_TYPE signbit(
float x) {
return __ocml_signbit_f32(x); }
480 void sincosf(
float x,
float* sptr,
float* cptr)
485 __ocml_sincos_f32(x, (__attribute__((address_space(5)))
float*) &tmp);
490 void sincospif(
float x,
float* sptr,
float* cptr)
495 __ocml_sincospi_f32(x, (__attribute__((address_space(5)))
float*) &tmp);
500 float sinf(
float x) {
return __ocml_sin_f32(x); }
503 float sinhf(
float x) {
return __ocml_sinh_f32(x); }
506 float sinpif(
float x) {
return __ocml_sinpi_f32(x); }
509 float sqrtf(
float x) {
return __ocml_sqrt_f32(x); }
512 float tanf(
float x) {
return __ocml_tan_f32(x); }
515 float tanhf(
float x) {
return __ocml_tanh_f32(x); }
518 float tgammaf(
float x) {
return __ocml_tgamma_f32(x); }
521 float truncf(
float x) {
return __ocml_trunc_f32(x); }
524 float y0f(
float x) {
return __ocml_y0_f32(x); }
527 float y1f(
float x) {
return __ocml_y1_f32(x); }
530 float ynf(
int n,
float x)
535 if (n == 0)
return y0f(x);
536 if (n == 1)
return y1f(x);
540 for (
int i = 1; i < n; ++i) {
541 float x2 = (2 * i) / x * x1 - x0;
552 float __cosf(
float x) {
return __ocml_native_cos_f32(x); }
555 float __exp10f(
float x) {
return __ocml_native_exp10_f32(x); }
558 float __expf(
float x) {
return __ocml_native_exp_f32(x); }
559 #if defined OCML_BASIC_ROUNDED_OPERATIONS 562 float __fadd_rd(
float x,
float y) {
return __ocml_add_rtn_f32(x, y); }
566 float __fadd_rn(
float x,
float y) {
return x + y; }
567 #if defined OCML_BASIC_ROUNDED_OPERATIONS 570 float __fadd_ru(
float x,
float y) {
return __ocml_add_rtp_f32(x, y); }
573 float __fadd_rz(
float x,
float y) {
return __ocml_add_rtz_f32(x, y); }
576 float __fdiv_rd(
float x,
float y) {
return __ocml_div_rtn_f32(x, y); }
580 float __fdiv_rn(
float x,
float y) {
return x / y; }
581 #if defined OCML_BASIC_ROUNDED_OPERATIONS 584 float __fdiv_ru(
float x,
float y) {
return __ocml_div_rtp_f32(x, y); }
587 float __fdiv_rz(
float x,
float y) {
return __ocml_div_rtz_f32(x, y); }
591 float __fdividef(
float x,
float y) {
return x / y; }
592 #if defined OCML_BASIC_ROUNDED_OPERATIONS 595 float __fmaf_rd(
float x,
float y,
float z)
597 return __ocml_fma_rtn_f32(x, y, z);
602 float __fmaf_rn(
float x,
float y,
float z)
604 return __ocml_fma_f32(x, y, z);
606 #if defined OCML_BASIC_ROUNDED_OPERATIONS 609 float __fmaf_ru(
float x,
float y,
float z)
611 return __ocml_fma_rtp_f32(x, y, z);
615 float __fmaf_rz(
float x,
float y,
float z)
617 return __ocml_fma_rtz_f32(x, y, z);
621 float __fmul_rd(
float x,
float y) {
return __ocml_mul_rtn_f32(x, y); }
625 float __fmul_rn(
float x,
float y) {
return x * y; }
626 #if defined OCML_BASIC_ROUNDED_OPERATIONS 629 float __fmul_ru(
float x,
float y) {
return __ocml_mul_rtp_f32(x, y); }
632 float __fmul_rz(
float x,
float y) {
return __ocml_mul_rtz_f32(x, y); }
635 float __frcp_rd(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
639 float __frcp_rn(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
640 #if defined OCML_BASIC_ROUNDED_OPERATIONS 643 float __frcp_ru(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
646 float __frcp_rz(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
650 float __frsqrt_rn(
float x) {
return __llvm_amdgcn_rsq_f32(x); }
651 #if defined OCML_BASIC_ROUNDED_OPERATIONS 654 float __fsqrt_rd(
float x) {
return __ocml_sqrt_rtn_f32(x); }
658 float __fsqrt_rn(
float x) {
return __ocml_native_sqrt_f32(x); }
659 #if defined OCML_BASIC_ROUNDED_OPERATIONS 662 float __fsqrt_ru(
float x) {
return __ocml_sqrt_rtp_f32(x); }
665 float __fsqrt_rz(
float x) {
return __ocml_sqrt_rtz_f32(x); }
668 float __fsub_rd(
float x,
float y) {
return __ocml_sub_rtn_f32(x, y); }
672 float __fsub_rn(
float x,
float y) {
return x - y; }
673 #if defined OCML_BASIC_ROUNDED_OPERATIONS 676 float __fsub_ru(
float x,
float y) {
return __ocml_sub_rtp_f32(x, y); }
679 float __fsub_rz(
float x,
float y) {
return __ocml_sub_rtz_f32(x, y); }
683 float __log10f(
float x) {
return __ocml_native_log10_f32(x); }
686 float __log2f(
float x) {
return __ocml_native_log2_f32(x); }
689 float __logf(
float x) {
return __ocml_native_log_f32(x); }
692 float __powf(
float x,
float y) {
return __ocml_pow_f32(x, y); }
695 float __saturatef(
float x) {
return (x < 0) ? 0 : ((x > 1) ? 1 : x); }
698 void __sincosf(
float x,
float* sptr,
float* cptr)
700 *sptr = __ocml_native_sin_f32(x);
701 *cptr = __ocml_native_cos_f32(x);
705 float __sinf(
float x) {
return __ocml_native_sin_f32(x); }
708 float __tanf(
float x) {
return __ocml_tan_f32(x); }
715 double abs(
double x) {
return __ocml_fabs_f64(x); }
718 double acos(
double x) {
return __ocml_acos_f64(x); }
721 double acosh(
double x) {
return __ocml_acosh_f64(x); }
724 double asin(
double x) {
return __ocml_asin_f64(x); }
727 double asinh(
double x) {
return __ocml_asinh_f64(x); }
730 double atan(
double x) {
return __ocml_atan_f64(x); }
733 double atan2(
double x,
double y) {
return __ocml_atan2_f64(x, y); }
736 double atanh(
double x) {
return __ocml_atanh_f64(x); }
739 double cbrt(
double x) {
return __ocml_cbrt_f64(x); }
742 double ceil(
double x) {
return __ocml_ceil_f64(x); }
745 double copysign(
double x,
double y) {
return __ocml_copysign_f64(x, y); }
748 double cos(
double x) {
return __ocml_cos_f64(x); }
751 double cosh(
double x) {
return __ocml_cosh_f64(x); }
754 double cospi(
double x) {
return __ocml_cospi_f64(x); }
757 double cyl_bessel_i0(
double x) {
return __ocml_i0_f64(x); }
760 double cyl_bessel_i1(
double x) {
return __ocml_i1_f64(x); }
763 double erf(
double x) {
return __ocml_erf_f64(x); }
766 double erfc(
double x) {
return __ocml_erfc_f64(x); }
769 double erfcinv(
double x) {
return __ocml_erfcinv_f64(x); }
772 double erfcx(
double x) {
return __ocml_erfcx_f64(x); }
775 double erfinv(
double x) {
return __ocml_erfinv_f64(x); }
778 double exp(
double x) {
return __ocml_exp_f64(x); }
781 double exp10(
double x) {
return __ocml_exp10_f64(x); }
784 double exp2(
double x) {
return __ocml_exp2_f64(x); }
787 double expm1(
double x) {
return __ocml_expm1_f64(x); }
790 double fabs(
double x) {
return __ocml_fabs_f64(x); }
793 double fdim(
double x,
double y) {
return __ocml_fdim_f64(x, y); }
796 double floor(
double x) {
return __ocml_floor_f64(x); }
799 double fma(
double x,
double y,
double z) {
return __ocml_fma_f64(x, y, z); }
802 double fmax(
double x,
double y) {
return __ocml_fmax_f64(x, y); }
805 double fmin(
double x,
double y) {
return __ocml_fmin_f64(x, y); }
808 double fmod(
double x,
double y) {
return __ocml_fmod_f64(x, y); }
811 double frexp(
double x,
int* nptr)
815 __ocml_frexp_f64(x, (__attribute__((address_space(5)))
int*) &tmp);
822 double hypot(
double x,
double y) {
return __ocml_hypot_f64(x, y); }
825 int ilogb(
double x) {
return __ocml_ilogb_f64(x); }
828 __RETURN_TYPE isfinite(
double x) {
return __ocml_isfinite_f64(x); }
831 __RETURN_TYPE isinf(
double x) {
return __ocml_isinf_f64(x); }
834 __RETURN_TYPE isnan(
double x) {
return __ocml_isnan_f64(x); }
837 double j0(
double x) {
return __ocml_j0_f64(x); }
840 double j1(
double x) {
return __ocml_j1_f64(x); }
843 double jn(
int n,
double x)
848 if (n == 0)
return j0f(x);
849 if (n == 1)
return j1f(x);
853 for (
int i = 1; i < n; ++i) {
854 double x2 = (2 * i) / x * x1 - x0;
863 double ldexp(
double x,
int e) {
return __ocml_ldexp_f64(x, e); }
866 double lgamma(
double x) {
return __ocml_lgamma_f64(x); }
869 long long int llrint(
double x) {
return __ocml_rint_f64(x); }
872 long long int llround(
double x) {
return __ocml_round_f64(x); }
875 double log(
double x) {
return __ocml_log_f64(x); }
878 double log10(
double x) {
return __ocml_log10_f64(x); }
881 double log1p(
double x) {
return __ocml_log1p_f64(x); }
884 double log2(
double x) {
return __ocml_log2_f64(x); }
887 double logb(
double x) {
return __ocml_logb_f64(x); }
890 long int lrint(
double x) {
return __ocml_rint_f64(x); }
893 long int lround(
double x) {
return __ocml_round_f64(x); }
896 double modf(
double x,
double* iptr)
900 __ocml_modf_f64(x, (__attribute__((address_space(5)))
double*) &tmp);
907 double nan(
const char* tagp)
913 uint64_t mantissa : 51;
915 uint32_t exponent : 11;
918 static_assert(
sizeof(
double) ==
sizeof(ieee_double),
"");
922 tmp.bits.exponent = ~0u;
924 tmp.bits.mantissa = __make_mantissa(tagp);
928 static_assert(
sizeof(uint64_t)==
sizeof(
double));
929 uint64_t val = __make_mantissa(tagp);
931 return *
reinterpret_cast<double*
>(&val);
936 double nearbyint(
double x) {
return __ocml_nearbyint_f64(x); }
939 double nextafter(
double x,
double y) {
return __ocml_nextafter_f64(x, y); }
942 double norm(
int dim,
const double* a)
945 while (dim--) { r += a[0] * a[0]; ++a; }
947 return __ocml_sqrt_f64(r);
951 double norm3d(
double x,
double y,
double z)
953 return __ocml_len3_f64(x, y, z);
957 double norm4d(
double x,
double y,
double z,
double w)
959 return __ocml_len4_f64(x, y, z, w);
963 double normcdf(
double x) {
return __ocml_ncdf_f64(x); }
966 double normcdfinv(
double x) {
return __ocml_ncdfinv_f64(x); }
969 double pow(
double x,
double y) {
return __ocml_pow_f64(x, y); }
972 double rcbrt(
double x) {
return __ocml_rcbrt_f64(x); }
975 double remainder(
double x,
double y) {
return __ocml_remainder_f64(x, y); }
978 double remquo(
double x,
double y,
int* quo)
982 __ocml_remquo_f64(x, y, (__attribute__((address_space(5)))
int*) &tmp);
989 double rhypot(
double x,
double y) {
return __ocml_rhypot_f64(x, y); }
992 double rint(
double x) {
return __ocml_rint_f64(x); }
995 double rnorm(
int dim,
const double* a)
998 while (dim--) { r += a[0] * a[0]; ++a; }
1000 return __ocml_rsqrt_f64(r);
1004 double rnorm3d(
double x,
double y,
double z)
1006 return __ocml_rlen3_f64(x, y, z);
1010 double rnorm4d(
double x,
double y,
double z,
double w)
1012 return __ocml_rlen4_f64(x, y, z, w);
1016 double round(
double x) {
return __ocml_round_f64(x); }
1019 double rsqrt(
double x) {
return __ocml_rsqrt_f64(x); }
1022 double scalbln(
double x,
long int n)
1024 return (n < INT_MAX) ? __ocml_scalbn_f64(x, n) : __ocml_scalb_f64(x, n);
1028 double scalbn(
double x,
int n) {
return __ocml_scalbn_f64(x, n); }
1031 __RETURN_TYPE signbit(
double x) {
return __ocml_signbit_f64(x); }
1034 double sin(
double x) {
return __ocml_sin_f64(x); }
1037 void sincos(
double x,
double* sptr,
double* cptr)
1041 __ocml_sincos_f64(x, (__attribute__((address_space(5)))
double*) &tmp);
1046 void sincospi(
double x,
double* sptr,
double* cptr)
1049 *sptr = __ocml_sincospi_f64(
1050 x, (__attribute__((address_space(5)))
double*) &tmp);
1055 double sinh(
double x) {
return __ocml_sinh_f64(x); }
1058 double sinpi(
double x) {
return __ocml_sinpi_f64(x); }
1061 double sqrt(
double x) {
return __ocml_sqrt_f64(x); }
1064 double tan(
double x) {
return __ocml_tan_f64(x); }
1067 double tanh(
double x) {
return __ocml_tanh_f64(x); }
1070 double tgamma(
double x) {
return __ocml_tgamma_f64(x); }
1073 double trunc(
double x) {
return __ocml_trunc_f64(x); }
1076 double y0(
double x) {
return __ocml_y0_f64(x); }
1079 double y1(
double x) {
return __ocml_y1_f64(x); }
1082 double yn(
int n,
double x)
1087 if (n == 0)
return j0f(x);
1088 if (n == 1)
return j1f(x);
1092 for (
int i = 1; i < n; ++i) {
1093 double x2 = (2 * i) / x * x1 - x0;
1102 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1105 double __dadd_rd(
double x,
double y) {
return __ocml_add_rtn_f64(x, y); }
1109 double __dadd_rn(
double x,
double y) {
return x + y; }
1110 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1113 double __dadd_ru(
double x,
double y) {
return __ocml_add_rtp_f64(x, y); }
1116 double __dadd_rz(
double x,
double y) {
return __ocml_add_rtz_f64(x, y); }
1119 double __ddiv_rd(
double x,
double y) {
return __ocml_div_rtn_f64(x, y); }
1123 double __ddiv_rn(
double x,
double y) {
return x / y; }
1124 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1127 double __ddiv_ru(
double x,
double y) {
return __ocml_div_rtp_f64(x, y); }
1130 double __ddiv_rz(
double x,
double y) {
return __ocml_div_rtz_f64(x, y); }
1133 double __dmul_rd(
double x,
double y) {
return __ocml_mul_rtn_f64(x, y); }
1137 double __dmul_rn(
double x,
double y) {
return x * y; }
1138 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1141 double __dmul_ru(
double x,
double y) {
return __ocml_mul_rtp_f64(x, y); }
1144 double __dmul_rz(
double x,
double y) {
return __ocml_mul_rtz_f64(x, y); }
1147 double __drcp_rd(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1151 double __drcp_rn(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1152 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1155 double __drcp_ru(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1158 double __drcp_rz(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1161 double __dsqrt_rd(
double x) {
return __ocml_sqrt_rtn_f64(x); }
1165 double __dsqrt_rn(
double x) {
return __ocml_sqrt_f64(x); }
1166 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1169 double __dsqrt_ru(
double x) {
return __ocml_sqrt_rtp_f64(x); }
1172 double __dsqrt_rz(
double x) {
return __ocml_sqrt_rtz_f64(x); }
1175 double __dsub_rd(
double x,
double y) {
return __ocml_sub_rtn_f64(x, y); }
1179 double __dsub_rn(
double x,
double y) {
return x - y; }
1180 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1183 double __dsub_ru(
double x,
double y) {
return __ocml_sub_rtp_f64(x, y); }
1186 double __dsub_rz(
double x,
double y) {
return __ocml_sub_rtz_f64(x, y); }
1189 double __fma_rd(
double x,
double y,
double z)
1191 return __ocml_fma_rtn_f64(x, y, z);
1196 double __fma_rn(
double x,
double y,
double z)
1198 return __ocml_fma_f64(x, y, z);
1200 #if defined OCML_BASIC_ROUNDED_OPERATIONS 1203 double __fma_ru(
double x,
double y,
double z)
1205 return __ocml_fma_rtp_f64(x, y, z);
1209 double __fma_rz(
double x,
double y,
double z)
1211 return __ocml_fma_rtz_f64(x, y, z);
1222 int sgn = x >> (
sizeof(int) * CHAR_BIT - 1);
1223 return (x ^ sgn) - sgn;
1229 long sgn = x >> (
sizeof(long) * CHAR_BIT - 1);
1230 return (x ^ sgn) - sgn;
1234 long long llabs(
long long x)
1236 long long sgn = x >> (
sizeof(
long long) * CHAR_BIT - 1);
1237 return (x ^ sgn) - sgn;
1240 #if defined(__cplusplus) 1243 long abs(
long x) {
return labs(x); }
1246 long long abs(
long long x) {
return llabs(x); }
1251 inline _Float16 fma(_Float16 x, _Float16 y, _Float16 z) {
1252 return __ocml_fma_f16(x, y, z);
1256 inline float fma(
float x,
float y,
float z) {
1257 return fmaf(x, y, z);
1260 #pragma push_macro("__DEF_FLOAT_FUN") 1261 #pragma push_macro("__DEF_FLOAT_FUN2") 1262 #pragma push_macro("__DEF_FLOAT_FUN2I") 1263 #pragma push_macro("__HIP_OVERLOAD") 1264 #pragma push_macro("__HIP_OVERLOAD2") 1267 template<
bool __B,
class __T =
void>
1277 #define __HIP_OVERLOAD1(__retty, __fn) \ 1278 template <typename __T> \ 1280 typename __hip_enable_if<std::numeric_limits<__T>::is_integer, \ 1283 return ::__fn((double)__x); \ 1289 #define __HIP_OVERLOAD2(__retty, __fn) \ 1290 template <typename __T1, typename __T2> \ 1291 __DEVICE__ typename __hip_enable_if< \ 1292 std::numeric_limits<__T1>::is_specialized && \ 1293 std::numeric_limits<__T2>::is_specialized, \ 1295 __fn(__T1 __x, __T2 __y) { \ 1296 return __fn((double)__x, (double)__y); \ 1300 #define __DEF_FUN1(retty, func) \ 1303 float func(float x) \ 1305 return func##f(x); \ 1307 __HIP_OVERLOAD1(retty, func) 1310 #define __DEF_FUNI(retty, func) \ 1313 retty func(float x) \ 1315 return func##f(x); \ 1317 __HIP_OVERLOAD1(retty, func) 1320 #define __DEF_FUN2(retty, func) \ 1323 float func(float x, float y) \ 1325 return func##f(x, y); \ 1327 __HIP_OVERLOAD2(retty, func) 1329 __DEF_FUN1(
double, acos)
1330 __DEF_FUN1(
double, acosh)
1331 __DEF_FUN1(
double, asin)
1332 __DEF_FUN1(
double, asinh)
1333 __DEF_FUN1(
double, atan)
1334 __DEF_FUN2(
double, atan2);
1335 __DEF_FUN1(
double, atanh)
1336 __DEF_FUN1(
double, cbrt)
1337 __DEF_FUN1(
double, ceil)
1338 __DEF_FUN2(
double, copysign);
1339 __DEF_FUN1(
double, cos)
1340 __DEF_FUN1(
double, cosh)
1341 __DEF_FUN1(
double, erf)
1342 __DEF_FUN1(
double, erfc)
1343 __DEF_FUN1(
double, exp)
1344 __DEF_FUN1(
double, exp2)
1345 __DEF_FUN1(
double, expm1)
1346 __DEF_FUN1(
double, fabs)
1347 __DEF_FUN2(
double, fdim);
1348 __DEF_FUN1(
double, floor)
1349 __DEF_FUN2(
double, fmax);
1350 __DEF_FUN2(
double, fmin);
1351 __DEF_FUN2(
double, fmod);
1353 __DEF_FUN2(
double, hypot);
1354 __DEF_FUNI(
int, ilogb)
1355 __HIP_OVERLOAD1(
bool, isfinite)
1356 __HIP_OVERLOAD2(
bool, isgreater);
1357 __HIP_OVERLOAD2(
bool, isgreaterequal);
1358 __HIP_OVERLOAD1(
bool, isinf);
1359 __HIP_OVERLOAD2(
bool, isless);
1360 __HIP_OVERLOAD2(
bool, islessequal);
1361 __HIP_OVERLOAD2(
bool, islessgreater);
1362 __HIP_OVERLOAD1(
bool, isnan);
1364 __HIP_OVERLOAD2(
bool, isunordered);
1365 __DEF_FUN1(
double, lgamma)
1366 __DEF_FUN1(
double, log)
1367 __DEF_FUN1(
double, log10)
1368 __DEF_FUN1(
double, log1p)
1369 __DEF_FUN1(
double, log2)
1370 __DEF_FUN1(
double, logb)
1371 __DEF_FUNI(
long long, llrint)
1372 __DEF_FUNI(
long long, llround)
1373 __DEF_FUNI(
long, lrint)
1374 __DEF_FUNI(
long, lround)
1375 __DEF_FUN1(
double, nearbyint);
1376 __DEF_FUN2(
double, nextafter);
1377 __DEF_FUN2(
double, pow);
1378 __DEF_FUN2(
double, remainder);
1379 __DEF_FUN1(
double, rint);
1380 __DEF_FUN1(
double, round);
1381 __HIP_OVERLOAD1(
bool, signbit)
1382 __DEF_FUN1(
double, sin)
1383 __DEF_FUN1(
double, sinh)
1384 __DEF_FUN1(
double, sqrt)
1385 __DEF_FUN1(
double, tan)
1386 __DEF_FUN1(
double, tanh)
1387 __DEF_FUN1(
double, tgamma)
1388 __DEF_FUN1(
double, trunc);
1391 #define __DEF_FLOAT_FUN2I(func) \ 1394 float func(float x, int y) \ 1396 return func##f(x, y); \ 1398 __DEF_FLOAT_FUN2I(scalbn)
1402 __DEVICE__
inline static T min(T arg1, T arg2) {
1403 return (arg1 < arg2) ? arg1 : arg2;
1406 __DEVICE__
inline static uint32_t min(uint32_t arg1, int32_t arg2) {
1407 return min(arg1, (uint32_t) arg2);
1428 __DEVICE__
inline static T max(T arg1, T arg2) {
1429 return (arg1 > arg2) ? arg1 : arg2;
1432 __DEVICE__
inline static uint32_t max(uint32_t arg1, int32_t arg2) {
1433 return max(arg1, (uint32_t) arg2);
1435 __DEVICE__
inline static uint32_t max(int32_t arg1, uint32_t arg2) {
1436 return max((uint32_t) arg1, arg2);
1453 __DEVICE__
inline int min(
int arg1,
int arg2) {
1454 return (arg1 < arg2) ? arg1 : arg2;
1456 __DEVICE__
inline int max(
int arg1,
int arg2) {
1457 return (arg1 > arg2) ? arg1 : arg2;
1462 float max(
float x,
float y) {
1468 double max(
double x,
double y) {
1474 float min(
float x,
float y) {
1480 double min(
double x,
double y) {
1484 __HIP_OVERLOAD2(
double, max)
1485 __HIP_OVERLOAD2(
double, min)
1489 __host__ inline static int min(
int arg1,
int arg2) {
1490 return std::min(arg1, arg2);
1493 __host__ inline static int max(
int arg1,
int arg2) {
1494 return std::max(arg1, arg2);
1498 #pragma pop_macro("__DEF_FLOAT_FUN") 1499 #pragma pop_macro("__DEF_FLOAT_FUN2") 1500 #pragma pop_macro("__DEF_FLOAT_FUN2I") 1501 #pragma pop_macro("__HIP_OVERLOAD") 1502 #pragma pop_macro("__HIP_OVERLOAD2") 1503 #pragma pop_macro("__DEVICE__") 1504 #pragma pop_macro("__RETURN_TYPE")
#define __host__
Definition: host_defines.h:41
Contains definitions of APIs for HIP runtime.
Defines the different newt vector types for HIP runtime.
Definition: math_functions.h:1268