33 #include "hip_fp16_math_fwd.h" 39 #include "kalmar_math.h" 42 #pragma push_macro("__DEVICE__") 43 #pragma push_macro("__RETURN_TYPE") 46 #define __DEVICE__ __device__ 47 #define __RETURN_TYPE int 48 #else // to be consistent with __clang_cuda_math_forward_declares 49 #define __DEVICE__ static __device__ 50 #define __RETURN_TYPE bool 55 uint64_t __make_mantissa_base8(
const char* tagp)
61 if (tmp >=
'0' && tmp <=
'7') r = (r * 8u) + tmp -
'0';
72 uint64_t __make_mantissa_base10(
const char* tagp)
78 if (tmp >=
'0' && tmp <=
'9') r = (r * 10u) + tmp -
'0';
89 uint64_t __make_mantissa_base16(
const char* tagp)
95 if (tmp >=
'0' && tmp <=
'9') r = (r * 16u) + tmp -
'0';
96 else if (tmp >=
'a' && tmp <=
'f') r = (r * 16u) + tmp -
'a' + 10;
97 else if (tmp >=
'A' && tmp <=
'F') r = (r * 16u) + tmp -
'A' + 10;
108 uint64_t __make_mantissa(
const char* tagp)
110 if (!tagp)
return 0u;
115 if (*tagp ==
'x' || *tagp ==
'X')
return __make_mantissa_base16(tagp);
116 else return __make_mantissa_base8(tagp);
119 return __make_mantissa_base10(tagp);
125 float abs(
float x) {
return __ocml_fabs_f32(x); }
128 float acosf(
float x) {
return __ocml_acos_f32(x); }
131 float acoshf(
float x) {
return __ocml_acosh_f32(x); }
134 float asinf(
float x) {
return __ocml_asin_f32(x); }
137 float asinhf(
float x) {
return __ocml_asinh_f32(x); }
140 float atan2f(
float x,
float y) {
return __ocml_atan2_f32(x, y); }
143 float atanf(
float x) {
return __ocml_atan_f32(x); }
146 float atanhf(
float x) {
return __ocml_atanh_f32(x); }
149 float cbrtf(
float x) {
return __ocml_cbrt_f32(x); }
152 float ceilf(
float x) {
return __ocml_ceil_f32(x); }
155 float copysignf(
float x,
float y) {
return __ocml_copysign_f32(x, y); }
158 float cosf(
float x) {
return __ocml_cos_f32(x); }
161 float coshf(
float x) {
return __ocml_cosh_f32(x); }
164 float cospif(
float x) {
return __ocml_cospi_f32(x); }
167 float cyl_bessel_i0f(
float x) {
return __ocml_i0_f32(x); }
170 float cyl_bessel_i1f(
float x) {
return __ocml_i1_f32(x); }
173 float erfcf(
float x) {
return __ocml_erfc_f32(x); }
176 float erfcinvf(
float x) {
return __ocml_erfcinv_f32(x); }
179 float erfcxf(
float x) {
return __ocml_erfcx_f32(x); }
182 float erff(
float x) {
return __ocml_erf_f32(x); }
185 float erfinvf(
float x) {
return __ocml_erfinv_f32(x); }
188 float exp10f(
float x) {
return __ocml_exp10_f32(x); }
191 float exp2f(
float x) {
return __ocml_exp2_f32(x); }
194 float expf(
float x) {
return __ocml_exp_f32(x); }
197 float expm1f(
float x) {
return __ocml_expm1_f32(x); }
200 float fabsf(
float x) {
return __ocml_fabs_f32(x); }
203 float fdimf(
float x,
float y) {
return __ocml_fdim_f32(x, y); }
206 float fdividef(
float x,
float y) {
return x / y; }
209 float floorf(
float x) {
return __ocml_floor_f32(x); }
212 float fmaf(
float x,
float y,
float z) {
return __ocml_fma_f32(x, y, z); }
215 float fmaxf(
float x,
float y) {
return __ocml_fmax_f32(x, y); }
218 float fminf(
float x,
float y) {
return __ocml_fmin_f32(x, y); }
221 float fmodf(
float x,
float y) {
return __ocml_fmod_f32(x, y); }
224 float frexpf(
float x,
int* nptr)
228 __ocml_frexp_f32(x, (__attribute__((address_space(5)))
int*) &tmp);
235 float hypotf(
float x,
float y) {
return __ocml_hypot_f32(x, y); }
238 int ilogbf(
float x) {
return __ocml_ilogb_f32(x); }
241 __RETURN_TYPE isfinite(
float x) {
return __ocml_isfinite_f32(x); }
244 __RETURN_TYPE isinf(
float x) {
return __ocml_isinf_f32(x); }
247 __RETURN_TYPE isnan(
float x) {
return __ocml_isnan_f32(x); }
250 float j0f(
float x) {
return __ocml_j0_f32(x); }
253 float j1f(
float x) {
return __ocml_j1_f32(x); }
256 float jnf(
int n,
float x)
260 if (n == 0)
return j0f(x);
261 if (n == 1)
return j1f(x);
265 for (
int i = 1; i < n; ++i) {
266 float x2 = (2 * i) / x * x1 - x0;
275 float ldexpf(
float x,
int e) {
return __ocml_ldexp_f32(x, e); }
278 float lgammaf(
float x) {
return __ocml_lgamma_f32(x); }
281 long long int llrintf(
float x) {
return __ocml_rint_f32(x); }
284 long long int llroundf(
float x) {
return __ocml_round_f32(x); }
287 float log10f(
float x) {
return __ocml_log10_f32(x); }
290 float log1pf(
float x) {
return __ocml_log1p_f32(x); }
293 float log2f(
float x) {
return __ocml_log2_f32(x); }
296 float logbf(
float x) {
return __ocml_logb_f32(x); }
299 float logf(
float x) {
return __ocml_log_f32(x); }
302 long int lrintf(
float x) {
return __ocml_rint_f32(x); }
305 long int lroundf(
float x) {
return __ocml_round_f32(x); }
308 float modff(
float x,
float* iptr)
312 __ocml_modf_f32(x, (__attribute__((address_space(5)))
float*) &tmp);
319 float nanf(
const char* tagp)
324 uint32_t mantissa : 22;
326 uint32_t exponent : 8;
330 static_assert(
sizeof(
float) ==
sizeof(ieee_float),
"");
334 tmp.bits.exponent = ~0u;
336 tmp.bits.mantissa = __make_mantissa(tagp);
342 float nearbyintf(
float x) {
return __ocml_nearbyint_f32(x); }
345 float nextafterf(
float x,
float y) {
return __ocml_nextafter_f32(x, y); }
348 float norm3df(
float x,
float y,
float z) {
return __ocml_len3_f32(x, y, z); }
351 float norm4df(
float x,
float y,
float z,
float w)
353 return __ocml_len4_f32(x, y, z, w);
357 float normcdff(
float x) {
return __ocml_ncdf_f32(x); }
360 float normcdfinvf(
float x) {
return __ocml_ncdfinv_f32(x); }
363 float normf(
int dim,
const float* a)
366 while (dim--) { r += a[0] * a[0]; ++a; }
368 return __ocml_sqrt_f32(r);
372 float powf(
float x,
float y) {
return __ocml_pow_f32(x, y); }
375 float rcbrtf(
float x) {
return __ocml_rcbrt_f32(x); }
378 float remainderf(
float x,
float y) {
return __ocml_remainder_f32(x, y); }
381 float remquof(
float x,
float y,
int* quo)
385 __ocml_remquo_f32(x, y, (__attribute__((address_space(5)))
int*) &tmp);
392 float rhypotf(
float x,
float y) {
return __ocml_rhypot_f32(x, y); }
395 float rintf(
float x) {
return __ocml_rint_f32(x); }
398 float rnorm3df(
float x,
float y,
float z)
400 return __ocml_rlen3_f32(x, y, z);
405 float rnorm4df(
float x,
float y,
float z,
float w)
407 return __ocml_rlen4_f32(x, y, z, w);
411 float rnormf(
int dim,
const float* a)
414 while (dim--) { r += a[0] * a[0]; ++a; }
416 return __ocml_rsqrt_f32(r);
420 float roundf(
float x) {
return __ocml_round_f32(x); }
423 float rsqrtf(
float x) {
return __ocml_rsqrt_f32(x); }
426 float scalblnf(
float x,
long int n)
428 return (n < INT_MAX) ? __ocml_scalbn_f32(x, n) : __ocml_scalb_f32(x, n);
432 float scalbnf(
float x,
int n) {
return __ocml_scalbn_f32(x, n); }
435 __RETURN_TYPE signbit(
float x) {
return __ocml_signbit_f32(x); }
438 void sincosf(
float x,
float* sptr,
float* cptr)
443 __ocml_sincos_f32(x, (__attribute__((address_space(5)))
float*) &tmp);
448 void sincospif(
float x,
float* sptr,
float* cptr)
453 __ocml_sincospi_f32(x, (__attribute__((address_space(5)))
float*) &tmp);
458 float sinf(
float x) {
return __ocml_sin_f32(x); }
461 float sinhf(
float x) {
return __ocml_sinh_f32(x); }
464 float sinpif(
float x) {
return __ocml_sinpi_f32(x); }
467 float sqrtf(
float x) {
return __ocml_sqrt_f32(x); }
470 float tanf(
float x) {
return __ocml_tan_f32(x); }
473 float tanhf(
float x) {
return __ocml_tanh_f32(x); }
476 float tgammaf(
float x) {
return __ocml_tgamma_f32(x); }
479 float truncf(
float x) {
return __ocml_trunc_f32(x); }
482 float y0f(
float x) {
return __ocml_y0_f32(x); }
485 float y1f(
float x) {
return __ocml_y1_f32(x); }
488 float ynf(
int n,
float x)
493 if (n == 0)
return y0f(x);
494 if (n == 1)
return y1f(x);
498 for (
int i = 1; i < n; ++i) {
499 float x2 = (2 * i) / x * x1 - x0;
510 float __cosf(
float x) {
return __llvm_amdgcn_cos_f32(x); }
513 float __exp10f(
float x) {
return __ocml_exp10_f32(x); }
516 float __expf(
float x) {
return __ocml_exp_f32(x); }
519 float __fadd_rd(
float x,
float y) {
return __ocml_add_rtp_f32(x, y); }
522 float __fadd_rn(
float x,
float y) {
return __ocml_add_rte_f32(x, y); }
525 float __fadd_ru(
float x,
float y) {
return __ocml_add_rtn_f32(x, y); }
528 float __fadd_rz(
float x,
float y) {
return __ocml_add_rtz_f32(x, y); }
531 float __fdiv_rd(
float x,
float y) {
return x / y; }
534 float __fdiv_rn(
float x,
float y) {
return x / y; }
537 float __fdiv_ru(
float x,
float y) {
return x / y; }
540 float __fdiv_rz(
float x,
float y) {
return x / y; }
543 float __fdividef(
float x,
float y) {
return x / y; }
546 float __fmaf_rd(
float x,
float y,
float z)
548 return __ocml_fma_rtp_f32(x, y, z);
552 float __fmaf_rn(
float x,
float y,
float z)
554 return __ocml_fma_rte_f32(x, y, z);
558 float __fmaf_ru(
float x,
float y,
float z)
560 return __ocml_fma_rtn_f32(x, y, z);
564 float __fmaf_rz(
float x,
float y,
float z)
566 return __ocml_fma_rtz_f32(x, y, z);
570 float __fmul_rd(
float x,
float y) {
return __ocml_mul_rtp_f32(x, y); }
573 float __fmul_rn(
float x,
float y) {
return __ocml_mul_rte_f32(x, y); }
576 float __fmul_ru(
float x,
float y) {
return __ocml_mul_rtn_f32(x, y); }
579 float __fmul_rz(
float x,
float y) {
return __ocml_mul_rtz_f32(x, y); }
582 float __frcp_rd(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
585 float __frcp_rn(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
588 float __frcp_ru(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
591 float __frcp_rz(
float x) {
return __llvm_amdgcn_rcp_f32(x); }
594 float __frsqrt_rn(
float x) {
return __llvm_amdgcn_rsq_f32(x); }
597 float __fsqrt_rd(
float x) {
return __ocml_sqrt_rtp_f32(x); }
600 float __fsqrt_rn(
float x) {
return __ocml_sqrt_rte_f32(x); }
603 float __fsqrt_ru(
float x) {
return __ocml_sqrt_rtn_f32(x); }
606 float __fsqrt_rz(
float x) {
return __ocml_sqrt_rtz_f32(x); }
609 float __fsub_rd(
float x,
float y) {
return __ocml_sub_rtp_f32(x, y); }
612 float __fsub_rn(
float x,
float y) {
return __ocml_sub_rte_f32(x, y); }
615 float __fsub_ru(
float x,
float y) {
return __ocml_sub_rtn_f32(x, y); }
618 float __fsub_rz(
float x,
float y) {
return __ocml_sub_rtz_f32(x, y); }
621 float __log10f(
float x) {
return __ocml_log10_f32(x); }
624 float __log2f(
float x) {
return __ocml_log2_f32(x); }
627 float __logf(
float x) {
return __ocml_log_f32(x); }
630 float __powf(
float x,
float y) {
return __ocml_pow_f32(x, y); }
633 float __saturatef(
float x) {
return (x < 0) ? 0 : ((x > 1) ? 1 : x); }
636 void __sincosf(
float x,
float* sptr,
float* cptr)
641 __ocml_sincos_f32(x, (__attribute__((address_space(5)))
float*) &tmp);
646 float __sinf(
float x) {
return __llvm_amdgcn_sin_f32(x); }
649 float __tanf(
float x) {
return __ocml_tan_f32(x); }
656 double abs(
double x) {
return __ocml_fabs_f64(x); }
659 double acos(
double x) {
return __ocml_acos_f64(x); }
662 double acosh(
double x) {
return __ocml_acosh_f64(x); }
665 double asin(
double x) {
return __ocml_asin_f64(x); }
668 double asinh(
double x) {
return __ocml_asinh_f64(x); }
671 double atan(
double x) {
return __ocml_atan_f64(x); }
674 double atan2(
double x,
double y) {
return __ocml_atan2_f64(x, y); }
677 double atanh(
double x) {
return __ocml_atanh_f64(x); }
680 double cbrt(
double x) {
return __ocml_cbrt_f64(x); }
683 double ceil(
double x) {
return __ocml_ceil_f64(x); }
686 double copysign(
double x,
double y) {
return __ocml_copysign_f64(x, y); }
689 double cos(
double x) {
return __ocml_cos_f64(x); }
692 double cosh(
double x) {
return __ocml_cosh_f64(x); }
695 double cospi(
double x) {
return __ocml_cospi_f64(x); }
698 double cyl_bessel_i0(
double x) {
return __ocml_i0_f64(x); }
701 double cyl_bessel_i1(
double x) {
return __ocml_i1_f64(x); }
704 double erf(
double x) {
return __ocml_erf_f64(x); }
707 double erfc(
double x) {
return __ocml_erfc_f64(x); }
710 double erfcinv(
double x) {
return __ocml_erfcinv_f64(x); }
713 double erfcx(
double x) {
return __ocml_erfcx_f64(x); }
716 double erfinv(
double x) {
return __ocml_erfinv_f64(x); }
719 double exp(
double x) {
return __ocml_exp_f64(x); }
722 double exp10(
double x) {
return __ocml_exp10_f64(x); }
725 double exp2(
double x) {
return __ocml_exp2_f64(x); }
728 double expm1(
double x) {
return __ocml_expm1_f64(x); }
731 double fabs(
double x) {
return __ocml_fabs_f64(x); }
734 double fdim(
double x,
double y) {
return __ocml_fdim_f64(x, y); }
737 double floor(
double x) {
return __ocml_floor_f64(x); }
740 double fma(
double x,
double y,
double z) {
return __ocml_fma_f64(x, y, z); }
743 double fmax(
double x,
double y) {
return __ocml_fmax_f64(x, y); }
746 double fmin(
double x,
double y) {
return __ocml_fmin_f64(x, y); }
749 double fmod(
double x,
double y) {
return __ocml_fmod_f64(x, y); }
752 double frexp(
double x,
int* nptr)
756 __ocml_frexp_f64(x, (__attribute__((address_space(5)))
int*) &tmp);
763 double hypot(
double x,
double y) {
return __ocml_hypot_f64(x, y); }
766 int ilogb(
double x) {
return __ocml_ilogb_f64(x); }
769 __RETURN_TYPE isfinite(
double x) {
return __ocml_isfinite_f64(x); }
772 __RETURN_TYPE isinf(
double x) {
return __ocml_isinf_f64(x); }
775 __RETURN_TYPE isnan(
double x) {
return __ocml_isnan_f64(x); }
778 double j0(
double x) {
return __ocml_j0_f64(x); }
781 double j1(
double x) {
return __ocml_j1_f64(x); }
784 double jn(
int n,
double x)
789 if (n == 0)
return j0f(x);
790 if (n == 1)
return j1f(x);
794 for (
int i = 1; i < n; ++i) {
795 double x2 = (2 * i) / x * x1 - x0;
804 double ldexp(
double x,
int e) {
return __ocml_ldexp_f64(x, e); }
807 double lgamma(
double x) {
return __ocml_lgamma_f64(x); }
810 long long int llrint(
double x) {
return __ocml_rint_f64(x); }
813 long long int llround(
double x) {
return __ocml_round_f64(x); }
816 double log(
double x) {
return __ocml_log_f64(x); }
819 double log10(
double x) {
return __ocml_log10_f64(x); }
822 double log1p(
double x) {
return __ocml_log1p_f64(x); }
825 double log2(
double x) {
return __ocml_log2_f64(x); }
828 double logb(
double x) {
return __ocml_logb_f64(x); }
831 long int lrint(
double x) {
return __ocml_rint_f64(x); }
834 long int lround(
double x) {
return __ocml_round_f64(x); }
837 double modf(
double x,
double* iptr)
841 __ocml_modf_f64(x, (__attribute__((address_space(5)))
double*) &tmp);
848 double nan(
const char* tagp)
853 uint64_t mantissa : 51;
855 uint32_t exponent : 11;
859 static_assert(
sizeof(
double) ==
sizeof(ieee_double),
"");
863 tmp.bits.exponent = ~0u;
865 tmp.bits.mantissa = __make_mantissa(tagp);
871 double nearbyint(
double x) {
return __ocml_nearbyint_f64(x); }
874 double nextafter(
double x,
double y) {
return __ocml_nextafter_f64(x, y); }
877 double norm(
int dim,
const double* a)
880 while (dim--) { r += a[0] * a[0]; ++a; }
882 return __ocml_sqrt_f64(r);
886 double norm3d(
double x,
double y,
double z)
888 return __ocml_len3_f64(x, y, z);
892 double norm4d(
double x,
double y,
double z,
double w)
894 return __ocml_len4_f64(x, y, z, w);
898 double normcdf(
double x) {
return __ocml_ncdf_f64(x); }
901 double normcdfinv(
double x) {
return __ocml_ncdfinv_f64(x); }
904 double pow(
double x,
double y) {
return __ocml_pow_f64(x, y); }
907 double rcbrt(
double x) {
return __ocml_rcbrt_f64(x); }
910 double remainder(
double x,
double y) {
return __ocml_remainder_f64(x, y); }
913 double remquo(
double x,
double y,
int* quo)
917 __ocml_remquo_f64(x, y, (__attribute__((address_space(5)))
int*) &tmp);
924 double rhypot(
double x,
double y) {
return __ocml_rhypot_f64(x, y); }
927 double rint(
double x) {
return __ocml_rint_f64(x); }
930 double rnorm(
int dim,
const double* a)
933 while (dim--) { r += a[0] * a[0]; ++a; }
935 return __ocml_rsqrt_f64(r);
939 double rnorm3d(
double x,
double y,
double z)
941 return __ocml_rlen3_f64(x, y, z);
945 double rnorm4d(
double x,
double y,
double z,
double w)
947 return __ocml_rlen4_f64(x, y, z, w);
951 double round(
double x) {
return __ocml_round_f64(x); }
954 double rsqrt(
double x) {
return __ocml_rsqrt_f64(x); }
957 double scalbln(
double x,
long int n)
959 return (n < INT_MAX) ? __ocml_scalbn_f64(x, n) : __ocml_scalb_f64(x, n);
963 double scalbn(
double x,
int n) {
return __ocml_scalbn_f64(x, n); }
966 __RETURN_TYPE signbit(
double x) {
return __ocml_signbit_f64(x); }
969 double sin(
double x) {
return __ocml_sin_f64(x); }
972 void sincos(
double x,
double* sptr,
double* cptr)
976 __ocml_sincos_f64(x, (__attribute__((address_space(5)))
double*) &tmp);
981 void sincospi(
double x,
double* sptr,
double* cptr)
984 *sptr = __ocml_sincospi_f64(
985 x, (__attribute__((address_space(5)))
double*) &tmp);
990 double sinh(
double x) {
return __ocml_sinh_f64(x); }
993 double sinpi(
double x) {
return __ocml_sinpi_f64(x); }
996 double sqrt(
double x) {
return __ocml_sqrt_f64(x); }
999 double tan(
double x) {
return __ocml_tan_f64(x); }
1002 double tanh(
double x) {
return __ocml_tanh_f64(x); }
1005 double tgamma(
double x) {
return __ocml_tgamma_f64(x); }
1008 double trunc(
double x) {
return __ocml_trunc_f64(x); }
1011 double y0(
double x) {
return __ocml_y0_f64(x); }
1014 double y1(
double x) {
return __ocml_y1_f64(x); }
1017 double yn(
int n,
double x)
1022 if (n == 0)
return j0f(x);
1023 if (n == 1)
return j1f(x);
1027 for (
int i = 1; i < n; ++i) {
1028 double x2 = (2 * i) / x * x1 - x0;
1039 double __dadd_rd(
double x,
double y) {
return __ocml_add_rtp_f64(x, y); }
1042 double __dadd_rn(
double x,
double y) {
return __ocml_add_rte_f64(x, y); }
1045 double __dadd_ru(
double x,
double y) {
return __ocml_add_rtn_f64(x, y); }
1048 double __dadd_rz(
double x,
double y) {
return __ocml_add_rtz_f64(x, y); }
1051 double __ddiv_rd(
double x,
double y) {
return x / y; }
1054 double __ddiv_rn(
double x,
double y) {
return x / y; }
1057 double __ddiv_ru(
double x,
double y) {
return x / y; }
1060 double __ddiv_rz(
double x,
double y) {
return x / y; }
1063 double __dmul_rd(
double x,
double y) {
return __ocml_mul_rtp_f64(x, y); }
1066 double __dmul_rn(
double x,
double y) {
return __ocml_mul_rte_f64(x, y); }
1069 double __dmul_ru(
double x,
double y) {
return __ocml_mul_rtn_f64(x, y); }
1072 double __dmul_rz(
double x,
double y) {
return __ocml_mul_rtz_f64(x, y); }
1075 double __drcp_rd(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1078 double __drcp_rn(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1081 double __drcp_ru(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1084 double __drcp_rz(
double x) {
return __llvm_amdgcn_rcp_f64(x); }
1087 double __dsqrt_rd(
double x) {
return __ocml_sqrt_rtp_f64(x); }
1090 double __dsqrt_rn(
double x) {
return __ocml_sqrt_rte_f64(x); }
1093 double __dsqrt_ru(
double x) {
return __ocml_sqrt_rtn_f64(x); }
1096 double __dsqrt_rz(
double x) {
return __ocml_sqrt_rtz_f64(x); }
1099 double __dsub_rd(
double x,
double y) {
return __ocml_sub_rtp_f64(x, y); }
1102 double __dsub_rn(
double x,
double y) {
return __ocml_sub_rte_f64(x, y); }
1105 double __dsub_ru(
double x,
double y) {
return __ocml_sub_rtn_f64(x, y); }
1108 double __dsub_rz(
double x,
double y) {
return __ocml_sub_rtz_f64(x, y); }
1111 double __fma_rd(
double x,
double y,
double z)
1113 return __ocml_fma_rtp_f64(x, y, z);
1117 double __fma_rn(
double x,
double y,
double z)
1119 return __ocml_fma_rte_f64(x, y, z);
1123 double __fma_ru(
double x,
double y,
double z)
1125 return __ocml_fma_rtn_f64(x, y, z);
1129 double __fma_rz(
double x,
double y,
double z)
1131 return __ocml_fma_rtz_f64(x, y, z);
1141 int sgn = x >> (
sizeof(int) * CHAR_BIT - 1);
1142 return (x ^ sgn) - sgn;
1148 long sgn = x >> (
sizeof(long) * CHAR_BIT - 1);
1149 return (x ^ sgn) - sgn;
1153 long long llabs(
long long x)
1155 long long sgn = x >> (
sizeof(
long long) * CHAR_BIT - 1);
1156 return (x ^ sgn) - sgn;
1159 #if defined(__cplusplus) 1162 long abs(
long x) {
return labs(x); }
1165 long long abs(
long long x) {
return llabs(x); }
1170 inline _Float16 fma(_Float16 x, _Float16 y, _Float16 z) {
1171 return __ocml_fma_f16(x, y, z);
1175 inline float fma(
float x,
float y,
float z) {
1176 return fmaf(x, y, z);
1179 #pragma push_macro("__DEF_FLOAT_FUN") 1180 #pragma push_macro("__DEF_FLOAT_FUN2") 1181 #pragma push_macro("__DEF_FLOAT_FUN2I") 1182 #pragma push_macro("__HIP_OVERLOAD") 1183 #pragma push_macro("__HIP_OVERLOAD2") 1186 template<
bool __B,
class __T =
void>
1196 #define __HIP_OVERLOAD1(__retty, __fn) \ 1197 template <typename __T> \ 1199 typename __hip_enable_if<std::numeric_limits<__T>::is_integer, \ 1202 return ::__fn((double)__x); \ 1208 #define __HIP_OVERLOAD2(__retty, __fn) \ 1209 template <typename __T1, typename __T2> \ 1210 __DEVICE__ typename __hip_enable_if< \ 1211 std::numeric_limits<__T1>::is_specialized && \ 1212 std::numeric_limits<__T2>::is_specialized, \ 1214 __fn(__T1 __x, __T2 __y) { \ 1215 return __fn((double)__x, (double)__y); \ 1219 #define __DEF_FUN1(retty, func) \ 1222 float func(float x) \ 1224 return func##f(x); \ 1226 __HIP_OVERLOAD1(retty, func) 1229 #define __DEF_FUNI(retty, func) \ 1232 retty func(float x) \ 1234 return func##f(x); \ 1236 __HIP_OVERLOAD1(retty, func) 1239 #define __DEF_FUN2(retty, func) \ 1242 float func(float x, float y) \ 1244 return func##f(x, y); \ 1246 __HIP_OVERLOAD2(retty, func) 1248 __DEF_FUN1(
double, acos)
1249 __DEF_FUN1(
double, acosh)
1250 __DEF_FUN1(
double, asin)
1251 __DEF_FUN1(
double, asinh)
1252 __DEF_FUN1(
double, atan)
1253 __DEF_FUN2(
double, atan2);
1254 __DEF_FUN1(
double, atanh)
1255 __DEF_FUN1(
double, cbrt)
1256 __DEF_FUN1(
double, ceil)
1257 __DEF_FUN2(
double, copysign);
1258 __DEF_FUN1(
double, cos)
1259 __DEF_FUN1(
double, cosh)
1260 __DEF_FUN1(
double, erf)
1261 __DEF_FUN1(
double, erfc)
1262 __DEF_FUN1(
double, exp)
1263 __DEF_FUN1(
double, exp2)
1264 __DEF_FUN1(
double, expm1)
1265 __DEF_FUN1(
double, fabs)
1266 __DEF_FUN2(
double, fdim);
1267 __DEF_FUN1(
double, floor)
1268 __DEF_FUN2(
double, fmax);
1269 __DEF_FUN2(
double, fmin);
1270 __DEF_FUN2(
double, fmod);
1272 __DEF_FUN2(
double, hypot);
1273 __DEF_FUNI(
int, ilogb)
1274 __HIP_OVERLOAD1(
bool, isfinite)
1275 __HIP_OVERLOAD2(
bool, isgreater);
1276 __HIP_OVERLOAD2(
bool, isgreaterequal);
1277 __HIP_OVERLOAD1(
bool, isinf);
1278 __HIP_OVERLOAD2(
bool, isless);
1279 __HIP_OVERLOAD2(
bool, islessequal);
1280 __HIP_OVERLOAD2(
bool, islessgreater);
1281 __HIP_OVERLOAD1(
bool, isnan);
1283 __HIP_OVERLOAD2(
bool, isunordered);
1284 __DEF_FUN1(
double, lgamma)
1285 __DEF_FUN1(
double, log)
1286 __DEF_FUN1(
double, log10)
1287 __DEF_FUN1(
double, log1p)
1288 __DEF_FUN1(
double, log2)
1289 __DEF_FUN1(
double, logb)
1290 __DEF_FUNI(
long long, llrint)
1291 __DEF_FUNI(
long long, llround)
1292 __DEF_FUNI(
long, lrint)
1293 __DEF_FUNI(
long, lround)
1294 __DEF_FUN1(
double, nearbyint);
1295 __DEF_FUN2(
double, nextafter);
1296 __DEF_FUN2(
double, pow);
1297 __DEF_FUN2(
double, remainder);
1298 __DEF_FUN1(
double, rint);
1299 __DEF_FUN1(
double, round);
1300 __HIP_OVERLOAD1(
bool, signbit)
1301 __DEF_FUN1(
double, sin)
1302 __DEF_FUN1(
double, sinh)
1303 __DEF_FUN1(
double, sqrt)
1304 __DEF_FUN1(
double, tan)
1305 __DEF_FUN1(
double, tanh)
1306 __DEF_FUN1(
double, tgamma)
1307 __DEF_FUN1(
double, trunc);
1310 #define __DEF_FLOAT_FUN2I(func) \ 1313 float func(float x, int y) \ 1315 return func##f(x, y); \ 1317 __DEF_FLOAT_FUN2I(scalbn)
1321 __DEVICE__
inline static T min(T arg1, T arg2) {
1322 return (arg1 < arg2) ? arg1 : arg2;
1325 __DEVICE__
inline static uint32_t min(uint32_t arg1, int32_t arg2) {
1326 return min(arg1, (uint32_t) arg2);
1347 __DEVICE__
inline static T max(T arg1, T arg2) {
1348 return (arg1 > arg2) ? arg1 : arg2;
1351 __DEVICE__
inline static uint32_t max(uint32_t arg1, int32_t arg2) {
1352 return max(arg1, (uint32_t) arg2);
1354 __DEVICE__
inline static uint32_t max(int32_t arg1, uint32_t arg2) {
1355 return max((uint32_t) arg1, arg2);
1372 __DEVICE__
inline int min(
int arg1,
int arg2) {
1373 return (arg1 < arg2) ? arg1 : arg2;
1375 __DEVICE__
inline int max(
int arg1,
int arg2) {
1376 return (arg1 > arg2) ? arg1 : arg2;
1381 float max(
float x,
float y) {
1387 double max(
double x,
double y) {
1393 float min(
float x,
float y) {
1399 double min(
double x,
double y) {
1403 __HIP_OVERLOAD2(
double, max)
1404 __HIP_OVERLOAD2(
double, min)
1408 __host__ inline static int min(
int arg1,
int arg2) {
1409 return std::min(arg1, arg2);
1412 __host__ inline static int max(
int arg1,
int arg2) {
1413 return std::max(arg1, arg2);
1417 #pragma pop_macro("__DEF_FLOAT_FUN") 1418 #pragma pop_macro("__DEF_FLOAT_FUN2") 1419 #pragma pop_macro("__DEF_FLOAT_FUN2I") 1420 #pragma pop_macro("__HIP_OVERLOAD") 1421 #pragma pop_macro("__HIP_OVERLOAD2") 1422 #pragma pop_macro("__DEVICE__") 1423 #pragma pop_macro("__RETURN_TYPE")
#define __host__
Definition: host_defines.h:41
Contains definitions of APIs for HIP runtime.
Definition: math_functions.h:1187