From b9ea40c30d1d32d0f31b047aa681c384fd1a2c98 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Tue, 19 Apr 2016 14:35:11 -0700 Subject: Don't take the address of a kernel on CUDA devices that don't support this feature. --- unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h') diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h index 821835cf3..8e7f5dddb 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h @@ -291,14 +291,17 @@ struct GpuDevice { int max_blocks_; }; -#ifndef __CUDA_ARCH__ +#if !defined(__CUDA_ARCH__) #define LAUNCH_CUDA_KERNEL(kernel, gridsize, blocksize, sharedmem, device, ...) \ (kernel) <<< (gridsize), (blocksize), (sharedmem), (device).stream() >>> (__VA_ARGS__); \ assert(cudaGetLastError() == cudaSuccess); -#else +#elif __CUDA_ARCH__ >= 350 #define LAUNCH_CUDA_KERNEL(kernel, ...) \ { const auto __attribute__((__unused__)) __makeTheKernelInstantiate = &(kernel); } \ - eigen_assert(false && "Cannot launch a kernel from another kernel" __CUDA_ARCH__); + eigen_assert(false && "Cannot launch a kernel from another kernel" __CUDA_ARCH__ kernel); +#else +#define LAUNCH_CUDA_KERNEL(kernel, ...) \ + eigen_assert(false && "Cannot launch a kernel from another kernel" __CUDA_ARCH__ kernel); #endif -- cgit v1.2.3