added half2 cmp and conv, data movement device functions

1. Added half2 comparision functions
2. Added conversion and data movement half apis

Change-Id: Ia33c0e957d9deb1f2b7a8fde8e22168f4d41b88b


[ROCm/clr commit: 646f566bbf]

projects/clr/hipamd/include/hip/hcc_detail/hip_fp16.h

@@ -36,7 +36,7 @@ typedef struct __attribute__((aligned(4))){
   };
 } __half2;
 
-struct holder{
+struct hipHalfHolder{
   union {
     __half h;
     unsigned short s;
@@ -45,7 +45,7 @@ struct holder{
 
 #define HINF 65504
 
-static struct holder hInf = {HINF};
+static struct hipHalfHolder __hInfValue = {HINF};
 
 extern "C" __half __hip_hc_ir_hadd_half(__half, __half);
 extern "C" __half __hip_hc_ir_hfma_half(__half, __half, __half);
@@ -180,11 +180,11 @@ __device__ static inline bool __hgt(__half a, __half b) {
 }
 
 __device__ static inline bool __hisinf(__half a) {
-  return a == hInf.s ? true : false;
+  return a == __hInfValue.h ? true : false;
 }
 
 __device__ static inline bool __hisnan(__half a) {
-  return a > hInf.s ? true : false;
+  return a > __hInfValue.h ? true : false;
 }
 
 __device__ static inline bool __hle(__half a, __half b) {
@@ -227,6 +227,389 @@ __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);
 }
 
+__device__ static inline __half2 __heq2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] == b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] == b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__ static inline __half2 __hge2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] >= b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] >= b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__ static inline __half2 __hgt2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] > b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] > b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__ static inline __half2 __hisnan2(__half2 a) {
+  __half2 c;
+  c.p[0] = (a.p[0] > __hInfValue.h) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] > __hInfValue.h) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__ static inline __half2 __hle2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] <= b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] <= b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__ static inline __half2 __hlt2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] < b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] < b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+__device__ static inline __half2 __hne2(__half2 a, __half2 b) {
+  __half2 c;
+  c.p[0] = (a.p[0] != b.p[0]) ? (__half)1 : (__half)0;
+  c.p[1] = (a.p[1] != b.p[1]) ? (__half)1 : (__half)0;
+  return c;
+}
+
+/*
+Conversion instructions
+*/
+
+__device__ static inline __half2 __float22half2_rn(const float2 a) {
+  __half2 b;
+  b.p[0] = (__half)a.x;
+  b.p[1] = (__half)a.y;
+  return b;
+}
+
+__device__ static inline __half __float2half(const float a) {
+  return (__half)a;
+}
+
+__device__ static inline __half2 __float2half2_rn(const float a) {
+  __half2 b;
+  b.p[0] = (__half)a;
+  b.p[1] = (__half)a;
+  return b;
+}
+
+__device__ static inline __half __float2half_rd(const float a) {
+  return (__half)a;
+}
+
+__device__ static inline __half __float2half_ru(const float a) {
+  return (__half)a;
+}
+
+__device__ static inline __half __float2half_rz(const float a) {
+  return (__half)a;
+}
+
+__device__ static inline __half2 __floats2half2_rn(const float a, const float b) {
+  __half2 c;
+  c.p[0] = (__half)a;
+  c.p[1] = (__half)b;
+  return c;
+}
+
+__device__ static inline float2 __half22float2(const __half2 a) {
+  float2 b;
+  b.x = (float)a.p[0];
+  b.y = (float)a.p[1];
+  return b;
+}
+
+__device__ static inline float __half2float(const __half a) {
+  return (float)a;
+}
+
+__device__ static inline __half2 half2half2(const __half a) {
+  __half2 b;
+  b.p[0] = a;
+  b.p[1] = a;
+  return b;
+}
+
+__device__ static inline int __half2int_rd(__half h) {
+  return (int)h;
+}
+
+__device__ static inline int __half2int_rn(__half h) {
+  return (int)h;
+}
+
+__device__ static inline int __half2int_ru(__half h) {
+  return (int)h;
+}
+
+__device__ static inline int __half2int_rz(__half h) {
+  return (int)h;
+}
+
+__device__ static inline long long int __half2ll_rd(__half h) {
+  return (long long int)h;
+}
+
+__device__ static inline long long int __half2ll_rn(__half h) {
+  return (long long int)h;
+}
+
+__device__ static inline long long int __half2ll_ru(__half h) {
+  return (long long int)h;
+}
+
+__device__ static inline long long int __half2ll_rz(__half h) {
+  return (long long int)h;
+}
+
+__device__ static inline short __half2short_rd(__half h) {
+  return (short)h;
+}
+
+__device__ static inline short __half2short_rn(__half h) {
+  return (short)h;
+}
+
+__device__ static inline short __half2short_ru(__half h) {
+  return (short)h;
+}
+
+__device__ static inline short __half2short_rz(__half h) {
+  return (short)h;
+}
+
+__device__ static inline unsigned int __half2uint_rd(__half h) {
+  return (unsigned int)h;
+}
+
+__device__ static inline unsigned int __half2uint_rn(__half h) {
+  return (unsigned int)h;
+}
+
+__device__ static inline unsigned int __half2uint_ru(__half h) {
+  return (unsigned int)h;
+}
+
+__device__ static inline unsigned int __half2uint_rz(__half h) {
+  return (unsigned int)h;
+}
+
+__device__ static inline unsigned long long int __half2ull_rd(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__ static inline unsigned long long int __half2ull_rn(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__ static inline unsigned long long int __half2ull_ru(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__ static inline unsigned long long int __half2ull_rz(__half h) {
+  return (unsigned long long)h;
+}
+
+__device__ static inline unsigned short int __half2ushort_rd(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__ static inline unsigned short int __half2ushort_rn(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__ static inline unsigned short int __half2ushort_ru(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__ static inline unsigned short int __half2ushort_rz(__half h) {
+  return (unsigned short int)h;
+}
+
+__device__ static inline short int __half_as_short(const __half h) {
+  hipHalfHolder hH;
+  hH.h = h;
+  return (short)hH.s;
+}
+
+__device__ static inline unsigned short int __half_as_ushort(const __half h) {
+  hipHalfHolder hH;
+  hH.h = h;
+  return hH.s;
+}
+
+__device__ static inline __half2 __halves2half2(const __half a, const __half b) {
+  __half2 c;
+  c.p[0] = a;
+  c.p[1] = b;
+  return c;
+}
+
+__device__ static inline float __high2float(const __half2 a) {
+  return (float)a.p[1];
+}
+
+__device__ static inline __half __high2half(const __half2 a) {
+  return a.p[1];
+}
+
+__device__ static inline __half2 __high2half2(const __half2 a) {
+  __half2 b;
+  b.p[0] = a.p[1];
+  b.p[1] = a.p[1];
+  return b;
+}
+
+__device__ static inline __half2 __highs2half2(const __half2 a, const __half2 b) {
+  __half2 c;
+  c.p[0] = a.p[1];
+  c.p[1] = b.p[1];
+  return c;
+}
+
+__device__ static inline __half __int2half_rd(int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __int2half_rn(int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __int2half_ru(int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __int2half_rz(int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ll2half_rd(long long int i){
+  return (__half)i;
+}
+
+__device__ static inline __half __ll2half_rn(long long int i){
+  return (__half)i;
+}
+
+__device__ static inline __half __ll2half_ru(long long int i){
+  return (__half)i;
+}
+
+__device__ static inline __half __ll2half_rz(long long int i){
+  return (__half)i;
+}
+
+__device__ static inline float __low2float(const __half2 a) {
+  return (float)a.p[0];
+}
+
+__device__ static inline __half __low2half(const __half2 a) {
+  return a.p[0];
+}
+
+__device__ static inline __half2 __low2half2(const __half2 a, const __half2 b) {
+  __half2 c;
+  c.p[0] = a.p[0];
+  c.p[1] = b.p[0];
+  return c;
+}
+
+__device__ static inline __half2 __low2half2(const __half2 a) {
+  __half2 b;
+  b.p[0] = a.p[0];
+  b.p[1] = a.p[0];
+  return b;
+}
+
+__device__ static inline __half2 __lowhigh2highlow(const __half2 a) {
+  __half2 b;
+  b.p[0] = a.p[1];
+  b.p[1] = a.p[0];
+  return b;
+}
+
+__device__ static inline __half2 __lows2half2(const __half2 a, const __half2 b) {
+  __half2 c;
+  c.p[0] = a.p[0];
+  c.p[1] = b.p[0];
+  return c;
+}
+
+__device__ static inline __half __short2half_rd(short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __short2half_rn(short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __short2half_ru(short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __short2half_rz(short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __uint2half_rd(unsigned int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __uint2half_rn(unsigned int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __uint2half_ru(unsigned int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __uint2half_rz(unsigned int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ull2half_rd(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ull2half_rn(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ull2half_ru(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ull2half_rz(unsigned long long int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ushort2half_rd(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ushort2half_rn(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ushort2half_ru(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ushort2half_rz(unsigned short int i) {
+  return (__half)i;
+}
+
+__device__ static inline __half __ushort_as_half(const unsigned short int i) {
+  hipHalfHolder hH;
+  hH.s = i;
+  return hH.h;
+}
+
 #endif
 
 #if __clang_major__ == 3

projects/clr/hipamd/tests/src/deviceLib/hipTestHalf.cpp

@@ -86,6 +86,16 @@ __global__ void CheckCmpHalf(hipLaunchParm lp, __half* In1, __half* In2, bool* O
   Out[7] = __hne(In1[7], In2[7]);
 }
 
+__global__ void CheckCmpHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __half2* Out) {
+  Out[0] = __heq2(In1[0], In2[0]);
+  Out[1] = __hge2(In1[1], In2[1]);
+  Out[2] = __hgt2(In1[2], In2[2]);
+  Out[4] = __hisnan2(In1[4]);
+  Out[5] = __hle2(In1[5], In2[5]);
+  Out[6] = __hlt2(In1[6], In2[6]);
+  Out[7] = __hne2(In1[7], In2[7]);
+}
+
 int main(){
 
 }