 #define M_PI 3.1415926535897932384626433832795
 #endif
 
-#if __CUDACC_VER_MAJOR__ < 9
-
-// float16
-class float16
-{
+#if __CUDACC_VER_MAJOR__ >= 9
+
+#include <cuda_fp16.h>
+
+#else  // #if __CUDACC_VER_MAJOR__ >= 9
+
+struct __half_raw {
+  unsigned short x;
+};
+
+struct half {
 private:
   unsigned short data_;
 public:
+  __device__ half() {}
+  __device__ half(const half &v) : data_(v.data_) {}
+  __device__ half(float v) : data_(__float2half_rn(v)) {}
+
+  explicit __device__ half(const __half_raw &v) : data_(v.x) {}
+  explicit __device__ half(bool v) : data_(__float2half_rn(float(v))) {}
+  explicit __device__ half(double v) : data_(__float2half_rn(float(v))) {}
+  explicit __device__ half(int v) : data_(__float2half_rn(float(v))) {}
+  explicit __device__ half(unsigned int v) : data_(__float2half_rn(float(v))) {}
+  explicit __device__ half(long long v) : data_(__float2half_rn(float(v))) {}
+  explicit __device__ half(unsigned long long v) : data_(__float2half_rn(float(v))) {}
+
+  __device__ operator float() const {return __half2float(data_);}
+  __device__ operator __half_raw() const {__half_raw ret = {data_}; return ret;}
+};
+
+#endif  // #if __CUDACC_VER_MAJOR__ >= 9
+
+class float16 {
+private:
+  half  data_;
+public:
   __device__ float16() {}
-  __device__ float16(float v) {
-    data_ = __float2half_rn(v);
-  }
-  explicit __device__ float16(bool v) {
-    data_ = __float2half_rn(static_cast<float>(v));
-  }
-  explicit __device__ float16(double v) {
-    data_ = __float2half_rn(static_cast<float>(v));
-  }
-  explicit __device__ float16(int v) {
-    data_ = __float2half_rn(static_cast<float>(v));
-  }
-  explicit __device__ float16(unsigned int v) {
-    data_ = __float2half_rn(static_cast<float>(v));
-  }
-  explicit __device__ float16(long long v) {
-    data_ = __float2half_rn(static_cast<float>(v));
-  }
-  explicit __device__ float16(unsigned long long v) {
-    data_ = __float2half_rn(static_cast<float>(v));
-  }
-
-  __device__ operator float() const {return __half2float(data_);}
+  __device__ float16(float v) : data_(v) {}
+
+  explicit __device__ float16(bool v) : data_(float(v)) {}
+  explicit __device__ float16(double v) : data_(v) {}
+  explicit __device__ float16(int v) : data_(v) {}
+  explicit __device__ float16(unsigned int v) : data_(v) {}
+  explicit __device__ float16(long long v) : data_(v) {}
+  explicit __device__ float16(unsigned long long v) : data_(v) {}
+
+  explicit __device__ float16(const half &v): data_(v) {}
+  explicit __device__ float16(const __half_raw &v): data_(v) {}
+
+  __device__ operator float() const {return float(data_);}
 
   static const unsigned short nan = 0x7e00u;
 
   __device__ int iszero() const {
-    return (data_ & 0x7fff) == 0;
-  }
-
-  __device__ int isnan() const {
-    return (data_ & 0x7c00u) == 0x7c00u && (data_ & 0x03ffu) != 0x0000u;
-  }
-
-  __device__ int isinf() const {
-    return (data_ & 0x7fffu) == 0x7c00u;
-  }
-
-  __device__ int isfinite() const {
-    return (data_ & 0x7c00u) != 0x7c00u;
-  }
-
-  __device__ int signbit() const
-  {
-    return (data_ & 0x8000u) != 0;
-  }
-
-  template<typename T>
-  inline __device__ float16& operator+=(const T& rhs)
-  {
-    *this = *this + rhs;
-    return *this;
-  }
-
-  template<typename T>
-  inline __device__ float16& operator-=(const T& rhs)
-  {
-    *this = *this - rhs;
-    return *this;
-  }
-
-  template<typename T>
-  inline __device__ float16& operator*=(const T& rhs)
-  {
-    *this = *this * rhs;
-    return *this;
-  }
-
-  template<typename T>
-  inline __device__ float16& operator/=(const T& rhs)
-  {
-    *this = *this / rhs;
-    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;
-    } else if (eq_nonan(x, y)) {
-      ret = x;
-    } else if (x.iszero()) {
-      ret.data_ = (y.data_ & 0x8000u) + 1;
-    } else if (!(x.data_ & 0x8000u)) {
-      if (static_cast<short>(x.data_) > static_cast<short>(y.data_)) {
-        ret.data_ = x.data_ - 1;
-      } else {
-        ret.data_ = x.data_ + 1;
-      }
-    } else if(!(y.data_ & 0x8000u) || (x.data_ & 0x7fffu) > (y.data_ & 0x7fffu)) {
-      ret.data_ = x.data_ - 1;
-    } else {
-      ret.data_ = x.data_ + 1;
-    }
-    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);
-  }
-};
-
-#else  // #if __CUDACC_VER_MAJOR__ < 9
-
-#include <cuda_fp16.h>
-
-// float16
-class float16
-{
-private:
-    half  data_;
-public:
-  __device__ float16() {}
-  __device__ float16(float v) {
-    data_ = __half(v);
-  }
-
-  explicit __device__ float16(bool v) {
-    data_ = __float2half_rn(static_cast<short>(v));
-  }
-  explicit __device__ float16(double v) {
-    data_ = __half(v);
-  }
-  explicit __device__ float16(int v) {
-    data_ = __half(v);
-  }
-  explicit __device__ float16(unsigned int v) {
-    data_ = __half(v);
-  }
-  explicit __device__ float16(long long v) {
-    data_ = __half(v);
-  }
-  explicit __device__ float16(unsigned long long v) {
-    data_ = __half(v);
-  }
-  explicit __device__ float16(half v) {
-    data_ = v;
-  }
-
-  __device__ operator float() const {return __half2float(data_);}
-
-  static const unsigned short nan = 0x7e00u;
-
-  __device__ int iszero() const {
-    __half_raw raw_ = __half_raw(data_);
-    return (raw_.x & 0x7fff) == 0;
+    return (__half_raw(data_).x & 0x7fffu) == 0;
   }
 
   __device__ int isnan() const {

      ret_raw_.x = (y_raw_.x & 0x8000u) + 1;
    } else if (!(x_raw_.x & 0x8000u)) {
      if (x_raw_.x > y_raw_.x) {
	ret_raw_.x = x_raw_.x - 1;
      } else {
	ret_raw_.x = x_raw_.x + 1;
      }
    } else if(!(y_raw_.x & 0x8000u) || (x_raw_.x & 0x7fffu) > (y_raw_.x & 0x7fffu)) {
      ret_raw_.x = x_raw_.x - 1;
    } else {
      ret_raw_.x = x_raw_.x + 1;
    }
    ret.data_ = ret_raw_;
    return ret;
  }

private:

  }

  __device__ int isinf() const {
    return (__half_raw(data_).x & 0x7fffu) == 0x7c00u;
  }

  __device__ int isfinite() const {
    return (__half_raw(data_).x & 0x7c00u) != 0x7c00u;
  }

  __device__ int signbit() const {
    return (__half_raw(data_).x & 0x8000u) != 0;
  }

  template<typename T>
  inline __device__ float16& operator+=(const T& rhs) {
    *this = *this + rhs;
    return *this;
  }

  template<typename T>
  inline __device__ float16& operator-=(const T& rhs) {
    *this = *this - rhs;
    return *this;
  }

  template<typename T>
  inline __device__ float16& operator*=(const T& rhs) {
    *this = *this * rhs;
    return *this;
  }

  template<typename T>
  inline __device__ float16& operator/=(const T& rhs) {
    *this = *this / rhs;
    return *this;
  }

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

  friend __device__ float16 nextafter(float16 x, float16 y) {
    __half_raw x_raw_ = __half_raw(x.data_);
    __half_raw y_raw_ = __half_raw(y.data_);
    __half_raw ret_raw_;
