From 6f2cfddc670835a3965afe252352c488680b846d Mon Sep 17 00:00:00 2001 From: Aditya Atluri Date: Mon, 16 Jan 2017 12:10:05 -0600 Subject: [PATCH] Added type conversion intrinsics 1. Added all type conversion intrinsics 2. NO TESTS have been added. (Will add in next commit) 3. Sanatized code in hip_runtime.h 4. Added passed() to hipTestHalf to make it pass on HIT Change-Id: I0987963c802fc7ff4d7e07d7b88d86da35da53c9 --- include/hip/hcc_detail/device_functions.h | 112 ++++++- include/hip/hcc_detail/hip_fp16.h | 4 + include/hip/hcc_detail/hip_runtime.h | 30 +- src/device_functions.cpp | 362 ++++++++++++++++++++-- src/device_util.cpp | 10 - tests/src/deviceLib/hipTestHalf.cpp | 2 +- 6 files changed, 462 insertions(+), 58 deletions(-) diff --git a/include/hip/hcc_detail/device_functions.h b/include/hip/hcc_detail/device_functions.h index e2b061c640..e6771650a2 100644 --- a/include/hip/hcc_detail/device_functions.h +++ b/include/hip/hcc_detail/device_functions.h @@ -23,15 +23,119 @@ THE SOFTWARE. #include #include -__device__ float __int_as_float (int x); +/* +Rounding modes are not yet supported in HIP +*/ -__device__ double __hiloint2double (int hi, int lo); +__device__ float __double2float_rd(double x); +__device__ float __double2float_rn(double x); +__device__ float __double2float_ru(double x); +__device__ float __double2float_rz(double x); + +__device__ int __double2hiint(double x); +__device__ int __double2loint(double x); + +__device__ int __double2int_rd(double x); +__device__ int __double2int_rn(double x); +__device__ int __double2int_ru(double x); +__device__ int __double2int_rz(double x); + +__device__ long long int __double2ll_rd(double x); +__device__ long long int __double2ll_rn(double x); +__device__ long long int __double2ll_ru(double x); +__device__ long long int __double2ll_rz(double x); + +__device__ unsigned int __double2uint_rd(double x); +__device__ unsigned int __double2uint_rn(double x); +__device__ unsigned int __double2uint_ru(double x); +__device__ unsigned int __double2uint_rz(double x); + +__device__ unsigned long long int __double2ull_rd(double x); +__device__ unsigned long long int __double2ull_rn(double x); +__device__ unsigned long long int __double2ull_ru(double x); +__device__ unsigned long long int __double2ull_rz(double x); + +__device__ long long int __double_as_longlong(double x); +/* +__device__ unsigned short __float2half_rn(float x); +__device__ float __half2float(unsigned short); + +The above device function are not a valid . +Use +__device__ __half __float2half_rn(float x); +__device__ float __half2float(__half); +from hip_fp16.h + +CUDA implements half as unsigned short whereas, HIP doesn't. + +*/ + +__device__ int float2int_rd(float x); +__device__ int float2int_rn(float x); +__device__ int float2int_ru(float x); +__device__ int float2int_rz(float x); + +__device__ long long int __float2ll_rd(float x); +__device__ long long int __float2ll_rn(float x); +__device__ long long int __float2ll_ru(float x); +__device__ long long int __float2ll_rz(float x); + +__device__ unsigned int __float2uint_rd(float x); +__device__ unsigned int __float2uint_rn(float x); +__device__ unsigned int __float2uint_ru(float x); +__device__ unsigned int __float2uint_rz(float x); + +__device__ unsigned long long int __float2ull_rd(float x); +__device__ unsigned long long int __float2ull_rn(float x); +__device__ unsigned long long int __float2ull_ru(float x); +__device__ unsigned long long int __float2ull_rz(float x); + +__device__ int __float_as_int(float x); +__device__ unsigned int __float_as_uint(float x); +__device__ double __hiloint2double(int hi, int lo); +__device__ double __int2double_rn(int x); + +__device__ float __int2float_rd(int x); +__device__ float __int2float_rn(int x); +__device__ float __int2float_ru(int x); +__device__ float __int2float_rz(int x); + +__device__ float __int_as_float(int x); + +__device__ double __ll2double_rd(long long int x); +__device__ double __ll2double_rn(long long int x); +__device__ double __ll2double_ru(long long int x); +__device__ double __ll2double_rz(long long int x); + +__device__ float __ll2float_rd(long long int x); +__device__ float __ll2float_rn(long long int x); +__device__ float __ll2float_ru(long long int x); +__device__ float __ll2float_rz(long long int x); + +__device__ double __longlong_as_double(long long int x); + +__device__ double __uint2double_rn(int x); + +__device__ float __uint2float_rd(unsigned int x); +__device__ float __uint2float_rn(unsigned int x); +__device__ float __uint2float_ru(unsigned int x); +__device__ float __uint2float_rz(unsigned int x); + +__device__ float __uint_as_float(unsigned int x); + +__device__ double __ull2double_rd(unsigned long long int x); +__device__ double __ull2double_rn(unsigned long long int x); +__device__ double __ull2double_ru(unsigned long long int x); +__device__ double __ull2double_rz(unsigned long long int x); + +__device__ float __ull2float_rd(unsigned long long int x); +__device__ float __ull2float_rn(unsigned long long int x); +__device__ float __ull2float_ru(unsigned long long int x); +__device__ float __ull2float_rz(unsigned long long int x); __device__ char4 __hip_hc_add8pk(char4, char4); __device__ char4 __hip_hc_sub8pk(char4, char4); __device__ char4 __hip_hc_mul8pk(char4, char4); -extern __device__ double __longlong_as_double(long long int x); -extern __device__ long long int __double_as_longlong(double x); #endif diff --git a/include/hip/hcc_detail/hip_fp16.h b/include/hip/hcc_detail/hip_fp16.h index eb63c7cd60..8ae6d0d85b 100644 --- a/include/hip/hcc_detail/hip_fp16.h +++ b/include/hip/hcc_detail/hip_fp16.h @@ -325,6 +325,10 @@ __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; } diff --git a/include/hip/hcc_detail/hip_runtime.h b/include/hip/hcc_detail/hip_runtime.h index f747b446d7..f595ff1c05 100644 --- a/include/hip/hcc_detail/hip_runtime.h +++ b/include/hip/hcc_detail/hip_runtime.h @@ -136,12 +136,10 @@ __device__ float cyl_bessel_i0f(float x); __device__ float cyl_bessel_i1f(float x); __device__ float erfcf(float x); __device__ float erfcinvf(float y); -__host__ float erfcinvf(float y); + __device__ float erfcxf(float x); -__host__ float erfcxf(float x); __device__ float erff(float x); __device__ float erfinvf(float y); -__host__ float erfinvf(float y); __device__ float exp2f(float x); __device__ float expm1f(float x); __device__ float fabsf(float x); @@ -174,34 +172,24 @@ __device__ float nanf(const char* tagp); __device__ float nearbyintf(float x); __device__ float nextafterf(float x, float y); __device__ float norm3df(float a, float b, float c); -__host__ float norm3df(float a, float b, float c); __device__ float norm4df(float a, float b, float c, float d); -__host__ float norm4df(float a, float b, float c, float d); __device__ float normcdff(float y); -__host__ float normcdff(float y); __device__ float normcdfinvf(float y); -__host__ float normcdfinvf(float y); __device__ float normf(int dim, const float *a); __device__ float rcbrtf(float x); -__host__ float rcbrtf(float x); __device__ float remainderf(float x, float y); __device__ float remquof(float x, float y, int *quo); __device__ float rhypotf(float x, float y); -__host__ float rhypotf(float x, float y); __device__ float rintf(float x); __device__ float rnorm3df(float a, float b, float c); -__host__ float rnorm3df(float a, float b, float c); __device__ float rnorm4df(float a, float b, float c, float d); -__host__ float rnorm4df(float a, float b, float c, float d); __device__ float rnormf(int dim, const float* a); -__host__ float rnormf(int dim, const float* a); __device__ float roundf(float x); __device__ float rsqrtf(float x); __device__ float scalblnf(float x, long int n); __device__ float scalbnf(float x, int n); __host__ __device__ unsigned signbit(float a); __device__ void sincospif(float x, float *sptr, float *cptr); -__host__ void sincospif(float x, float *sptr, float *cptr); __device__ float sinhf(float x); __device__ float sinpif(float x); __device__ float sqrtf(float x); @@ -214,8 +202,22 @@ __device__ float ynf(int n, float x); __host__ __device__ float cospif(float x); __host__ __device__ float sinpif(float x); -__device__ float sqrtf(float x); +// /__device__ float sqrtf(float x); __host__ __device__ float rsqrtf(float x); +__host__ float normcdff(float y); + +__host__ float erfcinvf(float y); +__host__ float erfcxf(float x); +__host__ float erfinvf(float y); +__host__ float norm3df(float a, float b, float c); +__host__ float normcdfinvf(float y); +__host__ float norm4df(float a, float b, float c, float d); +__host__ float rcbrtf(float x); +__host__ float rhypotf(float x, float y); +__host__ float rnorm3df(float a, float b, float c); +__host__ float rnormf(int dim, const float* a); +__host__ float rnorm4df(float a, float b, float c, float d); +__host__ void sincospif(float x, float *sptr, float *cptr); __device__ double acos(double x); __device__ double acosh(double x); diff --git a/src/device_functions.cpp b/src/device_functions.cpp index 0de0cf7f6b..01d18d8d28 100644 --- a/src/device_functions.cpp +++ b/src/device_functions.cpp @@ -19,38 +19,342 @@ THE SOFTWARE. #include -extern "C" float __hip_int_as_float(int); +struct holder64Bit{ + union{ + double d; + unsigned long int uli; + signed long int sli; + signed int si[2]; + unsigned int ui[2]; + }; +} __attribute__((aligned(8))); -typedef struct { - signed int hi; - signed int lo; -} __hip_signed_2; +struct holder32Bit { + union { + float f; + unsigned int ui; + signed int si; + }; +} __attribute__((aligned(4))); -typedef struct { -union { - double d; - long long int lli; - __hip_signed_2 s2; -} ; -} __hip_64bit_struct; +struct holder64Bit hold64; +struct holder32Bit hold32; -typedef struct { -union { - float f; - unsigned int ui; - signed int si; -}; -} __hip_32bit_struct; - -__device__ float __int_as_float (int x) { - __hip_32bit_struct s; - s.si = x; - return s.f; +__device__ float __double2float_rd(double x) +{ + return (double)x; +} +__device__ float __double2float_rn(double x) +{ + return (double)x; +} +__device__ float __double2float_ru(double x) +{ + return (double)x; +} +__device__ float __double2float_rz(double x) +{ + return (double)x; } -__device__ double __hiloint2double (int hi, int lo) { - __hip_64bit_struct s; - s.s2.hi = hi; - s.s2.lo = lo; - return s.d; + +__device__ int __double2hiint(double x) +{ + hold64.d = x; + return hold64.si[1]; +} +__device__ int __double2loint(double x) +{ + hold64.d = x; + return hold64.si[0]; +} + + +__device__ int __double2int_rd(double x) +{ + return (int)x; +} +__device__ int __double2int_rn(double x) +{ + return (int)x; +} +__device__ int __double2int_ru(double x) +{ + return (int)x; +} +__device__ int __double2int_rz(double x) +{ + return (int)x; +} + +__device__ long long int __double2ll_rd(double x) +{ + return (long long int)x; +} +__device__ long long int __double2ll_rn(double x) +{ + return (long long int)x; +} +__device__ long long int __double2ll_ru(double x) +{ + return (long long int)x; +} +__device__ long long int __double2ll_rz(double x) +{ + return (long long int)x; +} + + +__device__ unsigned int __double2uint_rd(double x) +{ + return (unsigned int)x; +} +__device__ unsigned int __double2uint_rn(double x) +{ + return (unsigned int)x; +} +__device__ unsigned int __double2uint_ru(double x) +{ + return (unsigned int)x; +} +__device__ unsigned int __double2uint_rz(double x) +{ + return (unsigned int)x; +} + +__device__ unsigned long long int __double2ull_rd(double x) +{ + return (unsigned long long int)x; +} +__device__ unsigned long long int __double2ull_rn(double x) +{ + return (unsigned long long int)x; +} +__device__ unsigned long long int __double2ull_ru(double x) +{ + return (unsigned long long int)x; +} +__device__ unsigned long long int __double2ull_rz(double x) +{ + return (unsigned long long int)x; +} + +__device__ long long int __double_as_longlong(double x) +{ + hold64.d = x; + return hold64.sli; +} + +__device__ int float2int_rd(float x) +{ + return (int)x; +} +__device__ int float2int_rn(float x) +{ + return (int)x; +} +__device__ int float2int_ru(float x) +{ + return (int)x; +} +__device__ int float2int_rz(float x) +{ + return (int)x; +} + +__device__ long long int __float2ll_rd(float x) +{ + return (long long int)x; +} +__device__ long long int __float2ll_rn(float x) +{ + return (long long int)x; +} +__device__ long long int __float2ll_ru(float x) +{ + return (long long int)x; +} +__device__ long long int __float2ll_rz(float x) +{ + return (long long int)x; +} + +__device__ unsigned int __float2uint_rd(float x) +{ + return (unsigned int)x; +} +__device__ unsigned int __float2uint_rn(float x) +{ + return (unsigned int)x; +} +__device__ unsigned int __float2uint_ru(float x) +{ + return (unsigned int)x; +} +__device__ unsigned int __float2uint_rz(float x) +{ + return (unsigned int)x; +} + +__device__ unsigned long long int __float2ull_rd(float x) +{ + return (unsigned long long int)x; +} +__device__ unsigned long long int __float2ull_rn(float x) +{ + return (unsigned long long int)x; +} +__device__ unsigned long long int __float2ull_ru(float x) +{ + return (unsigned long long int)x; +} +__device__ unsigned long long int __float2ull_rz(float x) +{ + return (unsigned long long int)x; +} + +__device__ int __float_as_int(float x) +{ + hold32.f = x; + return hold32.si; +} +__device__ unsigned int __float_as_uint(float x) +{ + hold32.f = x; + return hold32.ui; +} +__device__ double __hiloint2double(int hi, int lo) +{ + hold64.si[1] = hi; + hold64.si[0] = lo; + return hold64.d; +} +__device__ double __int2double_rn(int x) +{ + return (double)x; +} + +__device__ float __int2float_rd(int x) +{ + return (float)x; +} +__device__ float __int2float_rn(int x) +{ + return (float)x; +} +__device__ float __int2float_ru(int x) +{ + return (float)x; +} +__device__ float __int2float_rz(int x) +{ + return (float)x; +} + +__device__ float __int_as_float(int x) +{ + hold32.si = x; + return hold32.f; +} + +__device__ double __ll2double_rd(long long int x) +{ + return (double)x; +} +__device__ double __ll2double_rn(long long int x) +{ + return (double)x; +} +__device__ double __ll2double_ru(long long int x) +{ + return (double)x; +} +__device__ double __ll2double_rz(long long int x) +{ + return (double)x; +} + +__device__ float __ll2float_rd(long long int x) +{ + return (float)x; +} +__device__ float __ll2float_rn(long long int x) +{ + return (float)x; +} +__device__ float __ll2float_ru(long long int x) +{ + return (float)x; +} +__device__ float __ll2float_rz(long long int x) +{ + return (float)x; +} + +__device__ double __longlong_as_double(long long int x) +{ + hold64.sli = x; + return hold64.d; +} + +__device__ double __uint2double_rn(int x) +{ + return (double)x; +} + +__device__ float __uint2float_rd(unsigned int x) +{ + return (float)x; +} +__device__ float __uint2float_rn(unsigned int x) +{ + return (float)x; +} +__device__ float __uint2float_ru(unsigned int x) +{ + return (float)x; +} +__device__ float __uint2float_rz(unsigned int x) +{ + return (float)x; +} + +__device__ float __uint_as_float(unsigned int x) +{ + hold32.ui = x; + return hold32.f; +} + +__device__ double __ull2double_rd(unsigned long long int x) +{ + return (double)x; +} +__device__ double __ull2double_rn(unsigned long long int x) +{ + return (double)x; +} +__device__ double __ull2double_ru(unsigned long long int x) +{ + return (double)x; +} +__device__ double __ull2double_rz(unsigned long long int x) +{ + return (double)x; +} + +__device__ float __ull2float_rd(unsigned long long int x) +{ + return (float)x; +} +__device__ float __ull2float_rn(unsigned long long int x) +{ + return (float)x; +} +__device__ float __ull2float_ru(unsigned long long int x) +{ + return (float)x; +} +__device__ float __ull2float_rz(unsigned long long int x) +{ + return (float)x; } diff --git a/src/device_util.cpp b/src/device_util.cpp index db0f494af4..cf5d6ff5af 100644 --- a/src/device_util.cpp +++ b/src/device_util.cpp @@ -2179,16 +2179,6 @@ __device__ double __hip_fast_dsqrt_rz(double x) { return hc::fast_math::sqrt(x); } -__device__ double __longlong_as_double(long long int x) -{ - return static_cast(x); -} - -__device__ long long __double_as_longlong(double x) -{ - return static_cast(x); -} - __device__ void __threadfence_system(void){ // no-op } diff --git a/tests/src/deviceLib/hipTestHalf.cpp b/tests/src/deviceLib/hipTestHalf.cpp index 975a494b2a..12b7e2e270 100644 --- a/tests/src/deviceLib/hipTestHalf.cpp +++ b/tests/src/deviceLib/hipTestHalf.cpp @@ -128,5 +128,5 @@ __global__ void CheckCmpHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __ha } int main(){ - + passed(); }