P4 to Git Change 1589138 by cpaquot@cpaquot-ocl-lc-lnx on 2018/08/03 01:04:05

SWDEV-145570 - [HIP] Added __gnu_h2f_ieee and __gnu_f2h_ieee

Affected files ...

... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.def.in#9 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_hcc.map.in#10 edit
... //depot/stg/opencl/drivers/opencl/api/hip/hip_platform.cpp#16 edit

api/hip/hip_hcc.def.in

@@ -131,6 +131,8 @@ __hipRegisterFatBinary
 __hipRegisterFunction
 __hipRegisterVar
 __hipUnregisterFatBinary
+__gnu_f2h_ieee
+__gnu_h2f_ieee
 hipConfigureCall
 hipSetupArgument
 hipLaunchByPtr

api/hip/hip_hcc.map.in

@@ -131,6 +131,8 @@ global:
     __hipRegisterFunction;
     __hipRegisterVar;
     __hipUnregisterFatBinary;
+    __gnu_h2f_ieee;
+    __gnu_f2h_ieee;
     hipConfigureCall;
     hipSetupArgument;
     hipLaunchByPtr;

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<int>((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)