From a86ef4e57726fd8fed73421020f54f08062c3077 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 13 Dec 2016 20:20:58 -0600
Subject: [PATCH] added simple half math ops

Change-Id: I10b1d1023a9e5f2ba63f28c4a2bbe60ee49a8aee


[ROCm/clr commit: 01ed8e91e9c1f340c2dd62023dcc8115db756c42]
---
 .../hipamd/include/hip/hcc_detail/hip_fp16.h  | 37 +++++++++++++++++--
 projects/clr/hipamd/src/hip_ir.ll             | 17 ++++++++-
 2 files changed, 50 insertions(+), 4 deletions(-)

diff --git a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h b/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
index 735f915bd2..5acf114518 100644
--- a/projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/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
diff --git a/projects/clr/hipamd/src/hip_ir.ll b/projects/clr/hipamd/src/hip_ir.ll
index 078dc3eed5..623be19084 100644
--- a/projects/clr/hipamd/src/hip_ir.ll
+++ b/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 }