From e95456eee85a4969011834a02bae1f1d6b48089d Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 16 Jan 2017 12:32:35 -0600 Subject: [PATCH] moved most of the fp16 code inside hip_fp16.cpp 1. As we use holder data structure, we move all the cmp, math, cvt apis to cpp file 2. All the tests passed 3. Add more extensive testing for half Change-Id: I92c6399dace602a0a24432728e3f2a07124e6fb1 --- hipamd/include/hip/hcc_detail/hip_fp16.h | 565 +++++------------------ hipamd/src/hip_fp16.cpp | 468 +++++++++++++++++++ 2 files changed, 572 insertions(+), 461 deletions(-) diff --git a/hipamd/include/hip/hcc_detail/hip_fp16.h b/hipamd/include/hip/hcc_detail/hip_fp16.h index 8ae6d0d85b..755cb19f6d 100644 --- a/hipamd/include/hip/hcc_detail/hip_fp16.h +++ b/hipamd/include/hip/hcc_detail/hip_fp16.h @@ -36,17 +36,6 @@ typedef struct __attribute__((aligned(4))){ }; } __half2; -struct hipHalfHolder{ - union { - __half h; - unsigned short s; - }; -}; - -#define HINF 65504 - -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); extern "C" __half __hip_hc_ir_hmul_half(__half, __half); @@ -186,456 +175,6 @@ __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 == __hInfValue.h ? true : false; -} - -__device__ static inline bool __hisnan(__half a) { - return a > __hInfValue.h ? 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); -} - -__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_rn(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; -} __device__ static inline __half hceil(const __half h) { return __hip_hc_ir_hceil_half(h); @@ -792,6 +331,110 @@ __device__ static inline __half2 h2trunc(const __half2 h) { return a; } +__device__ bool __heq(__half a, __half b); +__device__ bool __hge(__half a, __half b); +__device__ bool __hgt(__half a, __half b); +__device__ bool __hisinf(__half a); +__device__ bool __hisnan(__half a); +__device__ bool __hle(__half a, __half b); +__device__ bool __hlt(__half a, __half b); +__device__ bool __hne(__half a, __half b); +/* +Half2 Comparision Functions +*/ + +__device__ bool __hbeq2(__half2 a, __half2 b); +__device__ bool __hbge2(__half2 a, __half2 b); +__device__ bool __hbgt2(__half2 a, __half2 b); +__device__ bool __hble2(__half2 a, __half2 b); +__device__ bool __hblt2(__half2 a, __half2 b); +__device__ bool __hbne2(__half2 a, __half2 b); +__device__ __half2 __heq2(__half2 a, __half2 b); +__device__ __half2 __hge2(__half2 a, __half2 b); +__device__ __half2 __hgt2(__half2 a, __half2 b); +__device__ __half2 __hisnan2(__half2 a); +__device__ __half2 __hle2(__half2 a, __half2 b); +__device__ __half2 __hlt2(__half2 a, __half2 b); +__device__ __half2 __hne2(__half2 a, __half2 b); + +/* +Conversion instructions +*/ +__device__ __half2 __float22half2_rn(const float2 a); +__device__ __half __float2half(const float a); +__device__ __half2 __float2half2_rn(const float a); +__device__ __half __float2half_rd(const float a); +__device__ __half __float2half_rn(const float a); +__device__ __half __float2half_ru(const float a); +__device__ __half __float2half_rz(const float a); +__device__ __half2 __floats2half2_rn(const float a, const float b); +__device__ float2 __half22float2(const __half2 a); +__device__ float __half2float(const __half a); +__device__ __half2 half2half2(const __half a); +__device__ int __half2int_rd(__half h); +__device__ int __half2int_rn(__half h); +__device__ int __half2int_ru(__half h); +__device__ int __half2int_rz(__half h); +__device__ long long int __half2ll_rd(__half h); +__device__ long long int __half2ll_rn(__half h); +__device__ long long int __half2ll_ru(__half h); +__device__ long long int __half2ll_rz(__half h); +__device__ short __half2short_rd(__half h); +__device__ short __half2short_rn(__half h); +__device__ short __half2short_ru(__half h); +__device__ short __half2short_rz(__half h); +__device__ unsigned int __half2uint_rd(__half h); +__device__ unsigned int __half2uint_rn(__half h); +__device__ unsigned int __half2uint_ru(__half h); +__device__ unsigned int __half2uint_rz(__half h); +__device__ unsigned long long int __half2ull_rd(__half h); +__device__ unsigned long long int __half2ull_rn(__half h); +__device__ unsigned long long int __half2ull_ru(__half h); +__device__ unsigned long long int __half2ull_rz(__half h); +__device__ unsigned short int __half2ushort_rd(__half h); +__device__ unsigned short int __half2ushort_rn(__half h); +__device__ unsigned short int __half2ushort_ru(__half h); +__device__ unsigned short int __half2ushort_rz(__half h); +__device__ short int __half_as_short(const __half h); +__device__ unsigned short int __half_as_ushort(const __half h); +__device__ __half2 __halves2half2(const __half a, const __half b); +__device__ float __high2float(const __half2 a); +__device__ __half __high2half(const __half2 a); +__device__ __half2 __high2half2(const __half2 a); +__device__ __half2 __highs2half2(const __half2 a, const __half2 b); +__device__ __half __int2half_rd(int i); +__device__ __half __int2half_rn(int i); +__device__ __half __int2half_ru(int i); +__device__ __half __int2half_rz(int i); +__device__ __half __ll2half_rd(long long int i); +__device__ __half __ll2half_rn(long long int i); +__device__ __half __ll2half_ru(long long int i); +__device__ __half __ll2half_rz(long long int i); +__device__ float __low2float(const __half2 a); + +__device__ __half __low2half(const __half2 a); +__device__ __half2 __low2half2(const __half2 a, const __half2 b); +__device__ __half2 __low2half2(const __half2 a); +__device__ __half2 __lowhigh2highlow(const __half2 a); +__device__ __half2 __lows2half2(const __half2 a, const __half2 b); +__device__ __half __short2half_rd(short int i); +__device__ __half __short2half_rn(short int i); +__device__ __half __short2half_ru(short int i); +__device__ __half __short2half_rz(short int i); +__device__ __half __uint2half_rd(unsigned int i); +__device__ __half __uint2half_rn(unsigned int i); +__device__ __half __uint2half_ru(unsigned int i); +__device__ __half __uint2half_rz(unsigned int i); +__device__ __half __ull2half_rd(unsigned long long int i); +__device__ __half __ull2half_rn(unsigned long long int i); +__device__ __half __ull2half_ru(unsigned long long int i); +__device__ __half __ull2half_rz(unsigned long long int i); +__device__ __half __ushort2half_rd(unsigned short int i); +__device__ __half __ushort2half_rn(unsigned short int i); +__device__ __half __ushort2half_ru(unsigned short int i); +__device__ __half __ushort2half_rz(unsigned short int i); +__device__ __half __ushort_as_half(const unsigned short int i); + #endif diff --git a/hipamd/src/hip_fp16.cpp b/hipamd/src/hip_fp16.cpp index 83e0a161c7..ac79ddba08 100644 --- a/hipamd/src/hip_fp16.cpp +++ b/hipamd/src/hip_fp16.cpp @@ -22,6 +22,472 @@ THE SOFTWARE. #include"hip/hip_fp16.h" +struct hipHalfHolder{ + union { + __half h; + unsigned short s; + }; +}; + +#define HINF 65504 + +static struct hipHalfHolder __hInfValue = {HINF}; +/* +Half comparision Functions +*/ + +__device__ bool __heq(__half a, __half b) { + return a == b ? true : false; +} + +__device__ bool __hge(__half a, __half b) { + return a >= b ? true : false; +} + +__device__ bool __hgt(__half a, __half b) { + return a > b ? true : false; +} + +__device__ bool __hisinf(__half a) { + return a == __hInfValue.h ? true : false; +} + +__device__ bool __hisnan(__half a) { + return a > __hInfValue.h ? true : false; +} + +__device__ bool __hle(__half a, __half b) { + return a <= b ? true : false; +} + +__device__ bool __hlt(__half a, __half b) { + return a < b ? true : false; +} + +__device__ bool __hne(__half a, __half b) { + return a != b ? true : false; +} + +/* +Half2 Comparision Functions +*/ + +__device__ bool __hbeq2(__half2 a, __half2 b) { + return (a.p[0] == b.p[0] ? true : false) && (a.p[1] == b.p[1] ? true : false); +} + +__device__ bool __hbge2(__half2 a, __half2 b) { + return (a.p[0] >= b.p[0] ? true : false) && (a.p[1] >= b.p[1] ? true : false); +} + +__device__ bool __hbgt2(__half2 a, __half2 b) { + return (a.p[0] > b.p[0] ? true : false) && (a.p[1] > b.p[1] ? true : false); +} + +__device__ bool __hble2(__half2 a, __half2 b) { + return (a.p[0] <= b.p[0] ? true : false) && (a.p[1] <= b.p[1] ? true : false); +} + +__device__ bool __hblt2(__half2 a, __half2 b) { + return (a.p[0] < b.p[0] ? true : false) && (a.p[1] < b.p[1] ? true : false); +} + +__device__ bool __hbne2(__half2 a, __half2 b) { + return (a.p[0] != b.p[0] ? true : false) && (a.p[1] != b.p[1] ? true : false); +} + +__device__ __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__ __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__ __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__ __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__ __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__ __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__ __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__ __half2 __float22half2_rn(const float2 a) { + __half2 b; + b.p[0] = (__half)a.x; + b.p[1] = (__half)a.y; + return b; +} + +__device__ __half __float2half(const float a) { + return (__half)a; +} + +__device__ __half2 __float2half2_rn(const float a) { + __half2 b; + b.p[0] = (__half)a; + b.p[1] = (__half)a; + return b; +} + +__device__ __half __float2half_rd(const float a) { + return (__half)a; +} + +__device__ __half __float2half_rn(const float a) { + return (__half)a; +} + +__device__ __half __float2half_ru(const float a) { + return (__half)a; +} + +__device__ __half __float2half_rz(const float a) { + return (__half)a; +} + +__device__ __half2 __floats2half2_rn(const float a, const float b) { + __half2 c; + c.p[0] = (__half)a; + c.p[1] = (__half)b; + return c; +} + +__device__ float2 __half22float2(const __half2 a) { + float2 b; + b.x = (float)a.p[0]; + b.y = (float)a.p[1]; + return b; +} + +__device__ float __half2float(const __half a) { + return (float)a; +} + +__device__ __half2 half2half2(const __half a) { + __half2 b; + b.p[0] = a; + b.p[1] = a; + return b; +} + +__device__ int __half2int_rd(__half h) { + return (int)h; +} + +__device__ int __half2int_rn(__half h) { + return (int)h; +} + +__device__ int __half2int_ru(__half h) { + return (int)h; +} + +__device__ int __half2int_rz(__half h) { + return (int)h; +} + +__device__ long long int __half2ll_rd(__half h) { + return (long long int)h; +} + +__device__ long long int __half2ll_rn(__half h) { + return (long long int)h; +} + +__device__ long long int __half2ll_ru(__half h) { + return (long long int)h; +} + +__device__ long long int __half2ll_rz(__half h) { + return (long long int)h; +} + +__device__ short __half2short_rd(__half h) { + return (short)h; +} + +__device__ short __half2short_rn(__half h) { + return (short)h; +} + +__device__ short __half2short_ru(__half h) { + return (short)h; +} + +__device__ short __half2short_rz(__half h) { + return (short)h; +} + +__device__ unsigned int __half2uint_rd(__half h) { + return (unsigned int)h; +} + +__device__ unsigned int __half2uint_rn(__half h) { + return (unsigned int)h; +} + +__device__ unsigned int __half2uint_ru(__half h) { + return (unsigned int)h; +} + +__device__ unsigned int __half2uint_rz(__half h) { + return (unsigned int)h; +} + +__device__ unsigned long long int __half2ull_rd(__half h) { + return (unsigned long long)h; +} + +__device__ unsigned long long int __half2ull_rn(__half h) { + return (unsigned long long)h; +} + +__device__ unsigned long long int __half2ull_ru(__half h) { + return (unsigned long long)h; +} + +__device__ unsigned long long int __half2ull_rz(__half h) { + return (unsigned long long)h; +} + +__device__ unsigned short int __half2ushort_rd(__half h) { + return (unsigned short int)h; +} + +__device__ unsigned short int __half2ushort_rn(__half h) { + return (unsigned short int)h; +} + +__device__ unsigned short int __half2ushort_ru(__half h) { + return (unsigned short int)h; +} + +__device__ unsigned short int __half2ushort_rz(__half h) { + return (unsigned short int)h; +} + +__device__ short int __half_as_short(const __half h) { + hipHalfHolder hH; + hH.h = h; + return (short)hH.s; +} + +__device__ unsigned short int __half_as_ushort(const __half h) { + hipHalfHolder hH; + hH.h = h; + return hH.s; +} + +__device__ __half2 __halves2half2(const __half a, const __half b) { + __half2 c; + c.p[0] = a; + c.p[1] = b; + return c; +} + +__device__ float __high2float(const __half2 a) { + return (float)a.p[1]; +} + +__device__ __half __high2half(const __half2 a) { + return a.p[1]; +} + +__device__ __half2 __high2half2(const __half2 a) { + __half2 b; + b.p[0] = a.p[1]; + b.p[1] = a.p[1]; + return b; +} + +__device__ __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__ __half __int2half_rd(int i) { + return (__half)i; +} + +__device__ __half __int2half_rn(int i) { + return (__half)i; +} + +__device__ __half __int2half_ru(int i) { + return (__half)i; +} + +__device__ __half __int2half_rz(int i) { + return (__half)i; +} + +__device__ __half __ll2half_rd(long long int i){ + return (__half)i; +} + +__device__ __half __ll2half_rn(long long int i){ + return (__half)i; +} + +__device__ __half __ll2half_ru(long long int i){ + return (__half)i; +} + +__device__ __half __ll2half_rz(long long int i){ + return (__half)i; +} + +__device__ float __low2float(const __half2 a) { + return (float)a.p[0]; +} + +__device__ __half __low2half(const __half2 a) { + return a.p[0]; +} + +__device__ __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__ __half2 __low2half2(const __half2 a) { + __half2 b; + b.p[0] = a.p[0]; + b.p[1] = a.p[0]; + return b; +} + +__device__ __half2 __lowhigh2highlow(const __half2 a) { + __half2 b; + b.p[0] = a.p[1]; + b.p[1] = a.p[0]; + return b; +} + +__device__ __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__ __half __short2half_rd(short int i) { + return (__half)i; +} + +__device__ __half __short2half_rn(short int i) { + return (__half)i; +} + +__device__ __half __short2half_ru(short int i) { + return (__half)i; +} + +__device__ __half __short2half_rz(short int i) { + return (__half)i; +} + +__device__ __half __uint2half_rd(unsigned int i) { + return (__half)i; +} + +__device__ __half __uint2half_rn(unsigned int i) { + return (__half)i; +} + +__device__ __half __uint2half_ru(unsigned int i) { + return (__half)i; +} + +__device__ __half __uint2half_rz(unsigned int i) { + return (__half)i; +} + +__device__ __half __ull2half_rd(unsigned long long int i) { + return (__half)i; +} + +__device__ __half __ull2half_rn(unsigned long long int i) { + return (__half)i; +} + +__device__ __half __ull2half_ru(unsigned long long int i) { + return (__half)i; +} + +__device__ __half __ull2half_rz(unsigned long long int i) { + return (__half)i; +} + +__device__ __half __ushort2half_rd(unsigned short int i) { + return (__half)i; +} + +__device__ __half __ushort2half_rn(unsigned short int i) { + return (__half)i; +} + +__device__ __half __ushort2half_ru(unsigned short int i) { + return (__half)i; +} + +__device__ __half __ushort2half_rz(unsigned short int i) { + return (__half)i; +} + +__device__ __half __ushort_as_half(const unsigned short int i) { + hipHalfHolder hH; + hH.s = i; + return hH.h; +} + + +/* +Soft Implementation. Use it for backup. +*/ + + static const unsigned sign_val = 0x8000; static const __half __half_value_one_float = {0x3C00}; static const __half __half_value_zero_float = {0x0}; @@ -375,4 +841,6 @@ __device__ __half2 __soft_low2half2(const __half2 a, const __half2 b){ return {a.p[0], b.p[0]}; } + + #endif