   #endif
 }
 
+// arithmetic functions
+
+inline __host__ __device__ half operator+(half a, half b) {
+  return THCNumerics<half>::add(a, b);
+}
+
+inline __host__ __device__ float operator+(float a, half b) {
+  return a + ScalarConvert<half, float>::to(b);
+}
+
+inline __host__ __device__ double operator+(double a, half b) {
+  return a + ScalarConvert<half, double>::to(b);
+}
+
+inline __host__ __device__ half operator-(half a) {
+  return THCNumerics<half>::neg(a);
+}
+
+inline __host__ __device__ half operator-(half a, half b) {
+  return THCNumerics<half>::add(a, THCNumerics<half>::neg(b));
+}
+
+inline __host__ __device__ half operator-(half a, int b) {
+  return THCNumerics<half>::add(a, THCNumerics<half>::neg(ScalarConvert<int, half>::to(b)));
+}
+
+inline __host__ __device__ float operator-(half a, float b) {
+  return ScalarConvert<half, float>::to(a) - b;
+}
+
+inline __host__ __device__ double operator-(half a, double b) {
+  return ScalarConvert<half, double>::to(a) - b;
+}
+
+inline __host__ __device__ half operator-(int a, half b) {
+  return THCNumerics<half>::add(ScalarConvert<int, half>::to(a), THCNumerics<half>::neg(b));
+}
+
+inline __host__ __device__ float operator-(float a, half b) {
+  return a - ScalarConvert<half, float>::to(b);
+}
+
+inline __host__ __device__ double operator-(double a, half b) {
+  return a - ScalarConvert<half, double>::to(b);
+}
+
+inline __host__ __device__ half operator*(half a, half b) {
+  return mul(a, b);
+}
+
+inline __host__ __device__ float operator*(half a, float b) {
+  return ScalarConvert<half, float>::to(a) * b;
+}
+
+inline __host__ __device__ double operator*(half a, double b) {
+  return ScalarConvert<half, double>::to(a) * b;
+}
+
+inline __host__ __device__ half operator*(half a, int b) {
+  return a * ScalarConvert<int, half>::to(b);
+}
+
+inline __host__ __device__ float operator*(float a, half b) {
+  return a * ScalarConvert<half, float>::to(b);
+}
+
+inline __host__ __device__ double operator*(double a, half b) {
+  return a * ScalarConvert<half, double>::to(b);
+}
+
+inline __host__ __device__ half operator/(half a, half b) {
+  return div(a, b);
+}
+
+inline __host__ __device__ float operator/(float a, half b) {
+  return a / ScalarConvert<half, float>::to(b);
+}
+
+inline __host__ __device__ double operator/(double a, half b) {
+  return a / ScalarConvert<half, double>::to(b);
+}
+
 inline __host__ __device__ half operator/(int a, half b) {
   return ScalarConvert<int, half>::to(a) / b;
 }
 
-inline __host__ __device__ float operator/(float a, half b) {
-  return a / ScalarConvert<half, float>::to(b);
-}
-
-inline __host__ __device__ double operator/(double a, half b) {
-  return a / ScalarConvert<half, double>::to(b);
-}
-
 inline __host__ __device__ float operator/(half a, float b) {
   return ScalarConvert<half, float>::to(a) / b;
 }
 
 inline __host__ __device__ double operator/(half a, double b) {
   return ScalarConvert<half, double>::to(a) / b;
-}
-
-inline __device__ half fastExpIfAvail(half a) {
-  return ScalarConvert<float, half>::to(__expf(ScalarConvert<half, float>::to(a)));
-}
-
-inline __device__ float fastExpIfAvail(float a) {
-  return __expf(a);
-}
-
-inline __device__ double fastExpIfAvail(double a) {
-  // should we convert down to float to be able to use __expf?
-  return THCNumerics<double>::exp(a);
-}
-
-// these should move to THCNumerics
-inline __host__ __device__ float fmaxType(float x, half y) {
-  return fmaxf(x, ScalarConvert<half, float>::to(y));
-}
-
-inline __host__ __device__ float fmaxType(float x, float y) {
-  return fmaxf(x, y);
-}
-
-inline __host__ __device__ double fmaxType(double x, double y) {
-  return fmax(x, y);
 }
 
 inline __host__ __device__ half& operator+=(half &lhs, const half &rhs) {

  #endif
}

