diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 002fbdd3ce..c54aef319b 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -69,6 +69,17 @@ 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");
 
+extern "C" int __hip_hc_ir_h2ceil_int(int);
+extern "C" int __hip_hc_ir_h2cos_int(int);
+extern "C" int __hip_hc_ir_h2exp2_int(int);
+extern "C" int __hip_hc_ir_h2floor_int(int);
+extern "C" int __hip_hc_ir_h2log2_int(int);
+extern "C" int __hip_hc_ir_h2rcp_int(int);
+extern "C" int __hip_hc_ir_h2rsqrt_int(int);
+extern "C" int __hip_hc_ir_h2sin_int(int);
+extern "C" int __hip_hc_ir_h2sqrt_int(int);
+extern "C" int __hip_hc_ir_h2trunc_int(int);
+
 __device__ static inline __half __hadd(const __half a, const __half b) {
   return __hip_hc_ir_hadd_half(a, b);
 }
@@ -682,6 +693,71 @@ __device__ static inline __half htrunc(const __half a) {
   return __hip_hc_ir_htrunc_half(a);
 }
 
+/*
+Half2 Math Operations
+*/
+
+__device__ static inline __half2 h2ceil(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2ceil_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2cos(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2cos_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2exp2(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2exp2_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2floor(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2floor_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2log2(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2log2_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2rcp(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2rcp_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2rsqrt(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2rsqrt_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2sin(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2sin_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2sqrt(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2sqrt_int(h.q);
+  return a;
+}
+
+__device__ static inline __half2 h2trunc(const __half2 h) {
+  __half2 a;
+  a.q = __hip_hc_ir_h2trunc_int(h.q);
+  return a;
+}
+
+
 #endif
 
 #if __clang_major__ == 3
diff --git a/hipamd/src/hip_ir.ll b/hipamd/src/hip_ir.ll
index 52460a38bb..739717c740 100644
--- a/hipamd/src/hip_ir.ll
+++ b/hipamd/src/hip_ir.ll
@@ -86,4 +86,64 @@ define i32 @__hip_hc_ir_hsub2_int(i32 %a, i32 %b) #1 {
   ret i32 %1
 }
 
+define i32 @__hip_hc_ir_h2ceil_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_ceil_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_ceil_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2cos_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_cos_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_cos_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2exp2_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_exp_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_exp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2floor_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_floor_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_floor_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2log2_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_log_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_log_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2rcp_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_rcp_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_rcp_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2rsqrt_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_rsq_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_rsq_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2sin_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_sin_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_sin_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2sqrt_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_sqrt_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_sqrt_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
+define i32 @__hip_hc_ir_h2trunc_int(i32 %a) #1 {
+  %1 = tail call i32 asm sideeffect "v_trunc_f16 $0, $1","=v,v"(i32 %a)
+  tail call void asm sideeffect "v_trunc_f16_sdwa $0, $1 dst_sel:WORD_1 dst_unused:UNUSED_PRESERVE src0_sel:WORD_1","v,v"(i32 %1, i32 %a)
+  ret i32 %1
+}
+
 attributes #1 = { alwaysinline nounwind }
diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 55fb48cb91..c194a42ca2 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -88,6 +88,21 @@ __global__ void CheckHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2
   Out[7] = __hsub2(In1[7], In2[7]);
   Out[8] = __hsub2_sat(In1[8], In2[8]);
   Out[9] = h2div(In1[9], In2[9]);
+  Out[10] = h2ceil(In1[10]);
+  Out[11] = h2cos(In1[11]);
+//  Out[12] = h2exp(In1[12]);
+//  Out[13] = h2exp10(In1[13]);
+  Out[14] = h2exp2(In1[14]);
+  Out[15] = h2floor(In1[15]);
+//  Out[16] = h2log(In1[16]);
+//  Out[17] = h2log10(In1[17]);
+  Out[18] = h2log2(In1[18]);
+  Out[19] = h2rcp(In1[19]);
+//  Out[20] = h2rint(In1[20]);
+  Out[21] = h2rsqrt(In1[21]);
+  Out[22] = h2sin(In1[22]);
+  Out[23] = h2sqrt(In1[23]);
+  Out[24] = h2trunc(In1[24]);
 }
 
 __global__ void CheckCmpHalf(hipLaunchParm lp, __half* In1, __half* In2, bool* Out) {
@@ -109,6 +124,7 @@ __global__ void CheckCmpHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __ha
   Out[5] = __hle2(In1[5], In2[5]);
   Out[6] = __hlt2(In1[6], In2[6]);
   Out[7] = __hne2(In1[7], In2[7]);
+
 }
 
 int main(){