added half math addition ISA support

Change-Id: I293b771f695b499b795d7e53f600c9e4fe2a2071

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

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
+

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 }