diff options
| author |  Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-01-11 11:43:37 -0800 |
|---|---|---|
| committer | Benoit Steiner <benoit.steiner.goog@gmail.com> | 2016-01-11 11:43:37 -0800 |
| commit | 2c3b13eded68aa0cef0185a7db9483bfc27d749b (patch) | |
| tree | 5dcd63ebd6736c037dfdab7043f4be47c2386917 /unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h | |
| parent | 2ccb1c86342203cbc25a587590149d2cf5175900 (diff) | |
| parent | 91678f489a89b1f4a393cd3b703e67922b5c4257 (diff) | |
Merged in jeremy_barnes/eigen/shader-model-3.0 (pull request PR-152)
Alternative way of forcing instantiation of device kernels without causing warnings or requiring device to device kernel invocations.
Diffstat (limited to 'unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h')
| -rw-r--r-- | unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h | 8 |
1 files changed, 7 insertions, 1 deletions
diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h index af140a68b..c74613873 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h @@ -238,9 +238,15 @@ struct GpuDevice { }; +#ifndef __CUDA_ARCH__ #define LAUNCH_CUDA_KERNEL(kernel, gridsize, blocksize, sharedmem, device, ...) \ - (kernel) <<< (gridsize), (blocksize), (sharedmem), (device).stream() >>> (__VA_ARGS__); \ + (kernel) <<< (gridsize), (blocksize), (sharedmem), (device).stream() >>> (__VA_ARGS__); \ assert(cudaGetLastError() == cudaSuccess); +#else +#define LAUNCH_CUDA_KERNEL(kernel, ...) \ + { static const auto __attribute__((__unused__)) __makeTheKernelInstantiate = &(kernel); } \ + eigen_assert(false && "Cannot launch a kernel from another kernel" __CUDA_ARCH__); +#endif // FIXME: Should be device and kernel specific. |