added simple half math ops

Change-Id: I10b1d1023a9e5f2ba63f28c4a2bbe60ee49a8aee


[ROCm/clr commit: 01ed8e91e9]

projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h

@@ -30,10 +30,41 @@ THE SOFTWARE.
 #ifdef HIP_HALF_HW_SUPPORT
 
 typedef __fp16 __half;
-extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b);
+extern "C" __half __hip_hadd_gfx803(__half a, __half b);
+extern "C" __half __hip_hfma_gfx803(__half a, __half b);
+extern "C" __half __hip_hmul_gfx803(__half a, __half b);
+extern "C" __half __hip_hsub_gfx803(__half a, __half b);
 
-__device__ inline __half __hadd(__half a, __half b){
-  return __hip_hadd_clang40_gfx803(a, b);
+__device__ inline __half __hadd(__half a, __half b) {
+  return __hip_hadd_gfx803(a, b);
+}
+
+__device__ inline __half __hadd_sat(__half a, __half b) {
+  return __hip_add_gfx803(a, b);
+}
+
+__device__ inline __half __hfma(__half a, __half b) {
+  return __hip_hfma_gfx803(a, b);
+}
+
+__device__ inline __half __hfma_sat(__half a, __half b) {
+  return __hip_hfma_gfx803(a, b);
+}
+
+__device__ inline __half __hmul(__half a, __half b) {
+  return __hip_hmul_gfx803(a, b);
+}
+
+__device__ inline __half __hmul_sat(__half a, __half b) {
+  return __hip_hmul_gfx803(a, b);
+}
+
+__device__ inline __half __hsub(__half a, __half b) {
+  return __hip_hsub_gfx803(a, b);
+}
+
+__device__ inline __half __hsub_sat(__half a, __half b) {
+  return __hip_hsub_gfx803(a, b);
 }
 
 #else

projects/clr/hipamd/src/hip_ir.ll

@@ -34,9 +34,24 @@ define linkonce_odr spir_func i32 @__rocm_hadd(i32 %in1, i32 %in2) {
   ret i32 %val
 }
 
-define linkonce_odr spir_func half @__hip_hadd_clang40_gfx803(half %a, half %b) {
+define linkonce_odr spir_func half @__hip_hadd_gfx803(half %a, half %b) #1 {
   %val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
   ret half %val
 }
 
+define linkonce_odr spir_func half @__hip_hfma_gfx803(half %a, half %b, half %c) #1 {
+  %val = tail call half asm "v_fma_f16 $0, $1, $2, $3","=v,v,v,v"(half %a, half %b, half %c)
+  ret half %val
+}
+
+define linkonce_odr spir_func half @__hip_hmul_gfx803(half %a, half %b) #1 {
+  %val = tail call half asm "v_mul_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
+  ret half %val
+}
+
+define linkonce_odr spir_func half @__hip_hsub_gfx803(half %a, half %b) #1 {
+  %val = tail call half asm "v_sub_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
+  ret half %val
+}
+
 attributes #1 = { alwaysinline nounwind }