From db0a1032b1caa4a53e7979d6452de272e566742f Mon Sep 17 00:00:00 2001 From: guangzlu <87220526+guangzlu@users.noreply.github.com> Date: Thu, 8 Dec 2022 07:43:02 +0800 Subject: [PATCH] modified half function in math_v2.hpp (#528) Co-authored-by: Chao Liu [ROCm/composable_kernel commit: ce87b4f76529b940b4ab93a76cc88828cb91d15a] --- include/ck/utility/math_v2.hpp | 18 ++++++++++++++++-- 1 file changed, 16 insertions(+), 2 deletions(-) diff --git a/include/ck/utility/math_v2.hpp b/include/ck/utility/math_v2.hpp index 84a057815f..dc97666b37 100644 --- a/include/ck/utility/math_v2.hpp +++ b/include/ck/utility/math_v2.hpp @@ -114,7 +114,16 @@ static inline __device__ int4_t abs(int4_t x) }; #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(x); + + uint16_t abs_xx = xx & 0x7fff; + + half_t abs_x = ck::bit_cast(abs_xx); + + return abs_x; +}; static inline __device__ bool isnan(float x) { return ::isnan(x); }; @@ -140,7 +149,12 @@ static inline __device__ bool isnan(int4_t x) }; #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(x); + + return (xx & 0x7FFF) > 0x7C00; +}; static inline __device__ float sqrt(float x) { return ::sqrtf(x); };