From b2a47641ce0ad0642d93db0030cbf8cd0bb7f2c0 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Wed, 29 Jun 2016 15:32:47 -0700 Subject: Made the code compile when using CUDA architecture < 300 --- unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h | 8 ++++++++ 1 file changed, 8 insertions(+) (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h') diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h index d9bbcd858..5e512490c 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorReductionCuda.h @@ -117,6 +117,7 @@ template __global__ void FullReductionKernel(Reducer reducer, const Self input, Index num_coeffs, typename Self::CoeffReturnType* output, unsigned int* semaphore) { +#if __CUDA_ARCH__ >= 300 // Initialize the output value const Index first_index = blockIdx.x * BlockSize * NumPerThread + threadIdx.x; if (gridDim.x == 1) { @@ -171,6 +172,9 @@ __global__ void FullReductionKernel(Reducer reducer, const Self input, Index num // Let the last block reset the semaphore atomicInc(semaphore, gridDim.x + 1); } +#else + assert(0 && "Shouldn't be called on unsupported device"); +#endif } @@ -355,6 +359,7 @@ template __global__ void InnerReductionKernel(Reducer reducer, const Self input, Index num_coeffs_to_reduce, Index num_preserved_coeffs, typename Self::CoeffReturnType* output) { +#if __CUDA_ARCH__ >= 300 eigen_assert(blockDim.y == 1); eigen_assert(blockDim.z == 1); eigen_assert(gridDim.y == 1); @@ -414,6 +419,9 @@ __global__ void InnerReductionKernel(Reducer reducer, const Self input, Index nu } } } +#else + assert(0 && "Shouldn't be called on unsupported device"); +#endif } #ifdef EIGEN_HAS_CUDA_FP16 -- cgit v1.2.3