diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 638cecefb4..002fbdd3ce 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -57,6 +57,18 @@ extern "C" int __hip_hc_ir_hfma2_int(int, int, int);
 extern "C" int __hip_hc_ir_hmul2_int(int, int);
 extern "C" int __hip_hc_ir_hsub2_int(int, int);
 
+extern "C" __half __hip_hc_ir_hceil_half(__half) __asm("llvm.ceil.f16");
+extern "C" __half __hip_hc_ir_hcos_half(__half) __asm("llvm.cos.f16");
+extern "C" __half __hip_hc_ir_hexp2_half(__half) __asm("llvm.exp2.f16");
+extern "C" __half __hip_hc_ir_hfloor_half(__half) __asm("llvm.floor.f16");
+extern "C" __half __hip_hc_ir_hlog2_half(__half) __asm("llvm.log2.f16");
+extern "C" __half __hip_hc_ir_hrcp_half(__half) __asm("llvm.amdgcn.rcp.f16");
+extern "C" __half __hip_hc_ir_hrint_half(__half) __asm("llvm.rint.f16");
+extern "C" __half __hip_hc_ir_hrsqrt_half(__half) __asm("llvm.sqrt.f16");
+extern "C" __half __hip_hc_ir_hsin_half(__half) __asm("llvm.sin.f16");
+extern "C" __half __hip_hc_ir_hsqrt_half(__half) __asm("llvm.sqrt.f16");
+extern "C" __half __hip_hc_ir_htrunc_half(__half) __asm("llvm.trunc.f16");
+
 __device__ static inline __half __hadd(const __half a, const __half b) {
   return __hip_hc_ir_hadd_half(a, b);
 }
@@ -610,6 +622,66 @@ __device__ static inline __half __ushort_as_half(const unsigned short int i) {
   return hH.h;
 }
 
+__device__ static inline __half hceil(const __half h) {
+  return __hip_hc_ir_hceil_half(h);
+}
+
+__device__ static inline __half hcos(const __half h) {
+  return __hip_hc_ir_hcos_half(h);
+}
+
+__device__ static inline __half hexp(const __half h) {
+  return __hip_hc_ir_hexp2_half(__hip_hc_ir_hmul_half(h, 1.442694));
+}
+
+__device__ static inline __half hexp10(const __half h) {
+  return __hip_hc_ir_hexp2_half(__hip_hc_ir_hmul_half(h, 3.3219281));
+}
+
+__device__ static inline __half hexp2(const __half h) {
+  return __hip_hc_ir_hexp2_half(h);
+}
+
+__device__ static inline __half hfloor(const __half h) {
+  return __hip_hc_ir_hfloor_half(h);
+}
+
+__device__ static inline __half hlog(const __half h) {
+  return __hip_hc_ir_hmul_half(__hip_hc_ir_hlog2_half(h), 0.693147);
+}
+
+__device__ static inline __half hlog10(const __half h) {
+  return __hip_hc_ir_hmul_half(__hip_hc_ir_hlog2_half(h),  0.301029);
+}
+
+__device__ static inline __half hlog2(const __half h) {
+  return __hip_hc_ir_hlog2_half(h);
+}
+/*
+__device__ static inline __half hrcp(const __half h) {
+  return __hip_hc_ir_hrcp_half(h);
+}
+*/
+__device__ static inline __half hrint(const __half h) {
+  return __hip_hc_ir_hrint_half(h);
+}
+
+__device__ static inline __half hrsqrt(const __half h) {
+  return __hip_hc_ir_hrsqrt_half(h);
+}
+
+__device__ static inline __half hsin(const __half h) {
+  return __hip_hc_ir_hsin_half(h);
+}
+
+__device__ static inline __half hsqrt(const __half a) {
+  return __hip_hc_ir_hsqrt_half(a);
+}
+
+__device__ static inline __half htrunc(const __half a) {
+  return __hip_hc_ir_htrunc_half(a);
+}
+
 #endif
 
 #if __clang_major__ == 3
diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 05900259f1..55fb48cb91 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -60,6 +60,21 @@ __global__ void CheckHalf(hipLaunchParm lp, __half* In1, __half* In2, __half* In
   Out[7] = __hsub(In1[7], In2[7]);
   Out[8] = __hsub_sat(In1[8], In2[8]);
   Out[9] = hdiv(In1[9], In2[9]);
+  Out[10] = hceil(In1[10]);
+  Out[11] = hcos(In1[11]);
+  Out[12] = hexp(In1[12]);
+  Out[13] = hexp10(In1[13]);
+  Out[14] = hexp2(In1[14]);
+  Out[15] = hfloor(In1[15]);
+  Out[16] = hlog(In1[16]);
+  Out[17] = hlog10(In1[17]);
+  Out[18] = hlog2(In1[18]);
+//  Out[19] = hrcp(In1[19]);
+  Out[20] = hrint(In1[20]);
+  Out[21] = hrsqrt(In1[21]);
+  Out[22] = hsin(In1[22]);
+  Out[23] = hsqrt(In1[23]);
+  Out[24] = htrunc(In1[24]);
 }
 
 __global__ void CheckHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2* In3, __half2* Out){