fixed compilation issues for vector types and math functions

1. Added math_functions.h to hip_runtime.h
2. Changed operator overloading classifier static to static inline
3. Added vector types test for gpu
4. Seperated __host__ and __device__ for math functions in headers

Change-Id: I499862fad5d7b10da686da9011d7ecefe523f8e2


[ROCm/clr commit: 02190736e3]

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

@@ -69,6 +69,8 @@ extern int HIP_TRACE_API;
 #include <hip/hcc_detail/hip_ldg.h>
 #endif
 #include <hip/hcc_detail/host_defines.h>
+#include <hip/hcc_detail/math_functions.h>
+
 // TODO-HCC remove old definitions ; ~1602 hcc supports __HCC_ACCELERATOR__ define.
 #if defined (__KALMAR_ACCELERATOR__) && !defined (__HCC_ACCELERATOR__)
 #define __HCC_ACCELERATOR__  __KALMAR_ACCELERATOR__

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

@@ -1228,20 +1228,20 @@ DECLOP_MAKE_FOUR_COMPONENT(signed long, longlong4);
 #if __cplusplus
 
 #define DECLOP_1VAR_2IN_1OUT(type, op) \
-__device__ __host__ static type operator op (const type& lhs, const type& rhs) { \
+__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs.x op rhs.x; \
   return ret; \
 }
 
 #define DECLOP_1VAR_SCALE_PRODUCT(type, type1) \
-__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
+__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
   type ret; \
   ret.x = lhs.x * rhs; \
   return ret; \
 } \
 \
-__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
+__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs * rhs.x; \
   return ret; \
@@ -1260,7 +1260,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
 }
 
 #define DECLOP_1VAR_POSTOP(type, op) \
