From fe38e9652baf47595200344ebf92ddeb9ff17c19 Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Fri, 13 Jan 2017 12:05:29 -0600
Subject: [PATCH] added math functions for half 1. Added math functions for
 half precision 2. HRCP is not available due to device code linking errors,
 will be enabled once it is fixed 3. Added math functions to half test file

Change-Id: Ie317ce70ef518a4fc3f27142143d01e0327f5df3
---
 hipamd/include/hip/hcc_detail/hip_fp16.h   | 72 ++++++++++++++++++++++
 hipamd/tests/src/deviceLib/hipTestHalf.cpp | 15 +++++
 2 files changed, 87 insertions(+)

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 638cecefb4..002fbdd3ce 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -57,6 +57,18 @@ extern "C" int __hip_hc_ir_hfma2_int(int, int, int);
 extern "C" int __hip_hc_ir_hmul2_int(int, int);
 extern "C" int __hip_hc_ir_hsub2_int(int, int);
 
+extern "C" __half __hip_hc_ir_hceil_half(__half) __asm("llvm.ceil.f16");
+extern "C" __half __hip_hc_ir_hcos_half(__half) __asm("llvm.cos.f16");
+extern "C" __half __hip_hc_ir_hexp2_half(__half) __asm("llvm.exp2.f16");
+extern "C" __half __hip_hc_ir_hfloor_half(__half) __asm("llvm.floor.f16");
+extern "C" __half __hip_hc_ir_hlog2_half(__half) __asm("llvm.log2.f16");
+extern "C" __half __hip_hc_ir_hrcp_half(__half) __asm("llvm.amdgcn.rcp.f16");
+extern "C" __half __hip_hc_ir_hrint_half(__half) __asm("llvm.rint.f16");
+extern "C" __half __hip_hc_ir_hrsqrt_half(__half) __asm("llvm.sqrt.f16");
+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");
+
 __device__ static inline __half __hadd(const __half a, const __half b) {
   return __hip_hc_ir_hadd_half(a, b);
 }
@@ -610,6 +622,66 @@ __device__ static inline __half __ushort_as_half(const unsigned short int i) {
   return hH.h;
 }
 
+__device__ static inline __half hceil(const __half h) {
+  return __hip_hc_ir_hceil_half(h);
+}
+
+__device__ static inline __half hcos(const __half h) {
+  return __hip_hc_ir_hcos_half(h);
+}
+
+__device__ static inline __half hexp(const __half h) {
+  return __hip_hc_ir_hexp2_half(__hip_hc_ir_hmul_half(h, 1.442694));
+}
+
+__device__ static inline __half hexp10(const __half h) {
+  return __hip_hc_ir_hexp2_half(__hip_hc_ir_hmul_half(h, 3.3219281));
+}
+
+__device__ static inline __half hexp2(const __half h) {
+  return __hip_hc_ir_hexp2_half(h);
+}
+
+__device__ static inline __half hfloor(const __half h) {
+  return __hip_hc_ir_hfloor_half(h);
+}
+
+__device__ static inline __half hlog(const __half h) {
+  return __hip_hc_ir_hmul_half(__hip_hc_ir_hlog2_half(h), 0.693147);
+}
+
+__device__ static inline __half hlog10(const __half h) {
+  return __hip_hc_ir_hmul_half(__hip_hc_ir_hlog2_half(h),  0.301029);
+}
+
+__device__ static inline __half hlog2(const __half h) {
+  return __hip_hc_ir_hlog2_half(h);
+}
+/*
+__device__ static inline __half hrcp(const __half h) {
+  return __hip_hc_ir_hrcp_half(h);
+}
+*/
+__device__ static inline __half hrint(const __half h) {
+  return __hip_hc_ir_hrint_half(h);
+}
+
+__device__ static inline __half hrsqrt(const __half h) {
+  return __hip_hc_ir_hrsqrt_half(h);
+}
+
+__device__ static inline __half hsin(const __half h) {
+  return __hip_hc_ir_hsin_half(h);
+}
+
+__device__ static inline __half hsqrt(const __half a) {
+  return __hip_hc_ir_hsqrt_half(a);
+}
+
+__device__ static inline __half htrunc(const __half a) {
+  return __hip_hc_ir_htrunc_half(a);
+}
+
 #endif
 
 #if __clang_major__ == 3
diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 05900259f1..55fb48cb91 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -60,6 +60,21 @@ __global__ void CheckHalf(hipLaunchParm lp, __half* In1, __half* In2, __half* In
   Out[7] = __hsub(In1[7], In2[7]);
   Out[8] = __hsub_sat(In1[8], In2[8]);
   Out[9] = hdiv(In1[9], In2[9]);
+  Out[10] = hceil(In1[10]);
+  Out[11] = hcos(In1[11]);
+  Out[12] = hexp(In1[12]);
+  Out[13] = hexp10(In1[13]);
+  Out[14] = hexp2(In1[14]);
+  Out[15] = hfloor(In1[15]);
+  Out[16] = hlog(In1[16]);
+  Out[17] = hlog10(In1[17]);
+  Out[18] = hlog2(In1[18]);
+//  Out[19] = hrcp(In1[19]);
+  Out[20] = hrint(In1[20]);
+  Out[21] = hrsqrt(In1[21]);
+  Out[22] = hsin(In1[22]);
+  Out[23] = hsqrt(In1[23]);
+  Out[24] = htrunc(In1[24]);
 }
 
 __global__ void CheckHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2* In3, __half2* Out){