From 89998d436fc9aa4bde00ee9a4cfa5c6e4e6deefd Mon Sep 17 00:00:00 2001
From: Aditya Atluri <Aditya.Atluri@amd.com>
Date: Thu, 12 Jan 2017 14:52:14 -0600
Subject: [PATCH] added comparision device functions for fp16 1. Added
 comparision device functions 2. Added test to check correct isa getting
 generated

Change-Id: I16732f5a1438bdce145f7bfcecd28198e3cc4b79
---
 hipamd/include/hip/hcc_detail/hip_fp16.h   | 75 ++++++++++++++++++++++
 hipamd/tests/src/deviceLib/hipTestHalf.cpp | 18 ++++--
 2 files changed, 86 insertions(+), 7 deletions(-)

diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h
index 2ef8d330e7..e4408556f1 100644
--- a/hipamd/include/hip/hcc_detail/hip_fp16.h
+++ b/hipamd/include/hip/hcc_detail/hip_fp16.h
@@ -36,6 +36,17 @@ typedef struct __attribute__((aligned(4))){
   };
 } __half2;
 
+struct holder{
+  union {
+    __half h;
+    unsigned short s;
+  };
+};
+
+#define HINF 65504
+
+static struct holder hInf = {HINF};
+
 extern "C" __half __hip_hc_ir_hadd_half(__half, __half);
 extern "C" __half __hip_hc_ir_hfma_half(__half, __half, __half);
 extern "C" __half __hip_hc_ir_hmul_half(__half, __half);
@@ -152,6 +163,70 @@ __device__ static inline __half2 h2div(__half2 a, __half2 b) {
   return c;
 }
 
+/*
+Half comparision Functions
+*/
+
+__device__ static inline bool __heq(__half a, __half b) {
+  return a == b ? true : false;
+}
+
+__device__ static inline bool __hge(__half a, __half b) {
+  return a >= b ? true : false;
+}
+
+__device__ static inline bool __hgt(__half a, __half b) {
+  return a > b ? true : false;
+}
+
+__device__ static inline bool __hisinf(__half a) {
+  return a == hInf.s ? true : false;
+}
+
+__device__ static inline bool __hisnan(__half a) {
+  return a > hInf.s ? true : false;
+}
+
+__device__ static inline bool __hle(__half a, __half b) {
+  return a <= b ? true : false;
+}
+
+__device__ static inline bool __hlt(__half a, __half b) {
+  return a < b ? true : false;
+}
+
+__device__ static inline bool __hne(__half a, __half b) {
+  return a != b ? true : false;
+}
+
+/*
+Half2 Comparision Functions
+*/
+
+__device__ static inline bool __hbeq2(__half2 a, __half2 b) {
+  return (a.p[0] == b.p[0] ? true : false) && (a.p[1] == b.p[1] ? true : false);
+}
+
+__device__ static inline bool __hbge2(__half2 a, __half2 b) {
+  return (a.p[0] >= b.p[0] ? true : false) && (a.p[1] >= b.p[1] ? true : false);
+}
+
+__device__ static inline bool __hbgt2(__half2 a, __half2 b) {
+  return (a.p[0] > b.p[0] ? true : false) && (a.p[1] > b.p[1] ? true : false);
+}
+
+__device__ static inline bool __hble2(__half2 a, __half2 b) {
+  return (a.p[0] <= b.p[0] ? true : false) && (a.p[1] <= b.p[1] ? true : false);
+}
+
+__device__ static inline bool __hblt2(__half2 a, __half2 b) {
+  return (a.p[0] < b.p[0] ? true : false) && (a.p[1] < b.p[1] ? true : false);
+}
+
+__device__ static inline bool __hbne2(__half2 a, __half2 b) {
+  return (a.p[0] != b.p[0] ? true : false) && (a.p[1] != b.p[1] ? true : false);
+}
+
 #endif
 
 #if __clang_major__ == 3
diff --git a/hipamd/tests/src/deviceLib/hipTestHalf.cpp b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
index 2c01c5cb72..46927c3902 100644
--- a/hipamd/tests/src/deviceLib/hipTestHalf.cpp
+++ b/hipamd/tests/src/deviceLib/hipTestHalf.cpp
@@ -49,13 +49,6 @@ THE SOFTWARE.
 #define hNeg1    0xBC00
 #define hNeg1PK 0xBC00BC00
 
-struct holder{
-union{
-  __half a;
-  unsigned short b;
-};
-};
-
 __global__ void CheckHalf(hipLaunchParm lp, __half* In1, __half* In2, __half* In3, __half* Out){
   Out[0] = __hadd(In1[0], In2[0]);
   Out[1] = __hadd_sat(In1[1], In2[1]);
@@ -82,6 +75,17 @@ __global__ void CheckHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2
   Out[9] = h2div(In1[9], In2[9]);
 }
 
+__global__ void CheckCmpHalf(hipLaunchParm lp, __half* In1, __half* In2, bool* Out) {
+  Out[0] = __heq(In1[0], In2[0]);
+  Out[1] = __hge(In1[1], In2[1]);
+  Out[2] = __hgt(In1[2], In2[2]);
+  Out[3] = __hisinf(In1[3]);
+  Out[4] = __hisnan(In1[4]);
+  Out[5] = __hle(In1[5], In2[5]);
+  Out[6] = __hlt(In1[6], In2[6]);
+  Out[7] = __hne(In1[7], In2[7]);
+}
+
 int main(){
 
 }