-__device__ __host__ static type operator op (type& val, int i) { \
+__device__ __host__ static inline type operator op (type& val, int i) { \
   type ret; \
   ret.x = val.x; \
   val.x op; \
@@ -1273,7 +1273,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
 }
 
 #define DECLOP_1VAR_1IN_1OUT(type, op) \
-__device__ __host__ static type operator op(type& rhs) { \
+__device__ __host__ static inline type operator op(type& rhs) { \
   type ret; \
   ret.x = op rhs.x; \
   return ret; \
@@ -1289,7 +1289,7 @@ __device__ __host__ static inline bool operator op (type& rhs) { \
 */
 
 #define DECLOP_2VAR_2IN_1OUT(type, op) \
-__device__ __host__ static type operator op (const type& lhs, const type& rhs) { \
+__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs.x op rhs.x; \
   ret.y = lhs.y op rhs.y; \
@@ -1297,14 +1297,14 @@ __device__ __host__ static type operator op (const type& lhs, const type& rhs) {
 }
 
 #define DECLOP_2VAR_SCALE_PRODUCT(type, type1) \
-__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
+__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
   type ret; \
   ret.x = lhs.x * rhs; \
   ret.y = lhs.y * rhs; \
   return ret; \
 } \
 \
-__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
+__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs * rhs.x; \
   ret.y = lhs * rhs.y; \
@@ -1326,7 +1326,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
 }
 
 #define DECLOP_2VAR_POSTOP(type, op) \
-__device__ __host__ static type operator op (type& val, int i) { \
+__device__ __host__ static inline type operator op (type& val, int i) { \
   type ret; \
   ret.x = val.x; \
   ret.y = val.y; \
@@ -1341,7 +1341,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
 }
 
 #define DECLOP_2VAR_1IN_1OUT(type, op) \
-__device__ __host__ static type operator op(type &rhs) { \
+__device__ __host__ static inline type operator op(type &rhs) { \
   type ret; \
   ret.x = op rhs.x; \
   ret.y = op rhs.y; \
@@ -1359,7 +1359,7 @@ __device__ __host__ static inline bool operator op (type &rhs) { \
 */
 
 #define DECLOP_3VAR_2IN_1OUT(type, op) \
-__device__ __host__ static type operator op (const type& lhs, const type& rhs) { \
+__device__ __host__ static inline type operator op (const type& lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs.x op rhs.x; \
   ret.y = lhs.y op rhs.y; \
@@ -1368,7 +1368,7 @@ __device__ __host__ static type operator op (const type& lhs, const type& rhs) {
 }
 
 #define DECLOP_3VAR_SCALE_PRODUCT(type, type1) \
-__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
+__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
   type ret; \
   ret.x = lhs.x * rhs; \
   ret.y = lhs.y * rhs; \
@@ -1376,7 +1376,7 @@ __device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
   return ret; \
 } \
 \
-__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
+__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs * rhs.x; \
   ret.y = lhs * rhs.y; \
@@ -1401,7 +1401,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
 }
 
 #define DECLOP_3VAR_POSTOP(type, op) \
-__device__ __host__ static type operator op (type& val, int i) { \
+__device__ __host__ static inline type operator op (type& val, int i) { \
   type ret; \
   ret.x = val.x; \
   ret.y = val.y; \
@@ -1418,7 +1418,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
 }
 
 #define DECLOP_3VAR_1IN_1OUT(type, op) \
-__device__ __host__ static type operator op(type &rhs) { \
+__device__ __host__ static inline type operator op(type &rhs) { \
   type ret; \
   ret.x = op rhs.x; \
   ret.y = op rhs.y; \
@@ -1437,7 +1437,7 @@ __device__ __host__ static inline bool operator op (type &rhs) { \
 */
 
 #define DECLOP_4VAR_2IN_1OUT(type, op) \
-__device__ __host__ static type operator op ( const type& lhs, const type& rhs) { \
+__device__ __host__ static inline type operator op ( const type& lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs.x op rhs.x; \
   ret.y = lhs.y op rhs.y; \
@@ -1447,7 +1447,7 @@ __device__ __host__ static type operator op ( const type& lhs, const type& rhs)
 }
 
 #define DECLOP_4VAR_SCALE_PRODUCT(type, type1) \
-__device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
+__device__ __host__ static inline type operator * (const type& lhs, type1 rhs) { \
   type ret; \
   ret.x = lhs.x * rhs; \
   ret.y = lhs.y * rhs; \
@@ -1456,7 +1456,7 @@ __device__ __host__ static type operator * (const type& lhs, type1 rhs) { \
   return ret; \
 } \
 \
-__device__ __host__ static type operator * (type1 lhs, const type& rhs) { \
+__device__ __host__ static inline type operator * (type1 lhs, const type& rhs) { \
   type ret; \
   ret.x = lhs * rhs.x; \
   ret.y = lhs * rhs.y; \
@@ -1484,7 +1484,7 @@ __device__ __host__ static inline type& operator op (type& val) { \
 }
 
 #define DECLOP_4VAR_POSTOP(type, op) \
-__device__ __host__ static type operator op (type& val, int i) { \
+__device__ __host__ static inline type operator op (type& val, int i) { \
   type ret; \
   ret.x = val.x; \
   ret.y = val.y; \
@@ -1503,7 +1503,7 @@ __device__ __host__ static inline bool operator op (type& lhs, type& rhs) { \
 }
 
 #define DECLOP_4VAR_1IN_1OUT(type, op) \
-__device__ __host__ static type operator op(type &rhs) { \
+__device__ __host__ static inline type operator op(type &rhs) { \
   type ret; \
   ret.x = op rhs.x; \
   ret.y = op rhs.y; \

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

@@ -36,7 +36,7 @@ __device__ float ceilf(float x);
 __device__ float copysignf(float x, float y);
 __device__ float cosf(float x);
 __device__ float coshf(float x);
-__device__ __host__ float cospif(float x);
+__device__ float cospif(float x);
 //__device__ float cyl_bessel_i0f(float x);
 //__device__ float cyl_bessel_i1f(float x);
 __device__ float erfcf(float x);
@@ -50,7 +50,7 @@ __device__ float expf(float x);
 __device__ float expm1f(float x);
 __device__ float fabsf(float x);
 __device__ float fdimf(float x, float y);
-__device__ __host__ float fdividef(float x, float y);
+__device__ float fdividef(float x, float y);
 __device__ float floorf(float x);
 __device__ float fmaf(float x, float y, float z);
 __device__ float fmaxf(float x, float y);
@@ -59,7 +59,7 @@ __device__ float fmodf(float x, float y);
 //__device__ float frexpf(float x, int* nptr);
 __device__ float hypotf(float x, float y);
 __device__ float ilogbf(float x);
-__device__ __host__ int isfinite(float a);
+__device__ int isfinite(float a);
 __device__ unsigned isinf(float a);
 __device__ unsigned isnan(float a);
 __device__ float j0f(float x);
@@ -93,15 +93,15 @@ __device__ float rnorm3df(float a, float b, float c);
 __device__ float rnorm4df(float a, float b, float c, float d);
 __device__ float rnormf(int dim, const float* a);
 __device__ float roundf(float x);
-__device__ __host__ float rsqrtf(float x);
+__device__ float rsqrtf(float x);
 __device__ float scalblnf(float x, long int n);
 __device__ float scalbnf(float x, int n);
-__device__ __host__ unsigned signbit(float a);
+__device__ int signbit(float a);
 __device__ void sincosf(float x, float *sptr, float *cptr);
 __device__ void sincospif(float x, float *sptr, float *cptr);
 __device__ float sinf(float x);
 __device__ float sinhf(float x);
-__device__ __host__ float sinpif(float x);
+__device__ float sinpif(float x);
 __device__ float sqrtf(float x);
 __device__ float tanf(float x);
 __device__ float tanhf(float x);
@@ -111,11 +111,12 @@ __device__ float y0f(float x);
 __device__ float y1f(float x);
 __device__ float ynf(int n, float x);
 
-
+__host__ float cospif(float x);
 __host__ float normcdff(float y);
 __host__ float erfcinvf(float y);
 __host__ float erfcxf(float x);
 __host__ float erfinvf(float y);
+__host__ float fdividef(float x, 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);
@@ -125,7 +126,10 @@ __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);
-
+__host__ int isfinite(float a);
+__host__ float rsqrtf(float x);
+__host__ float sinpif(float x);
+__host__ int signbit(float a);
 
 
 __device__ double acos(double x);
@@ -140,7 +144,7 @@ __device__ double ceil(double x);
 __device__ double copysign(double x, double y);
 __device__ double cos(double x);
 __device__ double cosh(double x);
-__device__ __host__ double cospi(double x);
+__device__ double cospi(double x);
 //__device__ double cyl_bessel_i0(double x);
 //__device__ double cyl_bessel_i1(double x);
 __device__ double erf(double x);
@@ -162,7 +166,7 @@ __device__ double fmod(double x, double y);
 //__device__ double frexp(double x, int *nptr);
 __device__ double hypot(double x, double y);
 __device__ double ilogb(double x);
-__device__ __host__ unsigned isfinite(double x);
+__device__ int isfinite(double x);
 __device__ unsigned isinf(double x);
 __device__ unsigned isnan(double x);
 __device__ double j0(double x);
@@ -198,15 +202,15 @@ __device__ double rnorm(int dim, const double* t);
 __device__ double rnorm3d(double a, double b, double c);
 __device__ double rnorm4d(double a, double b, double c, double d);
 __device__ double round(double x);
-__device__ __host__ double rsqrt(double x);
+__device__ double rsqrt(double x);
 __device__ double scalbln(double x, long int n);
 __device__ double scalbn(double x, int n);
-__device__ __host__ unsigned signbit(double a);
+__device__ int signbit(double a);
 __device__ double sin(double a);
 __device__ void sincos(double x, double *sptr, double *cptr);
 __device__ void sincospi(double x, double *sptr, double *cptr);
 __device__ double sinh(double x);
-__device__ __host__ double sinpi(double x);
+__device__ double sinpi(double x);
 __device__ double sqrt(double x);
 __device__ double tan(double x);
 __device__ double tanh(double x);
@@ -231,6 +235,11 @@ __host__ double rnorm(int dim, const double* t);
 __host__ double rnorm3d(double a, double b, double c);
 __host__ double rnorm4d(double a, double b, double c, double d);
 __host__ void sincospi(double x, double *sptr, double *cptr);
+__host__ double cospi(double x);
+__host__ int isfinite(double x);
+__host__ double rsqrt(double x);
+__host__ int signbit(double a);
+__host__ double sinpi(double x);
 
 // ENDPARSER
 

projects/clr/hipamd/src/math_functions.cpp

@@ -158,7 +158,7 @@ __device__ float ilogbf(float x)
 {
     return hc::precise_math::ilogbf(x);
 }
-__device__ unsigned isfinite(float a)
+__device__ int isfinite(float a)
 {
     return hc::precise_math::isfinite(a);
 }
@@ -334,7 +334,7 @@ __device__ float scalbnf(float x, int n)
 {
     return hc::precise_math::scalbnf(x, n);
 }
-__device__ unsigned signbit(float a)
+__device__ int signbit(float a)
 {
     return hc::precise_math::signbit(a);
 }
@@ -539,7 +539,7 @@ __device__ double ilogb(double x)
 {
     return hc::precise_math::ilogb(x);
 }
-__device__ unsigned isfinite(double x)
+__device__ int isfinite(double x)
 {
     return hc::precise_math::isfinite(x);
 }
@@ -712,7 +712,7 @@ __device__ double scalbn(double x, int n)
 {
     return hc::precise_math::scalbn(x, n);
 }
-__device__ unsigned signbit(double x)
+__device__ int signbit(double x)
 {
     return hc::precise_math::signbit(x);
 }
@@ -792,7 +792,7 @@ __host__ int signbit(float x)
   return std::signbit(x);
 }
 
-__host__ int sinpif(float x)
+__host__ float sinpif(float x)
 {
   return std::sin(x*HIP_PI);
 }
@@ -984,7 +984,7 @@ __host__ double cospi(double a)
   return std::cos(HIP_PI * a);
 }
 
-__host__ double isfinite(double a)
+__host__ int isfinite(double a)
 {
   return std::isfinite(a);
 }