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


[ROCm/clr commit: d496576b55]

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

@@ -23,15 +23,119 @@ THE SOFTWARE.
 #include <hip/hip_runtime.h>
 #include <hip/hip_vector_types.h>
 
-__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

projects/clr/hipamd/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;
 }

projects/clr/hipamd/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);

projects/clr/hipamd/src/device_functions.cpp

@@ -19,38 +19,342 @@ THE SOFTWARE.
 
 #include <hip/device_functions.h>
 
-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;
 }

projects/clr/hipamd/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<double>(x);
-}
-
-__device__ long long __double_as_longlong(double x)
-{
-  return static_cast<long long>(x);
-}
-
 __device__ void  __threadfence_system(void){
     // no-op
 }

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

@@ -128,5 +128,5 @@ __global__ void CheckCmpHalf2(hipLaunchParm lp, __half2* In1, __half2* In2, __ha
 }
 
 int main(){
-
+  passed();
 }