diff --git a/api/hip/hip_hcc.def.in b/api/hip/hip_hcc.def.in index 71e87d9a16..f93517c67a 100644 --- a/api/hip/hip_hcc.def.in +++ b/api/hip/hip_hcc.def.in @@ -131,6 +131,8 @@ __hipRegisterFatBinary __hipRegisterFunction __hipRegisterVar __hipUnregisterFatBinary +__gnu_f2h_ieee +__gnu_h2f_ieee hipConfigureCall hipSetupArgument hipLaunchByPtr diff --git a/api/hip/hip_hcc.map.in b/api/hip/hip_hcc.map.in index 48bb8eccba..c983f83430 100644 --- a/api/hip/hip_hcc.map.in +++ b/api/hip/hip_hcc.map.in @@ -131,6 +131,8 @@ global: __hipRegisterFunction; __hipRegisterVar; __hipUnregisterFatBinary; + __gnu_h2f_ieee; + __gnu_f2h_ieee; hipConfigureCall; hipSetupArgument; hipLaunchByPtr; diff --git a/api/hip/hip_platform.cpp b/api/hip/hip_platform.cpp index 05e03123a8..a6bb4e290c 100644 --- a/api/hip/hip_platform.cpp +++ b/api/hip/hip_platform.cpp @@ -509,4 +509,44 @@ void hipLaunchKernelGGLImpl( } +// conversion routines between float and half precision +static inline std::uint32_t f32_as_u32(float f) { union { float f; std::uint32_t u; } v; v.f = f; return v.u; } +static inline float u32_as_f32(std::uint32_t u) { union { float f; std::uint32_t u; } v; v.u = u; return v.f; } +static inline int clamp_int(int i, int l, int h) { return std::min(std::max(i, l), h); } + +// half � float, the f16 is in the low 16 bits of the input argument �a� +static inline float __convert_half_to_float(std::uint32_t a) noexcept { + std::uint32_t u = ((a << 13) + 0x70000000U) & 0x8fffe000U; + std::uint32_t v = f32_as_u32(u32_as_f32(u) * 0x1.0p+112f) + 0x38000000U; + u = (a & 0x7fff) != 0 ? v : u; + return u32_as_f32(u) * 0x1.0p-112f; +} + +// float � half with nearest even rounding +// The lower 16 bits of the result is the bit pattern for the f16 +static inline std::uint32_t __convert_float_to_half(float a) noexcept { + std::uint32_t u = f32_as_u32(a); + int e = static_cast((u >> 23) & 0xff) - 127 + 15; + std::uint32_t m = ((u >> 11) & 0xffe) | ((u & 0xfff) != 0); + std::uint32_t i = 0x7c00 | (m != 0 ? 0x0200 : 0); + std::uint32_t n = ((std::uint32_t)e << 12) | m; + std::uint32_t s = (u >> 16) & 0x8000; + int b = clamp_int(1-e, 0, 13); + std::uint32_t d = (0x1000 | m) >> b; + d |= (d << b) != (0x1000 | m); + std::uint32_t v = e < 1 ? d : n; + v = (v >> 2) + (((v & 0x7) == 3) | ((v & 0x7) > 5)); + v = e > 30 ? 0x7c00 : v; + v = e == 143 ? i : v; + return s | v; +} + +extern "C" float __gnu_h2f_ieee(unsigned short h){ + return __convert_half_to_float((std::uint32_t) h); +} + +extern "C" unsigned short __gnu_f2h_ieee(float f){ + return (unsigned short)__convert_float_to_half(f); +} + #endif // defined(ATI_OS_LINUX)