From a6fe6222c44490476e87dd0e2b7668a82247be7f Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Tue, 13 Dec 2016 09:18:34 -0600
Subject: [PATCH] added half math addition ISA support

Change-Id: I293b771f695b499b795d7e53f600c9e4fe2a2071
---
 include/hip/hcc_detail/hip_fp16.h | 16 ++++++++++++++++
 src/hip_fp16.cpp                  |  5 +++++
 src/hip_ir.ll                     |  5 +++++
 3 files changed, 26 insertions(+)

diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h
index bcf2605f28..445df78eb4 100644
--- a/include/hip/hcc_detail/hip_fp16.h
+++ b/include/hip/hcc_detail/hip_fp16.h
@@ -25,6 +25,10 @@ THE SOFTWARE.
 
 #include "hip/hip_runtime.h"
 
+#define __CLANG_VERSION__ __clang_major__ * 10 + __clang_minor__
+
+#if __CLANG_VERSION__ == 35
+
 typedef struct{
   unsigned x: 16;
 } __half;
@@ -175,3 +179,15 @@ __device__ __half2 __lowhigh2highlow(const __half2 a);
 __device__ __half2 __low2half2(const __half2 a, const __half2 b);
 
 #endif
+
+#if __CLANG_VERSION__ == 40
+
+typedef __fp16 __half;
+extern "C" __half __hip_hadd_clang40_gfx803(__half a, __half b);
+
+__device__ inline __half __hadd(__half a, __half b){
+  return __hip_hadd_clang40_gfx803(a, b);
+}
+
+#endif
+#endif
diff --git a/src/hip_fp16.cpp b/src/hip_fp16.cpp
index 1a9d04474f..5d01f73cf7 100644
--- a/src/hip_fp16.cpp
+++ b/src/hip_fp16.cpp
@@ -22,6 +22,8 @@ THE SOFTWARE.
 
 #include"hip/hip_fp16.h"
 
+#if __CLANG_VERSION__ == 35
+
 static const unsigned sign_val = 0x8000;
 static const __half __half_value_one_float = {0x3C00};
 static const __half __half_value_zero_float = {0x0};
@@ -372,3 +374,6 @@ __device__ __half2 __lowhigh2highlow(const __half2 a){
 __device__ __half2 __low2half2(const __half2 a, const __half2 b){
   return {a.q, b.q};
 }
+
+#endif
+
diff --git a/src/hip_ir.ll b/src/hip_ir.ll
index 21123dd7c0..078dc3eed5 100644
--- a/src/hip_ir.ll
+++ b/src/hip_ir.ll
@@ -34,4 +34,9 @@ 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) {
+  %val = tail call half asm "v_add_f16 $0, $1, $2","=v,v,v"(half %a, half %b)
+  ret half %val
+}
+
 attributes #1 = { alwaysinline nounwind }