Merge remote-tracking branch 'origin/master' into mahilleb/cuDNN5

CNTK.Cpp.props

@@ -24,6 +24,7 @@
     <NvmlLibPath>"c:\Program Files\NVIDIA Corporation\GDK\gdk_win7_amd64_release\nvml\lib"</NvmlLibPath>
 
     <CudaVersion/>
+    <CudaVersion Condition="Exists('$(CUDA_PATH_V8_0)') And '$(CudaVersion)' == ''">8.0</CudaVersion>
     <CudaVersion Condition="Exists('$(CUDA_PATH_V7_5)') And '$(CudaVersion)' == ''">7.5</CudaVersion>
     <CudaVersion Condition="Exists('$(CUDA_PATH_V7_0)') And '$(CudaVersion)' == ''">7.0</CudaVersion>
 
@@ -87,6 +88,12 @@
     <OpenCvBinPath>$(OpenCvPath)\x64\vc12\bin</OpenCvBinPath>
   </PropertyGroup>
 
+  <PropertyGroup Condition="'$(CudaVersion)' == '8.0'">
+    <CudaPath>$(CUDA_PATH_V8_0)</CudaPath>
+    <CudaRuntimeDll>cudart64_80.dll</CudaRuntimeDll>
+    <CudaDlls>cublas64_80.dll;cusparse64_80.dll;curand64_80.dll;$(CudaRuntimeDll)</CudaDlls>
+  </PropertyGroup>
+
   <PropertyGroup Condition="'$(CudaVersion)' == '7.5'">
     <CudaPath>$(CUDA_PATH_V7_5)</CudaPath>
     <CudaRuntimeDll>cudart64_75.dll</CudaRuntimeDll>

Source/Math/GPUMatrixCUDAKernels.cuh

@@ -42,12 +42,8 @@
 
 #define IDX2C(i, j, ld) (((j) * (ld)) + (i)) // 0 based indexing
 
-// TODO: This condition seems wrong, it should be:
-// !defined(__CUDA_ARCH__) || __CUDA_ARCH__ < 600
-// NVIDIA should fix their CUDA 8.0 headers
+// On older GPUs, CUDA atomicAdd() only exists for 'float'. This is the 'double' version.
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 600
-// CUDA atomicAdd() only exists for 'float'. This is the 'double' version.
-// TODO: This may need to be guarded by CUDA version; newer devices may support this.
 static __inline__ __device__ double atomicAdd(double* address, double val)
 {
     unsigned long long int* address_as_ull = (unsigned long long int*) address;