Remove some unnecessary static functions in carray.cuh
@@ -85,17 +85,13 @@ public:
return *this;
}
- static __device__ float16 copysign(float16 x, float16 y) {
+ friend __device__ float16 copysign(float16 x, float16 y) {
float16 ret;
ret.data_ = (x.data_ & 0x7fffu) | (y.data_ & 0x8000u);
return ret;
}
- static __device__ int eq_nonan(const float16 x, const float16 y) {
- return (x.data_ == y.data_ || ((x.data_ | y.data_) & 0x7fff) == 0);
- }
-
- static __device__ float16 nextafter(float16 x, float16 y) {
+ friend __device__ float16 nextafter(float16 x, float16 y) {
float16 ret;
if (!x.isfinite() || y.isnan()) {
ret.data_ = nan;
@@ -116,6 +112,11 @@ public:
}
return ret;
}
+
+private:
+ static __device__ int eq_nonan(const float16 x, const float16 y) {
+ return (x.data_ == y.data_ || ((x.data_ | y.data_) & 0x7fff) == 0);
+ }
};
__device__ float16 min(float16 x, float16 y) {
@@ -129,9 +130,6 @@ __device__ int isnan(float16 x) {return x.isnan();}
__device__ int isinf(float16 x) {return x.isinf();}
__device__ int isfinite(float16 x) {return x.isfinite();}
__device__ int signbit(float16 x) {return x.signbit();}
-__device__ float16 copysign(float16 x, float16 y) {return float16::copysign(x, y);}
-__device__ float16 nextafter(float16 x, float16 y) {return float16::nextafter(x, y);}
-
// CArray
#define CUPY_FOR(i, n) \
