From cbc01d5488d160bb0e0403e705946f136284096d Mon Sep 17 00:00:00 2001
From: Tao Sang <tao.sang@amd.com>
Date: Tue, 21 Jul 2020 18:37:05 -0400
Subject: [PATCH] Remove __gnu_h2f_ieee and __gnu_f2h_ieee

Change-Id: Ide24d245d851e20961020323e52f33322a33fff9


[ROCm/hip commit: ef9176db03697eb88728f459e93539e7d94e7fe7]
---
 projects/hip/rocclr/hip_hcc.def.in   |  2 --
 projects/hip/rocclr/hip_hcc.map.in   |  2 --
 projects/hip/rocclr/hip_platform.cpp | 40 ----------------------------
 3 files changed, 44 deletions(-)

diff --git a/projects/hip/rocclr/hip_hcc.def.in b/projects/hip/rocclr/hip_hcc.def.in
index 5e45689c1a..253352f2cb 100755
--- a/projects/hip/rocclr/hip_hcc.def.in
+++ b/projects/hip/rocclr/hip_hcc.def.in
@@ -176,8 +176,6 @@ __hipRegisterVar
 __hipRegisterSurface
 __hipRegisterTexture
 __hipUnregisterFatBinary
-__gnu_h2f_ieee
-__gnu_f2h_ieee
 hipConfigureCall
 hipSetupArgument
 hipLaunchByPtr
diff --git a/projects/hip/rocclr/hip_hcc.map.in b/projects/hip/rocclr/hip_hcc.map.in
index f355a4e14a..be83e6d134 100755
--- a/projects/hip/rocclr/hip_hcc.map.in
+++ b/projects/hip/rocclr/hip_hcc.map.in
@@ -176,8 +176,6 @@ global:
     __hipRegisterSurface;
     __hipRegisterTexture;
     __hipUnregisterFatBinary;
-    __gnu_h2f_ieee;
-    __gnu_f2h_ieee;
     hipConfigureCall;
     hipSetupArgument;
     hipLaunchByPtr;
diff --git a/projects/hip/rocclr/hip_platform.cpp b/projects/hip/rocclr/hip_platform.cpp
index cb3cdf7f97..0dafd268a4 100755
--- a/projects/hip/rocclr/hip_platform.cpp
+++ b/projects/hip/rocclr/hip_platform.cpp
@@ -635,46 +635,6 @@ hipError_t ihipLaunchKernel(const void* hostFunction,
                                     flags));
 }
 
-// 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
-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) * u32_as_f32(0x77800000U)/*0x1.0p+112f*/) + 0x38000000U;
-  u = (a & 0x7fff) != 0 ? v : u;
-  return u32_as_f32(u) * u32_as_f32(0x07800000U)/*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);
-}
-
 void PlatformState::init()
 {
   amd::ScopedLock lock(lock_);