     }
     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) {

__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) \

    return *this;
  }

  friend __device__ float16 copysign(float16 x, float16 y) {
    float16 ret;
    ret.data_ = (x.data_ & 0x7fffu) | (y.data_ & 0x8000u);
    return ret;
  }

  friend __device__ float16 nextafter(float16 x, float16 y) {
    float16 ret;
    if (!x.isfinite() || y.isnan()) {
      ret.data_ = nan;
