diff --git a/projects/hip/docs/markdown/hip-math-api.md b/projects/hip/docs/markdown/hip-math-api.md
index daf34fbfa3..6a8e90eb59 100644
--- a/projects/hip/docs/markdown/hip-math-api.md
+++ b/projects/hip/docs/markdown/hip-math-api.md
@@ -16,7 +16,7 @@ __device__ static inline float __dotf(float4 x, float4 y) {
 
 This helps python script to add the device function newly declared into markdown documentation (as it looks at functions with `;` at the end and `__device__` at the beginning)
 
-The next step would be to add Description to  `deviceFuncDesc`.
+The next step would be to add Description to  `deviceFuncDesc` dictionary in python script.
 From the above example, it can be writtern as,
 `deviceFuncDesc['__dotf'] = 'This functions takes 2 4 component float vector and outputs dot product across them'`
 
@@ -2716,6 +2716,166 @@ __device__ float __ull2float_rz(unsigned long long int x);
 **Description:**  Supported
 
 
+### __hadd
+```cpp 
+__device__ static __half __hadd(const __half a, const __half b);
+
+```
+**Description:**  Supported
+
+
+### __hadd_sat
+```cpp 
+__device__ static __half __hadd_sat(__half a, __half b);
+
+```
+**Description:**  Supported
+
+
+### __hfma
+```cpp 
+__device__ static __half __hfma(__half a, __half b, __half c);
+
+```
+**Description:**  Supported
+
+
+### __hfma_sat
+```cpp 
+__device__ static __half __hfma_sat(__half a, __half b, __half c);
+
+```
+**Description:**  Supported
+
+
+### __hmul
+```cpp 
+__device__ static __half __hmul(__half a, __half b);
+
+```
+**Description:**  Supported
+
+
+### __hmul_sat
+```cpp 
+__device__ static __half __hmul_sat(__half a, __half b);
+
+```
+**Description:**  Supported
+
+
+### __hneg
+```cpp 
+__device__ static __half __hneg(__half a);
+
+```
+**Description:**  Supported
+
+
+### __hsub
+```cpp 
+__device__ static __half __hsub(__half a, __half b);
+
+```
+**Description:**  Supported
+
+
+### __hsub_sat
+```cpp 
+__device__ static __half __hsub_sat(__half a, __half b);
+
+```
+**Description:**  Supported
+
+
+### hdiv
+```cpp 
+__device__ static __half hdiv(__half a, __half b);
+
+```
+**Description:**  Supported
+
+
+### __hadd2
+```cpp 
+__device__ static __half2 __hadd2(__half2 a, __half2 b);
+
+```
+**Description:**  Supported
+
+
+### __hadd2_sat
+```cpp 
+__device__ static __half2 __hadd2_sat(__half2 a, __half2 b);
+
+```
+**Description:**  Supported
+
+
+### __hfma2
+```cpp 
+__device__ static __half2 __hfma2(__half2 a, __half2 b, __half2 c);
+
+```
+**Description:**  Supported
+
+
+### __hfma2_sat
+```cpp 
+__device__ static __half2 __hfma2_sat(__half2 a, __half2 b, __half2 c);
+
+```
+**Description:**  Supported
+
+
+### __hmul2
+```cpp 
+__device__ static __half2 __hmul2(__half2 a, __half2 b);
+
+```
+**Description:**  Supported
+
+
+### __hmul2_sat
+```cpp 
+__device__ static __half2 __hmul2_sat(__half2 a, __half2 b);
+
+```
+**Description:**  Supported
+
+
+### __hsub2
+```cpp 
+__device__ static __half2 __hsub2(__half2 a, __half2 b);
+
+```
+**Description:**  Supported
+
+
+### __hneg2
+```cpp 
+__device__ static __half2 __hneg2(__half2 a);
+
+```
+**Description:**  Supported
+
+
+### __hsub2_sat
+```cpp 
+__device__ static __half2 __hsub2_sat(__half2 a, __half2 b);
+
+```
+**Description:**  Supported
+
+
+### h2div
+```cpp 
+__device__ static __half2 h2div(__half2 a, __half2 b);
+
+```
+**Description:**  Supported
+
+
 ### __heq
 ```cpp 
 __device__  bool __heq(__half a, __half b);
@@ -2884,6 +3044,214 @@ __device__  __half2 __hne2(__half2 a, __half2 b);
 **Description:**  Supported
 
 
+### hceil
+```cpp 
+__device__ static __half hceil(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hcos
+```cpp 
+__device__ static __half hcos(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hexp
+```cpp 
+__device__ static __half hexp(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hexp10
+```cpp 
+__device__ static __half hexp10(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hexp2
+```cpp 
+__device__ static __half hexp2(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hfloor
+```cpp 
+__device__ static __half hfloor(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hlog
+```cpp 
+__device__ static __half hlog(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hlog10
+```cpp 
+__device__ static __half hlog10(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hlog2
+```cpp 
+__device__ static __half hlog2(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hrcp
+```cpp 
+//__device__ static __half hrcp(const __half h);
+
+```
+**Description:**  **NOT Supported**
+
+
+### hrint
+```cpp 
+__device__ static __half hrint(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hsin
+```cpp 
+__device__ static __half hsin(const __half h);
+
+```
+**Description:**  Supported
+
+
+### hsqrt
+```cpp 
+__device__ static __half hsqrt(const __half a);
+
+```
+**Description:**  Supported
+
+
+### htrunc
+```cpp 
+__device__ static __half htrunc(const __half a);
+
+```
+**Description:**  Supported
+
+
+### h2ceil
+```cpp 
+__device__ static __half2 h2ceil(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2exp
+```cpp 
+__device__ static __half2 h2exp(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2exp10
+```cpp 
+__device__ static __half2 h2exp10(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2exp2
+```cpp 
+__device__ static __half2 h2exp2(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2floor
+```cpp 
+__device__ static __half2 h2floor(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2log
+```cpp 
+__device__ static __half2 h2log(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2log10
+```cpp 
+__device__ static __half2 h2log10(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2log2
+```cpp 
+__device__ static __half2 h2log2(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2rcp
+```cpp 
+__device__ static __half2 h2rcp(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2rsqrt
+```cpp 
+__device__ static __half2 h2rsqrt(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2sin
+```cpp 
+__device__ static __half2 h2sin(const __half2 h);
+
+```
+**Description:**  Supported
+
+
+### h2sqrt
+```cpp 
+__device__ static __half2 h2sqrt(const __half2 h);
+
+```
+**Description:**  Supported
+
+
 ### __float22half2_rn
 ```cpp 
 __device__  __half2 __float22half2_rn(const float2 a);
diff --git a/projects/hip/include/hip/hcc_detail/hip_fp16.h b/projects/hip/include/hip/hcc_detail/hip_fp16.h
index 73049eb5fb..67d1fe4e06 100644
--- a/projects/hip/include/hip/hcc_detail/hip_fp16.h
+++ b/projects/hip/include/hip/hcc_detail/hip_fp16.h
@@ -36,6 +36,181 @@ typedef struct __attribute__((aligned(4))){
   };
 } __half2;
 
+/*
+Half Arithmetic Functions
+*/
+__device__ static __half __hadd(const __half a, const __half b);
+__device__ static __half __hadd_sat(__half a, __half b);
+__device__ static __half __hfma(__half a, __half b, __half c);
+__device__ static __half __hfma_sat(__half a, __half b, __half c);
+__device__ static __half __hmul(__half a, __half b);
+__device__ static __half __hmul_sat(__half a, __half b);
+__device__ static __half __hneg(__half a);
+__device__ static __half __hsub(__half a, __half b);
+__device__ static __half __hsub_sat(__half a, __half b);
+__device__ static __half hdiv(__half a, __half b);
+
+/*
+Half2 Arithmetic Functions
+*/
+
+__device__ static __half2 __hadd2(__half2 a, __half2 b);
+__device__ static __half2 __hadd2_sat(__half2 a, __half2 b);
+__device__ static __half2 __hfma2(__half2 a, __half2 b, __half2 c);
+__device__ static __half2 __hfma2_sat(__half2 a, __half2 b, __half2 c);
+__device__ static __half2 __hmul2(__half2 a, __half2 b);
+__device__ static __half2 __hmul2_sat(__half2 a, __half2 b);
+__device__ static __half2 __hsub2(__half2 a, __half2 b);
+__device__ static __half2 __hneg2(__half2 a);
+__device__ static __half2 __hsub2_sat(__half2 a, __half2 b);
+__device__ static __half2 h2div(__half2 a, __half2 b);
+
+/*
+Half Comparision Functions
+*/
+
+__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);
+
+/*
+Half Math Functions
+*/
+
+__device__ static __half hceil(const __half h);
+__device__ static __half hcos(const __half h);
+__device__ static __half hexp(const __half h);
+__device__ static __half hexp10(const __half h);
+__device__ static __half hexp2(const __half h);
+__device__ static __half hfloor(const __half h);
+__device__ static __half hlog(const __half h);
+__device__ static __half hlog10(const __half h);
+__device__ static __half hlog2(const __half h);
+//__device__ static __half hrcp(const __half h);
+__device__ static __half hrint(const __half h);
+__device__ static __half hsin(const __half h);
+__device__ static __half hsqrt(const __half a);
+__device__ static __half htrunc(const __half a);
+
+/*
+Half2 Math Functions
+*/
+
+__device__ static __half2 h2ceil(const __half2 h);
+__device__ static __half2 h2exp(const __half2 h);
+__device__ static __half2 h2exp10(const __half2 h);
+__device__ static __half2 h2exp2(const __half2 h);
+__device__ static __half2 h2floor(const __half2 h);
+__device__ static __half2 h2log(const __half2 h);
+__device__ static __half2 h2log10(const __half2 h);
+__device__ static __half2 h2log2(const __half2 h);
+__device__ static __half2 h2rcp(const __half2 h);
+__device__ static __half2 h2rsqrt(const __half2 h);
+__device__ static __half2 h2sin(const __half2 h);
+__device__ static __half2 h2sqrt(const __half2 h);
+
+/*
+Half Conversion And Data Movement
+*/
+
+__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);
+
 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);
@@ -331,109 +506,7 @@ __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