Merge commit '224f5eabf5cfb3a19abc1819f7dac230500b6bdb'

Author: soumith

torch/lib/THC/THCGeneral.c

@@ -916,3 +916,20 @@ void THCHeapUpdate(THCState *state, ptrdiff_t size) {
 
 #include "THCStorage.c"
 #include "THCAllocator.c"
+
+/* from THCHalf.h */
+
+half THC_float2half(float f)
+{
+ half h;
+ TH_float2halfbits(&f, &h.x);
+ return h;
+}
+
+float THC_half2float(half h)
+{
+ float f;
+ TH_halfbits2float(&h.x, &f);
+ return f;
+}
+

torch/lib/THC/THCHalf.cu

@@ -33,96 +33,6 @@ void THCHalf2Float(THCState *state, float *out, half *in, ptrdiff_t len) {
  in, in + len, out, __half2floatOp());
 }
 
-// FixMe: could call TH_half2float
-// and convert types here, but maybe slower?
-float THC_half2float(half h)
-{
- unsigned sign = ((h.x >> 15) & 1);
- unsigned exponent = ((h.x >> 10) & 0x1f);
- unsigned mantissa = ((h.x & 0x3ff) << 13);
-
- if (exponent == 0x1f) { /* NaN or Inf */
- mantissa = (mantissa ? (sign = 0, 0x7fffff) : 0);
- exponent = 0xff;
- } else if (!exponent) { /* Denorm or Zero */
- if (mantissa) {
- unsigned int msb;
- exponent = 0x71;
- do {
- msb = (mantissa & 0x400000);
- mantissa <<= 1; /* normalize */
- --exponent;
- } while (!msb);
- mantissa &= 0x7fffff; /* 1.mantissa is implicit */
- }
- } else {
- exponent += 0x70;
- }
-
- int temp = ((sign << 31) | (exponent << 23) | mantissa);
-
- float x;
- memcpy(&x,&temp,sizeof(float));
- return x;
-}
-
-half THC_float2half(float f)
-{
- half ret;
-
- unsigned x;
- memcpy(&x,&f,sizeof(f));
- unsigned u = (x & 0x7fffffff), remainder, shift, lsb, lsb_s1, lsb_m1;
- unsigned sign, exponent, mantissa;
-
- // Get rid of +NaN/-NaN case first.
- if (u > 0x7f800000) {
- ret.x = 0x7fffU;
- return ret;
- }
-
- sign = ((x >> 16) & 0x8000);
-
- // Get rid of +Inf/-Inf, +0/-0.
- if (u > 0x477fefff) {
- ret.x = sign | 0x7c00U;
- return ret;
- }
- if (u < 0x33000001) {
- ret.x = (sign | 0x0000);
- return ret;
- }
-
- exponent = ((u >> 23) & 0xff);
- mantissa = (u & 0x7fffff);
-
- if (exponent > 0x70) {
- shift = 13;
- exponent -= 0x70;
- } else {
- shift = 0x7e - exponent;
- exponent = 0;
- mantissa |= 0x800000;
- }
- lsb = (1 << shift);
- lsb_s1 = (lsb >> 1);
- lsb_m1 = (lsb - 1);
-
- // Round to nearest even.
- remainder = (mantissa & lsb_m1);
- mantissa >>= shift;
- if (remainder > lsb_s1 || (remainder == lsb_s1 && (mantissa & 0x1))) {
- ++mantissa;
- if (!(mantissa & 0x3ff)) {
- ++exponent;
- mantissa = 0;
- }
- }
-
- ret.x = (sign | (exponent << 10) | mantissa);
- return ret;
-}
-
 THC_EXTERNC int THC_nativeHalfInstructions(THCState *state) {
  cudaDeviceProp* prop =
  THCState_getCurrentDeviceProperties(state);

torch/lib/THC/THCTensorTopK.cuh

@@ -112,21 +112,33 @@ struct TopKTypeConfig<double> {
  }
 };
 
+#ifdef CUDA_HALF_TENSOR
 template <>
 struct TopKTypeConfig<half> {
  typedef unsigned int RadixType;
 
  static inline __device__ RadixType convert(half v) {
+#if defined(__CUDACC_VER__) && __CUDACC_VER__ >= 80000
  RadixType x = __half_as_ushort(v);
  RadixType mask = -((x >> 15)) | 0x8000;
  return (x ^ mask);
+#else
+ assert(false);
+ return 0u;
+#endif
  }
 
  static inline __device__ half deconvert(RadixType v) {
+#if defined(__CUDACC_VER__) && __CUDACC_VER__ >= 80000
  RadixType mask = ((v >> 15) - 1) | 0x8000;
  return __ushort_as_half(v ^ mask);
+#else
+ assert(false);
+ return ScalarConvert<int, half>::to(0);
+#endif
  }
 };
+#endif // CUDA_HALF_TENSOR
 
 // This function counts the distribution of all input values in a
 // slice we are selecting by radix digit at `radixDigitPos`, but only