Updates for CUDA 9

Author: csarofeen

lib/THCUNN/BatchNormalization.cu

@@ -5,7 +5,7 @@
 
 #include "THCDeviceTensor.cuh"
 #include "THCDeviceTensorUtils.cuh"
-
+#include "THCDeviceUtils.cuh"
 const int WARP_SIZE = 32;
 
 // The maximum number of threads in a block
@@ -80,7 +80,7 @@ template <typename T>
 static __device__ __forceinline__ T warpSum(T val) {
 #if __CUDA_ARCH__ >= 300
   for (int i = 0; i < getMSB(WARP_SIZE); ++i) {
-    val += __shfl_xor(val, 1 << i, WARP_SIZE);
+    val += WARP_SHFL_XOR(val, 1 << i, WARP_SIZE);
   }
 #else
   __shared__ T values[MAX_BLOCK_SIZE];

lib/THCUNN/CMakeLists.txt

@@ -51,6 +51,10 @@ if(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
   endif(CMAKE_CXX_COMPILER_VERSION VERSION_GREATER "4.9.3")
 endif(CMAKE_CXX_COMPILER_ID STREQUAL "GNU")
 
+if(CUDA_VERSION VERSION_GREATER "8.0")
+  LIST(APPEND CUDA_NVCC_FLAGS "-D__CUDA_NO_HALF_OPERATORS__")
+endif(CUDA_VERSION VERSION_GREATER "8.0")
+
 IF(MSVC)
   LIST(APPEND CUDA_NVCC_FLAGS "-Xcompiler /wd4819")
   ADD_DEFINITIONS(-DTH_EXPORTS)