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 /unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h | |
| parent | 884c07505868c0167467c5c3de207724b24f12ab (diff) | |
Don't take the address of a kernel on CUDA devices that don't support this feature.
Diffstat (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h')
| -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 |