Updates for CUDA 9

Author: csarofeen

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];

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)