inline __host__ __device__ float operator*(half a, float b) {
  return ScalarConvert<half, float>::to(a) * b;
}

inline __host__ __device__ double operator*(half a, double b) {
  return ScalarConvert<half, double>::to(a) * b;
}

inline __host__ __device__ half operator*(half a, int b) {
  return a * ScalarConvert<int, half>::to(b);
}

inline __host__ __device__ float operator*(float a, half b) {
  return a * ScalarConvert<half, float>::to(b);

inline __host__ __device__ double operator*(double a, half b) {
  return a * ScalarConvert<half, double>::to(b);
}

inline __host__ __device__ bool operator<(half a, half b) {
  return THCNumerics<half>::lt(a, b);
}

inline __host__ __device__ bool operator<=(half a, half b) {
  return THCNumerics<half>::le(a, b);
}

inline __host__ __device__ bool operator<=(half a, int b) {
  return THCNumerics<half>::le(a, ScalarConvert<int, half>::to(b));
}

inline __host__ __device__ bool operator<(half a, int b) {
  return THCNumerics<half>::lt(a, ScalarConvert<int, half>::to(b));
}

inline __host__ __device__ bool operator>(half a, half b) {
  return THCNumerics<half>::gt(a, b);
}

inline __host__ __device__ bool operator>(half a, int b) {
  return THCNumerics<half>::gt(a, ScalarConvert<int, half>::to(b));
}

inline __host__ __device__ bool operator>=(half a, half b) {
  return THCNumerics<half>::ge(a, b);
}

inline __host__ __device__ float& operator+=(float &lhs, const half &rhs) {
  lhs = lhs + rhs;
  return lhs;
}

inline __host__ __device__ half abs(half a) {
  return THCNumerics<half>::abs(a);
}

inline __host__ __device__ half exp(half a) {
  return THCNumerics<half>::exp(a);
}

inline __host__ __device__ half log1p(half a) {
  return THCNumerics<half>::log1p(a);
}

inline __host__ __device__ half sqrt(half a) {
  return THCNumerics<half>::sqrt(a);
}

inline __host__ __device__ half tanh(half a) {
  return THCNumerics<half>::tanh(a);
}

// This implementation could move to THCNumerics
inline __host__ __device__ half operator/(half a, half b) {
  #ifdef __CUDA_ARCH__
  #ifdef CUDA_HALF_INSTRUCTIONS
    return __hdiv(a, b);

  return lhs;
}

inline __host__ __device__ float& operator+=(float &lhs, const half &rhs) {
  lhs = lhs + rhs;
  return lhs;
}

inline __host__ __device__ half abs(half a) {
  return THCNumerics<half>::abs(a);
}

inline __host__ __device__ half exp(half a) {
  return THCNumerics<half>::exp(a);
}

inline __host__ __device__ half log1p(half a) {
  return THCNumerics<half>::log1p(a);
}

inline __host__ __device__ half sqrt(half a) {
  return THCNumerics<half>::sqrt(a);
}

inline __host__ __device__ half tanh(half a) {
  return THCNumerics<half>::tanh(a);
}

// comparison functions

inline __host__ __device__ bool operator<(half a, half b) {
  return THCNumerics<half>::lt(a, b);
}

inline __host__ __device__ bool operator<=(half a, half b) {
  return THCNumerics<half>::le(a, b);
}

inline __host__ __device__ bool operator<=(half a, int b) {
  return THCNumerics<half>::le(a, ScalarConvert<int, half>::to(b));
}

inline __host__ __device__ bool operator<(half a, int b) {
  return THCNumerics<half>::lt(a, ScalarConvert<int, half>::to(b));
}

inline __host__ __device__ bool operator>(half a, half b) {
  return THCNumerics<half>::gt(a, b);
}

inline __host__ __device__ bool operator>(half a, int b) {
  return THCNumerics<half>::gt(a, ScalarConvert<int, half>::to(b));
}

inline __host__ __device__ bool operator>=(half a, half b) {
  return THCNumerics<half>::ge(a, b);
}

#endif
#endif
