Fix modf/frexp/remquo in devicelib

These had incoherent function signatures, with pointers using
different addrspaces. apparently the compilation command chain
decided to deal with it by function casts but the intel gpu driver
doesn't seem to like them.

This should fix all signatures to use generic AS.

bitcode/devicelib.cl

@@ -46,8 +46,7 @@
 #pragma OPENCL EXTENSION cl_khr_int64_base_atomics : enable
 #pragma OPENCL EXTENSION cl_khr_int64_extended_atomics : enable
 
-#define DEFAULT_AS
-#define PRIVATE_AS __private
+#define DEFAULT_AS __generic
 
 #define CL_NAME(N) opencl_##N
 #define CL_NAME2(N, S) opencl_##N##_##S
@@ -176,8 +175,6 @@ DEF_OPENCL2F(fmax)
 DEF_OPENCL2F(fmin)
 DEF_OPENCL2F(fmod)
 
-float OVLD frexp(float f, PRIVATE_AS int *i);
-double OVLD frexp(double f, PRIVATE_AS int *i);
 EXPORT float CL_NAME2(frexp, f)(float x, DEFAULT_AS int *i) {
   int tmp;
   float ret = frexp(x, &tmp);
@@ -198,8 +195,6 @@ DEF_OPENCL1B(isfinite)
 DEF_OPENCL1B(isinf)
 DEF_OPENCL1B(isnan)
 
-float OVLD ldexp(float f, int k);
-double OVLD ldexp(double f, int k);
 EXPORT float CL_NAME2(ldexp, f)(float x, int k) { return ldexp(x, k); }
 EXPORT double CL_NAME2(ldexp, d)(double x, int k) { return ldexp(x, k); }
 
@@ -210,9 +205,6 @@ DEF_OPENCL1F(log2)
 DEF_OPENCL1F(logb)
 DEF_OPENCL1F(log)
 
-// modf
-float OVLD modf(float f, PRIVATE_AS float *i);
-double OVLD modf(double f, PRIVATE_AS double *i);
 EXPORT float CL_NAME2(modf, f)(float x, DEFAULT_AS float *i) {
   float tmp;
   float ret = modf(x, &tmp);
@@ -248,8 +240,6 @@ DEF_OPENCL2F(remainder)
 DEFOCML_OPENCL1F(rcbrt)
 
 // remquo
-float OVLD remquo(float x,   float y,  PRIVATE_AS int *quo);
-double OVLD remquo(double x, double y, PRIVATE_AS int *quo);
 EXPORT float CL_NAME2(remquo, f)(float x, float y, DEFAULT_AS int *quo) {
   int tmp;
   float rem = remquo(x, y, &tmp);
@@ -290,19 +280,16 @@ DEF_OPENCL1F(trunc)
 
 
 // sincos
-float OVLD sincos(float x, PRIVATE_AS float *cosval);
-double OVLD sincos(double x, PRIVATE_AS double *cosval);
-
 EXPORT float CL_NAME2(sincos, f)(float x, DEFAULT_AS float *cos) {
-  PRIVATE_AS float tmp;
-  PRIVATE_AS float sin = sincos(x, &tmp);
+  float tmp;
+  float sin = sincos(x, &tmp);
   *cos = tmp;
   return sin;
 }
 
 EXPORT double CL_NAME2(sincos, d)(double x, DEFAULT_AS double *cos) {
-  PRIVATE_AS double tmp;
-  PRIVATE_AS double sin = sincos(x, &tmp);
+  double tmp;
+  double sin = sincos(x, &tmp);
   *cos = tmp;
   return sin;
 }

include/hip/devicelib/double_precision/dp_math.hh

@@ -32,18 +32,18 @@ extern __device__ double rint(double x);
 extern __device__ double round(double x);
 extern __device__ long int convert_long(double x);
 
-__device__ long int lrint(double x) { return convert_long(rint(x)); }
-__device__ long int lround(double x) { return convert_long(round(x)); }
-
-__device__ long long int llrint(double x) { return lrint(x); }
-__device__ long long int llround(double x) { return lround(x); }
-
 extern __device__ double rnorm3d(double a, double b, double c);
 extern __device__ double rnorm4d(double a, double b, double c, double d);
 
 extern __device__ double lgamma ( double  x );
 }
 
+static inline __device__ long int lrint(double x) { return convert_long(rint(x)); }
+static inline __device__ long int lround(double x) { return convert_long(round(x)); }
+
+static inline __device__ long long int llrint(double x) { return lrint(x); }
+static inline __device__ long long int llround(double x) { return lround(x); }
+
 // __device__ double acos(double x)
 // __device__​ double acosh ( double  x )
 // __device__​ double asin ( double  x )

include/hip/devicelib/single_precision/sp_math.hh

@@ -32,23 +32,24 @@ extern __device__ float rint(float x);
 extern __device__ float round(float x);
 extern __device__ long int convert_long(float x);
 
-__device__ float rintf(float x) { return rint(x); }
-__device__ float roundf(float x) { return round(x); }
-
-__device__ long int lrintf(float x) { return convert_long(rint(x)); }
-__device__ long int lroundf(float x) { return convert_long(round(x)); }
-
-__device__ long long int llrintf(float x) { return lrintf(x); }
-__device__ long long int llroundf(float x) { return lroundf(x); }
-
 extern __device__ float rnorm3df(float a, float b, float c);
 extern __device__ float rnorm4df(float a, float b, float c, float d);
 
 extern __device__  float lgamma ( float  x );
-__device__  float lgammaf ( float  x ) { return (lgamma(x)); };
 
 }
 
+static inline __device__ float rintf(float x) { return rint(x); }
+static inline __device__ float roundf(float x) { return round(x); }
+
+static inline __device__ long int lrintf(float x) { return convert_long(rint(x)); }
+static inline __device__ long int lroundf(float x) { return convert_long(round(x)); }
+
+static inline __device__ long long int llrintf(float x) { return lrintf(x); }
+static inline __device__ long long int llroundf(float x) { return lroundf(x); }
+
+static inline __device__  float lgammaf ( float  x ) { return (lgamma(x)); };
+
 // __device__ float acosf(float x)
 // __device__  float acoshf ( float  x )
 // __device__  float asinf ( float  x )