6 #ifndef __HIP_DEVICE_COMPILE__
17 #if CK_WORKAROUND_SWDEV_383542
18 extern "C" __device__
float __ocml_native_recip_f32(
float);
22 #ifndef CK_CODE_GEN_RTC
23 static inline __host__
float abs(
float x) {
return std::abs(x); };
25 static inline __host__
double abs(
double x) {
return std::abs(x); };
31 return (x ^ sgn) - sgn;
34 static inline __host__ int32_t abs(int32_t x)
36 int32_t sgn = x >> (32 - 1);
38 return (x ^ sgn) - sgn;
43 uint16_t xx = ck::bit_cast<uint16_t>(x);
45 uint16_t abs_xx = xx & 0x7fff;
47 half_t abs_x = ck::bit_cast<half_t>(abs_xx);
52 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
56 return (x ^ sgn) - sgn;
60 static inline __host__
bool isnan(
float x) {
return std::isnan(x); };
62 static inline __host__
bool isnan(
double x) {
return std::isnan(x); };
64 static inline __host__
bool isnan(
int8_t x)
70 static inline __host__
bool isnan(int32_t x)
76 static inline __host__
bool isnan(
half_t x)
78 uint16_t xx = ck::bit_cast<uint16_t>(x);
80 return (xx & 0x7FFF) > 0x7C00;
83 static inline __host__
bool isnan(
f8_t x) {
return ck::fp8_is_nan(x); };
85 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
86 static inline __host__
bool isnan(
int4_t x)
95 return static_cast<half_t>(std::sqrt(
static_cast<float>(x)));
98 static inline __host__
float sqrt(
float x) {
return std::sqrt(x); };
100 static inline __host__
double sqrt(
double x) {
return std::sqrt(x); };
102 template <
typename T>
105 return ck::type_convert<T>(std::tanhf(ck::type_convert<float>(x)));
111 return std::tanhf(x);
120 template <
typename T>
123 return ck::type_convert<T>(std::acosf(ck::type_convert<float>(x)));
129 return std::acosf(x);
138 template <
typename T>
139 inline __host__ T
neg(T x)
141 return ck::type_convert<T>(-(ck::type_convert<float>(x)));
168 template <
typename T>
171 return ck::type_convert<T>(std::atanf(ck::type_convert<float>(x)));
177 return std::atanf(x);
186 template <
typename T>
187 inline __host__ T
sin(T x)
189 return ck::type_convert<T>(std::sinf(ck::type_convert<float>(x)));
204 template <
typename T>
207 return ck::type_convert<T>(std::asinf(ck::type_convert<float>(x)));
213 return std::asinf(x);
222 template <
typename T>
225 return ck::type_convert<T>(std::asinhf(ck::type_convert<float>(x)));
231 return std::asinhf(x);
240 template <
typename T>
241 inline __host__ T
cos(T x)
243 return ck::type_convert<T>(std::cosf(ck::type_convert<float>(x)));
258 template <
typename T>
261 return ck::type_convert<T>(std::acoshf(ck::type_convert<float>(x)));
267 return std::acoshf(x);
276 template <
typename T>
277 inline __host__ T
tan(T x)
279 return ck::type_convert<T>(std::tanf(ck::type_convert<float>(x)));
294 template <
typename T>
297 return ck::type_convert<T>(std::atanhf(ck::type_convert<float>(x)));
303 return std::atanhf(x);
312 template <
typename T>
315 return ck::type_convert<T>(std::sinhf(ck::type_convert<float>(x)));
321 return std::sinhf(x);
330 template <
typename T>
333 return ck::type_convert<T>(std::ceilf(ck::type_convert<float>(x)));
339 return std::ceilf(x);
348 template <
typename T>
351 return ck::type_convert<T>(std::coshf(ck::type_convert<float>(x)));
357 return std::coshf(x);
366 template <
typename T>
369 return ck::type_convert<T>(std::floorf(ck::type_convert<float>(x)));
375 return std::floorf(x);
384 template <
typename T>
385 inline __host__ T
rcp(T x)
387 return ck::type_convert<T>(1.f / ck::type_convert<float>(x));
390 template <
typename T>
391 inline __host__ T
exp(T x)
393 return ck::type_convert<T>(std::expf(ck::type_convert<float>(x)));
408 template <
typename T>
409 inline __host__ T
log(T x)
411 return ck::type_convert<T>(std::logf(ck::type_convert<float>(x)));
426 template <
typename T>
427 inline __host__ T
pow(T x, T gamma)
429 return ck::type_convert<T>(
430 std::powf(ck::type_convert<float>(x), ck::type_convert<float>(gamma)));
436 return std::powf(x, gamma);
445 template <
typename T>
448 return ck::type_convert<T>(std::expm1f(ck::type_convert<float>(x)));
454 return std::expm1f(x);
465 static inline __device__
float abs(
float x) { return ::abs(x); };
467 static inline __device__
double abs(
double x) { return ::abs(x); };
471 int8_t sgn = x >> (8 - 1);
473 return (x ^ sgn) - sgn;
476 static inline __device__ int32_t abs(int32_t x)
478 int32_t sgn = x >> (32 - 1);
480 return (x ^ sgn) - sgn;
483 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
486 int4_t sgn = x >> (4 - 1);
488 return (x ^ sgn) - sgn;
494 uint16_t xx = ck::bit_cast<uint16_t>(x);
496 uint16_t abs_xx = xx & 0x7fff;
498 half_t abs_x = ck::bit_cast<half_t>(abs_xx);
503 static inline __device__
bool isnan(
float x) { return ::isnan(x); };
505 static inline __device__
bool isnan(
double x) { return ::isnan(x); };
507 static inline __device__
bool isnan(
int8_t x)
513 static inline __device__
bool isnan(int32_t x)
519 #ifdef CK_EXPERIMENTAL_BIT_INT_EXTENSION_INT4
520 static inline __device__
bool isnan(
int4_t x)
527 static inline __device__
bool isnan(
half_t x)
529 uint16_t xx = ck::bit_cast<uint16_t>(x);
531 return (xx & 0x7FFF) > 0x7C00;
534 static inline __device__
bool isnan(
f8_t x) {
return ck::fp8_is_nan(x); };
538 return static_cast<half_t>(__builtin_amdgcn_sqrtf(
static_cast<float>(x)));
541 static inline __device__
float sqrt(
float x) {
return __builtin_amdgcn_sqrtf(x); };
543 static inline __device__
double sqrt(
double x) {
return __builtin_amdgcn_sqrt(x); };
545 template <
typename T>
548 return ck::type_convert<T>(::tanhf(ck::type_convert<float>(x)));
563 template <
typename T>
566 return ck::type_convert<T>(::acosf(ck::type_convert<float>(x)));
581 template <
typename T>
582 inline __device__ T
neg(T x)
584 return ck::type_convert<T>(-(ck::type_convert<float>(x)));
614 return __hneg(
static_cast<__half
>(x));
617 template <
typename T>
620 return ck::type_convert<T>(::atanf(ck::type_convert<float>(x)));
635 template <
typename T>
636 inline __device__ T
sin(T x)
638 return ck::type_convert<T>(::sinf(ck::type_convert<float>(x)));
656 return hsin(
static_cast<__half
>(x));
659 template <
typename T>
662 return ck::type_convert<T>(::asinf(ck::type_convert<float>(x)));
677 template <
typename T>
680 return ck::type_convert<T>(::asinhf(ck::type_convert<float>(x)));
695 template <
typename T>
698 return ck::type_convert<T>(::acoshf(ck::type_convert<float>(x)));
713 template <
typename T>
714 inline __device__ T
tan(T x)
716 return ck::type_convert<T>(::tanf(ck::type_convert<float>(x)));
731 template <
typename T>
734 return ck::type_convert<T>(::atanhf(ck::type_convert<float>(x)));
749 template <
typename T>
752 return ck::type_convert<T>(::sinhf(ck::type_convert<float>(x)));
767 template <
typename T>
770 return ck::type_convert<T>(::ceilf(ck::type_convert<float>(x)));
788 return hceil(
static_cast<__half
>(x));
791 template <
typename T>
794 return ck::type_convert<T>(::coshf(ck::type_convert<float>(x)));
809 template <
typename T>
812 return ck::type_convert<T>(::floorf(ck::type_convert<float>(x)));
830 return hfloor(
static_cast<__half
>(x));
833 template <
typename T>
834 inline __device__ T
rcp(T x)
836 #if !CK_WORKAROUND_SWDEV_383542
839 return __ocml_native_recip_f32(x);
843 template <
typename T>
844 inline __device__ T
exp(T x)
846 return ck::type_convert<T>(__ocml_exp_f32(ck::type_convert<float>(x)));
852 return hexp(
static_cast<__half
>(x));
858 return __ocml_exp_f32(x);
867 template <
typename T>
868 inline __device__ T
log(T x)
870 return ck::type_convert<T>(__logf(ck::type_convert<float>(x)));
876 return hlog(
static_cast<__half
>(x));
891 template <
typename T>
892 inline __device__ T
pow(T x, T gamma)
894 return ck::type_convert<T>(powf(ck::type_convert<float>(x), ck::type_convert<float>(gamma)));
898 inline __device__
float pow<float>(
float x,
float gamma)
900 return powf(x, gamma);
904 inline __device__
double pow<double>(
double x,
double gamma)
906 return pow(x, gamma);
909 template <
typename T>
912 return ck::type_convert<T>(expm1f(ck::type_convert<float>(x)));
__device__ T asin(T x)
Definition: math_v2.hpp:660
__host__ T log(T x)
Definition: math_v2.hpp:409
__host__ int8_t neg< int8_t >(int8_t x)
Definition: math_v2.hpp:163
__host__ T cosh(T x)
Definition: math_v2.hpp:349
__host__ double asin< double >(double x)
Definition: math_v2.hpp:217
__host__ T exp(T x)
Definition: math_v2.hpp:391
__host__ double atan< double >(double x)
Definition: math_v2.hpp:181
__host__ double sinh< double >(double x)
Definition: math_v2.hpp:325
__host__ float floor< float >(float x)
Definition: math_v2.hpp:373
__host__ double ceil< double >(double x)
Definition: math_v2.hpp:343
__host__ float acosh< float >(float x)
Definition: math_v2.hpp:265
__host__ float log< float >(float x)
Definition: math_v2.hpp:415
__host__ T rcp(T x)
Definition: math_v2.hpp:385
__host__ double expm1< double >(double x)
Definition: math_v2.hpp:458
__host__ float asinh< float >(float x)
Definition: math_v2.hpp:229
__host__ T tan(T x)
Definition: math_v2.hpp:277
__device__ T sinh(T x)
Definition: math_v2.hpp:750
__host__ double log< double >(double x)
Definition: math_v2.hpp:421
__device__ half_t floor< half_t >(half_t x)
Definition: math_v2.hpp:828
__host__ T sin(T x)
Definition: math_v2.hpp:187
__host__ float asin< float >(float x)
Definition: math_v2.hpp:211
__host__ T expm1(T x)
Definition: math_v2.hpp:446
__host__ double exp< double >(double x)
Definition: math_v2.hpp:403
__device__ half_t log< half_t >(half_t x)
Definition: math_v2.hpp:874
__host__ double tan< double >(double x)
Definition: math_v2.hpp:289
__device__ T asinh(T x)
Definition: math_v2.hpp:678
__device__ half_t sin< half_t >(half_t x)
Definition: math_v2.hpp:654
__device__ T acos(T x)
Definition: math_v2.hpp:564
__host__ T atan(T x)
Definition: math_v2.hpp:169
__host__ float sin< float >(float x)
Definition: math_v2.hpp:193
__host__ float pow< float >(float x, float gamma)
Definition: math_v2.hpp:434
__host__ T acos(T x)
Definition: math_v2.hpp:121
__host__ double acos< double >(double x)
Definition: math_v2.hpp:133
__host__ double cos< double >(double x)
Definition: math_v2.hpp:253
__device__ T expm1(T x)
Definition: math_v2.hpp:910
__host__ double acosh< double >(double x)
Definition: math_v2.hpp:271
__host__ double atanh< double >(double x)
Definition: math_v2.hpp:307
__device__ T log(T x)
Definition: math_v2.hpp:868
__host__ T ceil(T x)
Definition: math_v2.hpp:331
__host__ T asin(T x)
Definition: math_v2.hpp:205
__host__ T pow(T x, T gamma)
Definition: math_v2.hpp:427
__host__ T neg(T x)
Definition: math_v2.hpp:139
__host__ double pow< double >(double x, double gamma)
Definition: math_v2.hpp:440
__device__ T sin(T x)
Definition: math_v2.hpp:636
__device__ half_t neg< half_t >(half_t x)
Definition: math_v2.hpp:612
__host__ float cos< float >(float x)
Definition: math_v2.hpp:247
__host__ int32_t neg< int32_t >(int32_t x)
Definition: math_v2.hpp:157
__host__ float tan< float >(float x)
Definition: math_v2.hpp:283
__host__ float expm1< float >(float x)
Definition: math_v2.hpp:452
__device__ T atanh(T x)
Definition: math_v2.hpp:732
__device__ T ceil(T x)
Definition: math_v2.hpp:768
__host__ double floor< double >(double x)
Definition: math_v2.hpp:379
__host__ float acos< float >(float x)
Definition: math_v2.hpp:127
__device__ T exp(T x)
Definition: math_v2.hpp:844
__host__ float exp< float >(float x)
Definition: math_v2.hpp:397
__host__ float atanh< float >(float x)
Definition: math_v2.hpp:301
__host__ float neg< float >(float x)
Definition: math_v2.hpp:145
__host__ double asinh< double >(double x)
Definition: math_v2.hpp:235
__host__ T cos(T x)
Definition: math_v2.hpp:241
__host__ T floor(T x)
Definition: math_v2.hpp:367
__host__ double cosh< double >(double x)
Definition: math_v2.hpp:361
__device__ T cosh(T x)
Definition: math_v2.hpp:792
__host__ float tanh< float >(float x)
Definition: math_v2.hpp:109
__device__ T tan(T x)
Definition: math_v2.hpp:714
__host__ float atan< float >(float x)
Definition: math_v2.hpp:175
__host__ float cosh< float >(float x)
Definition: math_v2.hpp:355
__host__ T asinh(T x)
Definition: math_v2.hpp:223
__host__ float sinh< float >(float x)
Definition: math_v2.hpp:319
__device__ T acosh(T x)
Definition: math_v2.hpp:696
__host__ float ceil< float >(float x)
Definition: math_v2.hpp:337
__device__ half_t exp< half_t >(half_t x)
Definition: math_v2.hpp:850
__host__ T atanh(T x)
Definition: math_v2.hpp:295
__host__ double neg< double >(double x)
Definition: math_v2.hpp:151
__host__ T tanh(T x)
Definition: math_v2.hpp:103
__host__ double tanh< double >(double x)
Definition: math_v2.hpp:115
__device__ T tanh(T x)
Definition: math_v2.hpp:546
__host__ double sin< double >(double x)
Definition: math_v2.hpp:199
__device__ T pow(T x, T gamma)
Definition: math_v2.hpp:892
__host__ T acosh(T x)
Definition: math_v2.hpp:259
__device__ half_t ceil< half_t >(half_t x)
Definition: math_v2.hpp:786
__device__ T floor(T x)
Definition: math_v2.hpp:810
__device__ T atan(T x)
Definition: math_v2.hpp:618
__host__ T sinh(T x)
Definition: math_v2.hpp:313
int8_t int8_t
Definition: int8.hpp:20
f8_fnuz_t f8_t
Definition: amd_ck_fp8.hpp:990
_Float16 half_t
Definition: data_type.hpp:25
_BitInt(4) int4_t
Definition: data_type.hpp:26