diff options
author | Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-04-19 14:35:11 -0700 |
---|---|---|
committer | Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-04-19 14:35:11 -0700 |
commit | b9ea40c30d1d32d0f31b047aa681c384fd1a2c98 (patch) | |
tree | 75f4f6483a58103c19a940579cfc6ec8e0e1492b | |
parent | 884c07505868c0167467c5c3de207724b24f12ab (diff) |
Don't take the address of a kernel on CUDA devices that don't support this feature.
-rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h | 9 |
1 files changed, 6 insertions, 3 deletions
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 |