"...composable_kernel-1.git" did not exist on "08f16bd4dd3bde087f38e097f98376d3be27f3f4"
Unverified Commit ce87b4f7 authored by guangzlu's avatar guangzlu Committed by GitHub
Browse files

modified half function in math_v2.hpp (#528)


Co-authored-by: default avatarChao Liu <chao.liu2@amd.com>
parent d072790f
...@@ -114,7 +114,16 @@ static inline __device__ int4_t abs(int4_t x) ...@@ -114,7 +114,16 @@ static inline __device__ int4_t abs(int4_t x)
}; };
#endif #endif
static inline __device__ half_t abs(half_t x) { return ::__habs(x); }; static inline __device__ half_t abs(half_t x)
{
uint16_t xx = ck::bit_cast<uint16_t>(x);
uint16_t abs_xx = xx & 0x7fff;
half_t abs_x = ck::bit_cast<half_t>(abs_xx);
return abs_x;
};
static inline __device__ bool isnan(float x) { return ::isnan(x); }; static inline __device__ bool isnan(float x) { return ::isnan(x); };
...@@ -140,7 +149,12 @@ static inline __device__ bool isnan(int4_t x) ...@@ -140,7 +149,12 @@ static inline __device__ bool isnan(int4_t x)
}; };
#endif #endif
static inline __device__ bool isnan(half_t x) { return ::__hisnan(x); }; static inline __device__ bool isnan(half_t x)
{
uint16_t xx = ck::bit_cast<uint16_t>(x);
return (xx & 0x7FFF) > 0x7C00;
};
static inline __device__ float sqrt(float x) { return ::sqrtf(x); }; static inline __device__ float sqrt(float x) { return ::sqrtf(x); };
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment