diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 8ae6d0d85b..755cb19f6d 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -36,17 +36,6 @@ typedef struct __attribute__((aligned(4))){
   };
 } __half2;
 
-struct hipHalfHolder{
-  union {
-    __half h;
-    unsigned short s;
-  };
-};
-
-#define HINF 65504
-
-static struct hipHalfHolder __hInfValue = {HINF};
-
 extern "C" __half __hip_hc_ir_hadd_half(__half, __half);
 extern "C" __half __hip_hc_ir_hfma_half(__half, __half, __half);
 extern "C" __half __hip_hc_ir_hmul_half(__half, __half);
@@ -186,456 +175,6 @@ __device__ static inline __half2 h2div(__half2 a, __half2 b) {
   return c;
 }
 
-/*
-Half comparision Functions
-*/
-
-__device__ static inline bool __heq(__half a, __half b) {
-  return a == b ? true : false;
-}
-
-__device__ static inline bool __hge(__half a, __half b) {
-  return a >= b ? true : false;
-}
-
-__device__ static inline bool __hgt(__half a, __half b) {
-  return a > b ? true : false;
-}
-
-__device__ static inline bool __hisinf(__half a) {
-  return a == __hInfValue.h ? true : false;
-}
-
-__device__ static inline bool __hisnan(__half a) {
-  return a > __hInfValue.h ? true : false;
-}
-
-__device__ static inline bool __hle(__half a, __half b) {
-  return a <= b ? true : false;
-}
-
-__device__ static inline bool __hlt(__half a, __half b) {
-  return a < b ? true : false;
-}
-
-__device__ static inline bool __hne(__half a, __half b) {
-  return a != b ? true : false;
-}
-
-/*
-Half2 Comparision Functions
-*/
-
-__device__ static inline bool __hbeq2(__half2 a, __half2 b) {
-  return (a.p[0] == b.p[0] ? true : false) && (a.p[1] == b.p[1] ? true : false);
-}
-
-__device__ static inline bool __hbge2(__half2 a, __half2 b) {
-  return (a.p[0] >= b.p[0] ? true : false) && (a.p[1] >= b.p[1] ? true : false);
-}
-
-__device__ static inline bool __hbgt2(__half2 a, __half2 b) {
-  return (a.p[0] > b.p[0] ? true : false) && (a.p[1] > b.p[1] ? true : false);
-}
-
-__device__ static inline bool __hble2(__half2 a, __half2 b) {
-  return (a.p[0] <= b.p[0] ? true : false) && (a.p[1] <= b.p[1] ? true : false);
-}
-
-__device__ static inline bool __hblt2(__half2 a, __half2 b) {
-  return (a.p[0] < b.p[0] ? true : false) && (a.p[1] < b.p[1] ? true : false);
-}
-
-__device__ static inline bool __hbne2(__half2 a, __half2 b) {
-  return (a.p[0] != b.p[0] ? true : false) && (a.p[1] != b.p[1] ? true : false);
-}
-
-__device__ static inline __half2 __heq2(__half2 a, __half2 b) {
-  __half2 c;
-  c.p[0] = (a.p[0] == b.p[0]) ? (__half)1 : (__half)0;
-  c.p[1] = (a.p[1] == b.p[1]) ? (__half)1 : (__half)0;
-  return c;
-}
-
-__device__ static inline __half2 __hge2(__half2 a, __half2 b) {
-  __half2 c;
-  c.p[0] = (a.p[0] >= b.p[0]) ? (__half)1 : (__half)0;
-  c.p[1] = (a.p[1] >= b.p[1]) ? (__half)1 : (__half)0;
-  return c;
-}
-
-__device__ static inline __half2 __hgt2(__half2 a, __half2 b) {
-  __half2 c;
-  c.p[0] = (a.p[0] > b.p[0]) ? (__half)1 : (__half)0;
-  c.p[1] = (a.p[1] > b.p[1]) ? (__half)1 : (__half)0;
-  return c;
-}
-
-__device__ static inline __half2 __hisnan2(__half2 a) {
-  __half2 c;
-  c.p[0] = (a.p[0] > __hInfValue.h) ? (__half)1 : (__half)0;
-  c.p[1] = (a.p[1] > __hInfValue.h) ? (__half)1 : (__half)0;
-  return c;
-}
-
-__device__ static inline __half2 __hle2(__half2 a, __half2 b) {
-  __half2 c;
-  c.p[0] = (a.p[0] <= b.p[0]) ? (__half)1 : (__half)0;
-  c.p[1] = (a.p[1] <= b.p[1]) ? (__half)1 : (__half)0;
-  return c;
-}
-
-__device__ static inline __half2 __hlt2(__half2 a, __half2 b) {
-  __half2 c;
-  c.p[0] = (a.p[0] < b.p[0]) ? (__half)1 : (__half)0;
-  c.p[1] = (a.p[1] < b.p[1]) ? (__half)1 : (__half)0;
-  return c;
-}
-
-__device__ static inline __half2 __hne2(__half2 a, __half2 b) {
-  __half2 c;
-  c.p[0] = (a.p[0] != b.p[0]) ? (__half)1 : (__half)0;
-  c.p[1] = (a.p[1] != b.p[1]) ? (__half)1 : (__half)0;
-  return c;
-}
-
-/*
-Conversion instructions
-*/
-
-__device__ static inline __half2 __float22half2_rn(const float2 a) {
-  __half2 b;
-  b.p[0] = (__half)a.x;
-  b.p[1] = (__half)a.y;
-  return b;
-}
-
-__device__ static inline __half __float2half(const float a) {
-  return (__half)a;
-}
-
-__device__ static inline __half2 __float2half2_rn(const float a) {
-  __half2 b;
-  b.p[0] = (__half)a;
-  b.p[1] = (__half)a;
-  return b;
-}
-
-__device__ static inline __half __float2half_rd(const float a) {
-  return (__half)a;
-}
-
-__device__ static inline __half __float2half_rn(const float a) {
-  return (__half)a;
-}
-
-__device__ static inline __half __float2half_ru(const float a) {
-  return (__half)a;
-}
-
-__device__ static inline __half __float2half_rz(const float a) {
-  return (__half)a;
-}
-
-__device__ static inline __half2 __floats2half2_rn(const float a, const float b) {
-  __half2 c;
-  c.p[0] = (__half)a;
-  c.p[1] = (__half)b;
-  return c;
-}
-
-__device__ static inline float2 __half22float2(const __half2 a) {
-  float2 b;
-  b.x = (float)a.p[0];
-  b.y = (float)a.p[1];
-  return b;
-}
-
-__device__ static inline float __half2float(const __half a) {
-  return (float)a;
-}
-
-__device__ static inline __half2 half2half2(const __half a) {
-  __half2 b;
-  b.p[0] = a;
-  b.p[1] = a;
-  return b;
-}
-
-__device__ static inline int __half2int_rd(__half h) {
-  return (int)h;
-}
-
-__device__ static inline int __half2int_rn(__half h) {
-  return (int)h;
-}
-
-__device__ static inline int __half2int_ru(__half h) {
-  return (int)h;
-}
-
-__device__ static inline int __half2int_rz(__half h) {
-  return (int)h;
-}
-
-__device__ static inline long long int __half2ll_rd(__half h) {
-  return (long long int)h;
-}
-
-__device__ static inline long long int __half2ll_rn(__half h) {
-  return (long long int)h;
-}
-
-__device__ static inline long long int __half2ll_ru(__half h) {
-  return (long long int)h;
-}
-
-__device__ static inline long long int __half2ll_rz(__half h) {
-  return (long long int)h;
-}
-
-__device__ static inline short __half2short_rd(__half h) {
-  return (short)h;
-}
-
-__device__ static inline short __half2short_rn(__half h) {
-  return (short)h;
-}
-
-__device__ static inline short __half2short_ru(__half h) {
-  return (short)h;
-}
-
-__device__ static inline short __half2short_rz(__half h) {
-  return (short)h;
-}
-
-__device__ static inline unsigned int __half2uint_rd(__half h) {
-  return (unsigned int)h;
-}
-
-__device__ static inline unsigned int __half2uint_rn(__half h) {
-  return (unsigned int)h;
-}
-
-__device__ static inline unsigned int __half2uint_ru(__half h) {
-  return (unsigned int)h;
-}
-
-__device__ static inline unsigned int __half2uint_rz(__half h) {
-  return (unsigned int)h;
-}
-
-__device__ static inline unsigned long long int __half2ull_rd(__half h) {
-  return (unsigned long long)h;
-}
-
-__device__ static inline unsigned long long int __half2ull_rn(__half h) {
-  return (unsigned long long)h;
-}
-
-__device__ static inline unsigned long long int __half2ull_ru(__half h) {
-  return (unsigned long long)h;
-}
-
-__device__ static inline unsigned long long int __half2ull_rz(__half h) {
-  return (unsigned long long)h;
-}
-
-__device__ static inline unsigned short int __half2ushort_rd(__half h) {
-  return (unsigned short int)h;
-}
-
-__device__ static inline unsigned short int __half2ushort_rn(__half h) {
-  return (unsigned short int)h;
-}
-
-__device__ static inline unsigned short int __half2ushort_ru(__half h) {
-  return (unsigned short int)h;
-}
-
-__device__ static inline unsigned short int __half2ushort_rz(__half h) {
-  return (unsigned short int)h;
-}
-
-__device__ static inline short int __half_as_short(const __half h) {
-  hipHalfHolder hH;
-  hH.h = h;
-  return (short)hH.s;
-}
-
-__device__ static inline unsigned short int __half_as_ushort(const __half h) {
-  hipHalfHolder hH;
-  hH.h = h;
-  return hH.s;
-}
-
-__device__ static inline __half2 __halves2half2(const __half a, const __half b) {
-  __half2 c;
-  c.p[0] = a;
-  c.p[1] = b;
-  return c;
-}
-
-__device__ static inline float __high2float(const __half2 a) {
-  return (float)a.p[1];
-}
-
-__device__ static inline __half __high2half(const __half2 a) {
-  return a.p[1];
-}
-
-__device__ static inline __half2 __high2half2(const __half2 a) {
-  __half2 b;
-  b.p[0] = a.p[1];
-  b.p[1] = a.p[1];
-  return b;
-}
-
-__device__ static inline __half2 __highs2half2(const __half2 a, const __half2 b) {
-  __half2 c;
-  c.p[0] = a.p[1];
-  c.p[1] = b.p[1];
-  return c;
-}
-
-__device__ static inline __half __int2half_rd(int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __int2half_rn(int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __int2half_ru(int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __int2half_rz(int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ll2half_rd(long long int i){
-  return (__half)i;
-}
-
-__device__ static inline __half __ll2half_rn(long long int i){
-  return (__half)i;
-}
-
-__device__ static inline __half __ll2half_ru(long long int i){
-  return (__half)i;
-}
-
-__device__ static inline __half __ll2half_rz(long long int i){
-  return (__half)i;
-}
-
-__device__ static inline float __low2float(const __half2 a) {
-  return (float)a.p[0];
-}
-
-__device__ static inline __half __low2half(const __half2 a) {
-  return a.p[0];
-}
-
-__device__ static inline __half2 __low2half2(const __half2 a, const __half2 b) {
-  __half2 c;
-  c.p[0] = a.p[0];
-  c.p[1] = b.p[0];
-  return c;
-}
-
-__device__ static inline __half2 __low2half2(const __half2 a) {
-  __half2 b;
-  b.p[0] = a.p[0];
-  b.p[1] = a.p[0];
-  return b;
-}
-
-__device__ static inline __half2 __lowhigh2highlow(const __half2 a) {
-  __half2 b;
-  b.p[0] = a.p[1];
-  b.p[1] = a.p[0];
-  return b;
-}
-
-__device__ static inline __half2 __lows2half2(const __half2 a, const __half2 b) {
-  __half2 c;
-  c.p[0] = a.p[0];
-  c.p[1] = b.p[0];
-  return c;
-}
-
-__device__ static inline __half __short2half_rd(short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __short2half_rn(short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __short2half_ru(short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __short2half_rz(short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __uint2half_rd(unsigned int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __uint2half_rn(unsigned int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __uint2half_ru(unsigned int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __uint2half_rz(unsigned int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ull2half_rd(unsigned long long int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ull2half_rn(unsigned long long int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ull2half_ru(unsigned long long int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ull2half_rz(unsigned long long int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ushort2half_rd(unsigned short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ushort2half_rn(unsigned short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ushort2half_ru(unsigned short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ushort2half_rz(unsigned short int i) {
-  return (__half)i;
-}
-
-__device__ static inline __half __ushort_as_half(const unsigned short int i) {
-  hipHalfHolder hH;
-  hH.s = i;
-  return hH.h;
-}
 
 __device__ static inline __half hceil(const __half h) {
   return __hip_hc_ir_hceil_half(h);
@@ -792,6 +331,110 @@ __device__ static inline __half2 h2trunc(const __half2 h) {
   return a;
 }
 
+__device__  bool __heq(__half a, __half b);
+__device__  bool __hge(__half a, __half b);
+__device__  bool __hgt(__half a, __half b);
+__device__  bool __hisinf(__half a);
+__device__  bool __hisnan(__half a);
+__device__  bool __hle(__half a, __half b);
+__device__  bool __hlt(__half a, __half b);
+__device__  bool __hne(__half a, __half b);
+/*
+Half2 Comparision Functions
+*/
+
+__device__  bool __hbeq2(__half2 a, __half2 b);
+__device__  bool __hbge2(__half2 a, __half2 b);
+__device__  bool __hbgt2(__half2 a, __half2 b);
+__device__  bool __hble2(__half2 a, __half2 b);
+__device__  bool __hblt2(__half2 a, __half2 b);
+__device__  bool __hbne2(__half2 a, __half2 b);
+__device__  __half2 __heq2(__half2 a, __half2 b);
+__device__  __half2 __hge2(__half2 a, __half2 b);
+__device__  __half2 __hgt2(__half2 a, __half2 b);
+__device__  __half2 __hisnan2(__half2 a);
+__device__  __half2 __hle2(__half2 a, __half2 b);
+__device__  __half2 __hlt2(__half2 a, __half2 b);
+__device__  __half2 __hne2(__half2 a, __half2 b);
+
+/*
+Conversion instructions
+*/
+__device__  __half2 __float22half2_rn(const float2 a);
+__device__  __half __float2half(const float a);
+__device__  __half2 __float2half2_rn(const float a);
+__device__  __half __float2half_rd(const float a);
+__device__  __half __float2half_rn(const float a);
+__device__  __half __float2half_ru(const float a);
+__device__  __half __float2half_rz(const float a);
+__device__  __half2 __floats2half2_rn(const float a, const float b);
+__device__  float2 __half22float2(const __half2 a);
+__device__  float __half2float(const __half a);
+__device__  __half2 half2half2(const __half a);
+__device__  int __half2int_rd(__half h);
+__device__  int __half2int_rn(__half h);
+__device__  int __half2int_ru(__half h);
+__device__  int __half2int_rz(__half h);
+__device__  long long int __half2ll_rd(__half h);
+__device__  long long int __half2ll_rn(__half h);
+__device__  long long int __half2ll_ru(__half h);
+__device__  long long int __half2ll_rz(__half h);
+__device__  short __half2short_rd(__half h);
+__device__  short __half2short_rn(__half h);
+__device__  short __half2short_ru(__half h);
+__device__  short __half2short_rz(__half h);
+__device__  unsigned int __half2uint_rd(__half h);
+__device__  unsigned int __half2uint_rn(__half h);
+__device__  unsigned int __half2uint_ru(__half h);
+__device__  unsigned int __half2uint_rz(__half h);
+__device__  unsigned long long int __half2ull_rd(__half h);
+__device__  unsigned long long int __half2ull_rn(__half h);
+__device__  unsigned long long int __half2ull_ru(__half h);
+__device__  unsigned long long int __half2ull_rz(__half h);
+__device__  unsigned short int __half2ushort_rd(__half h);
+__device__  unsigned short int __half2ushort_rn(__half h);
+__device__  unsigned short int __half2ushort_ru(__half h);
+__device__  unsigned short int __half2ushort_rz(__half h);
+__device__  short int __half_as_short(const __half h);
+__device__  unsigned short int __half_as_ushort(const __half h);
+__device__  __half2 __halves2half2(const __half a, const __half b);
+__device__  float __high2float(const __half2 a);
+__device__  __half __high2half(const __half2 a);
+__device__  __half2 __high2half2(const __half2 a);
+__device__  __half2 __highs2half2(const __half2 a, const __half2 b);
+__device__  __half __int2half_rd(int i);
+__device__  __half __int2half_rn(int i);
+__device__  __half __int2half_ru(int i);
+__device__  __half __int2half_rz(int i);
+__device__  __half __ll2half_rd(long long int i);
+__device__  __half __ll2half_rn(long long int i);
+__device__  __half __ll2half_ru(long long int i);
+__device__  __half __ll2half_rz(long long int i);
+__device__  float __low2float(const __half2 a);
+
+__device__ __half __low2half(const __half2 a);
+__device__ __half2 __low2half2(const __half2 a, const __half2 b);
+__device__ __half2 __low2half2(const __half2 a);
+__device__ __half2 __lowhigh2highlow(const __half2 a);
+__device__ __half2 __lows2half2(const __half2 a, const __half2 b);
+__device__  __half __short2half_rd(short int i);
+__device__  __half __short2half_rn(short int i);
+__device__  __half __short2half_ru(short int i);
+__device__  __half __short2half_rz(short int i);
+__device__  __half __uint2half_rd(unsigned int i);
+__device__  __half __uint2half_rn(unsigned int i);
+__device__  __half __uint2half_ru(unsigned int i);
+__device__  __half __uint2half_rz(unsigned int i);
+__device__  __half __ull2half_rd(unsigned long long int i);
+__device__  __half __ull2half_rn(unsigned long long int i);
+__device__  __half __ull2half_ru(unsigned long long int i);
+__device__  __half __ull2half_rz(unsigned long long int i);
+__device__  __half __ushort2half_rd(unsigned short int i);
+__device__  __half __ushort2half_rn(unsigned short int i);
+__device__  __half __ushort2half_ru(unsigned short int i);
+__device__  __half __ushort2half_rz(unsigned short int i);
+__device__  __half __ushort_as_half(const unsigned short int i);
+
 
 #endif
 
diff --git a/hipamd/src/hip_fp16.cpp b/hipamd/src/hip_fp16.cpp
index 83e0a161c7..ac79ddba08 100644
--- a/hipamd/src/hip_fp16.cpp
+++ b/hipamd/src/hip_fp16.cpp
@@ -22,6 +22,472 @@ THE SOFTWARE.
 
 #include"hip/hip_fp16.h"
 
+struct hipHalfHolder{
+  union {
+    __half h;
+    unsigned short s;
+  };
+};
+
+#define HINF 65504
+
+static struct hipHalfHolder __hInfValue = {HINF};
+/*
+Half comparision Functions
+*/
+
+__device__  bool __heq(__half a, __half b) {
+  return a == b ? true : false;
+}
+
+__device__  bool __hge(__half a, __half b) {
+  return a >= b ? true : false;
+}
+
+__device__  bool __hgt(__half a, __half b) {
+  return a > b ? true : false;
+}
+
+__device__  bool __hisinf(__half a) {
+  return a == __hInfValue.h ? true : false;
+}
+
+__device__  bool __hisnan(__half a) {
+  return a > __hInfValue.h ? true : false;
+}
+
+__device__  bool __hle(__half a, __half b) {
+  return a <= b ? true : false;
+}
+
+__device__  bool __hlt(__half a, __half b) {
+  return a < b ? true : false;
+}
+
+__device__  bool __hne(__half a, __half b) {
+  return a != b ? true : false;
+}
+
+/*
+Half2 Comparision Functions
+*/
+
+__device__  bool __hbeq2(__half2 a, __half2 b) {
+  return (a.p[0] == b.p[0] ? true : false) && (a.p[1] == b.p[1] ? true : false);
+}
+
+__device__  bool __hbge2(__half2 a, __half2 b) {
+  return (a.p[0] >= b.p[0] ? true : false) && (a.p[1] >= b.p[1] ? true : false);
+}
+
+__device__  bool __hbgt2(__half2 a, __half2 b) {
+  return (a.p[0] > b.p[0] ? true : false) && (a.p[1] > b.p[1] ? true : false);
+}
+
+__device__  bool __hble2(__half2 a, __half2 b) {
+  return (a.p[0] <= b.p[0] ? true : false) && (a.p[1] <= b.p[1] ? true : false);
+}
+
+__device__  bool __hblt2(__half2 a, __half2 b) {
+  return (a.p[0] < b.p[0] ? true : false) && (a.p[1] < b.p[1] ? true : false);
+}
+
+__device__  bool __hbne2(__half2 a, __half2 b) {
+  return (a.p[0] != b.p[0] ? true : false) && (a.p[1] != b.p[1] ? true : false);
+}
+
+__device__  __half2 __heq2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] == b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] == b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__  __half2 __hge2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] >= b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] >= b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__  __half2 __hgt2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] > b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] > b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__  __half2 __hisnan2(__half2 a) {
+  __half2 c;
+  c.p[0] = (a.p[0] > __hInfValue.h) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] > __hInfValue.h) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__  __half2 __hle2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] <= b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] <= b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__  __half2 __hlt2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] < b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] < b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__  __half2 __hne2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] != b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] != b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+/*
+Conversion instructions
+*/
+__device__  __half2 __float22half2_rn(const float2 a) {
+  __half2 b;
+  b.p[0] = (__half)a.x;
+  b.p[1] = (__half)a.y;
+  return b;
+}
+
+__device__  __half __float2half(const float a) {
+  return (__half)a;
+}
+
+__device__  __half2 __float2half2_rn(const float a) {
+  __half2 b;
+  b.p[0] = (__half)a;
+  b.p[1] = (__half)a;
+  return b;
+}
+
+__device__  __half __float2half_rd(const float a) {
+  return (__half)a;
+}
+
+__device__  __half __float2half_rn(const float a) {
+  return (__half)a;
+}
+
+__device__  __half __float2half_ru(const float a) {
+  return (__half)a;
+}
+
+__device__  __half __float2half_rz(const float a) {
+  return (__half)a;
+}
+
+__device__  __half2 __floats2half2_rn(const float a, const float b) {
+  __half2 c;
+  c.p[0] = (__half)a;
+  c.p[1] = (__half)b;
+  return c;
+}
+
+__device__  float2 __half22float2(const __half2 a) {
+  float2 b;
+  b.x = (float)a.p[0];
+  b.y = (float)a.p[1];
+  return b;
+}
+
+__device__  float __half2float(const __half a) {
+  return (float)a;
+}
+
+__device__  __half2 half2half2(const __half a) {
+  __half2 b;
+  b.p[0] = a;
+  b.p[1] = a;
+  return b;
+}
+
+__device__  int __half2int_rd(__half h) {
+  return (int)h;
+}
+
+__device__  int __half2int_rn(__half h) {
+  return (int)h;
+}
+
+__device__  int __half2int_ru(__half h) {
+  return (int)h;
+}
+
+__device__  int __half2int_rz(__half h) {
+  return (int)h;
+}
+
+__device__  long long int __half2ll_rd(__half h) {
+  return (long long int)h;
+}
+
+__device__  long long int __half2ll_rn(__half h) {
+  return (long long int)h;
+}
+
+__device__  long long int __half2ll_ru(__half h) {
+  return (long long int)h;
+}
+
+__device__  long long int __half2ll_rz(__half h) {
+  return (long long int)h;
+}
+
+__device__  short __half2short_rd(__half h) {
+  return (short)h;
+}
+
+__device__  short __half2short_rn(__half h) {
+  return (short)h;
+}
+
+__device__  short __half2short_ru(__half h) {
+  return (short)h;
+}
+
+__device__  short __half2short_rz(__half h) {
+  return (short)h;
+}
+
+__device__  unsigned int __half2uint_rd(__half h) {
+  return (unsigned int)h;
+}
+
+__device__  unsigned int __half2uint_rn(__half h) {
+  return (unsigned int)h;
+}
+
+__device__  unsigned int __half2uint_ru(__half h) {
+  return (unsigned int)h;
+}
+
+__device__  unsigned int __half2uint_rz(__half h) {
+  return (unsigned int)h;
+}
+
+__device__  unsigned long long int __half2ull_rd(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__  unsigned long long int __half2ull_rn(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__  unsigned long long int __half2ull_ru(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__  unsigned long long int __half2ull_rz(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__  unsigned short int __half2ushort_rd(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__  unsigned short int __half2ushort_rn(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__  unsigned short int __half2ushort_ru(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__  unsigned short int __half2ushort_rz(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__  short int __half_as_short(const __half h) {
+  hipHalfHolder hH;
+  hH.h = h;
+  return (short)hH.s;
+}
+
+__device__  unsigned short int __half_as_ushort(const __half h) {
+  hipHalfHolder hH;
+  hH.h = h;
+  return hH.s;
+}
+
+__device__  __half2 __halves2half2(const __half a, const __half b) {
+  __half2 c;
+  c.p[0] = a;
+  c.p[1] = b;
+  return c;
+}
+
+__device__  float __high2float(const __half2 a) {
+  return (float)a.p[1];
+}
+
+__device__  __half __high2half(const __half2 a) {
+  return a.p[1];
+}
+
+__device__  __half2 __high2half2(const __half2 a) {
+  __half2 b;
+  b.p[0] = a.p[1];
+  b.p[1] = a.p[1];
+  return b;
+}
+
+__device__  __half2 __highs2half2(const __half2 a, const __half2 b) {
+  __half2 c;
+  c.p[0] = a.p[1];
+  c.p[1] = b.p[1];
+  return c;
+}
+
+__device__  __half __int2half_rd(int i) {
+  return (__half)i;
+}
+
+__device__  __half __int2half_rn(int i) {
+  return (__half)i;
+}
+
+__device__  __half __int2half_ru(int i) {
+  return (__half)i;
+}
+
+__device__  __half __int2half_rz(int i) {
+  return (__half)i;
+}
+
+__device__  __half __ll2half_rd(long long int i){
+  return (__half)i;
+}
+
+__device__  __half __ll2half_rn(long long int i){
+  return (__half)i;
+}
+
+__device__  __half __ll2half_ru(long long int i){
+  return (__half)i;
+}
+
+__device__  __half __ll2half_rz(long long int i){
+  return (__half)i;
+}
+
+__device__  float __low2float(const __half2 a) {
+  return (float)a.p[0];
+}
+
+__device__  __half __low2half(const __half2 a) {
+  return a.p[0];
+}
+
+__device__  __half2 __low2half2(const __half2 a, const __half2 b) {
+  __half2 c;
+  c.p[0] = a.p[0];
+  c.p[1] = b.p[0];
+  return c;
+}
+
+__device__  __half2 __low2half2(const __half2 a) {
+  __half2 b;
+  b.p[0] = a.p[0];
+  b.p[1] = a.p[0];
+  return b;
+}
+
+__device__  __half2 __lowhigh2highlow(const __half2 a) {
+  __half2 b;
+  b.p[0] = a.p[1];
+  b.p[1] = a.p[0];
+  return b;
+}
+
+__device__  __half2 __lows2half2(const __half2 a, const __half2 b) {
+  __half2 c;
+  c.p[0] = a.p[0];
+  c.p[1] = b.p[0];
+  return c;
+}
+
+__device__  __half __short2half_rd(short int i) {
+  return (__half)i;
+}
+
+__device__  __half __short2half_rn(short int i) {
+  return (__half)i;
+}
+
+__device__  __half __short2half_ru(short int i) {
+  return (__half)i;
+}
+
+__device__  __half __short2half_rz(short int i) {
+  return (__half)i;
+}
+
+__device__  __half __uint2half_rd(unsigned int i) {
+  return (__half)i;
+}
+
+__device__  __half __uint2half_rn(unsigned int i) {
+  return (__half)i;
+}
+
+__device__  __half __uint2half_ru(unsigned int i) {
+  return (__half)i;
+}
+
+__device__  __half __uint2half_rz(unsigned int i) {
+  return (__half)i;
+}
+
+__device__  __half __ull2half_rd(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__  __half __ull2half_rn(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__  __half __ull2half_ru(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__  __half __ull2half_rz(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__  __half __ushort2half_rd(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__  __half __ushort2half_rn(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__  __half __ushort2half_ru(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__  __half __ushort2half_rz(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__  __half __ushort_as_half(const unsigned short int i) {
+  hipHalfHolder hH;
+  hH.s = i;
+  return hH.h;
+}
+
+
+/*
+Soft Implementation. Use it for backup.
+*/
+
+
 static const unsigned sign_val = 0x8000;
 static const __half __half_value_one_float = {0x3C00};
 static const __half __half_value_zero_float = {0x0};
@@ -375,4 +841,6 @@ __device__ __half2 __soft_low2half2(const __half2 a, const __half2 b){
   return {a.p[0], b.p[0]};
 }
 
+
+
 #endif