GitHub-Proxy/eigen
1
0
Fork 0
mirror of https://gitlab.com/libeigen/eigen.git synced 2025-05-01 16:24:28 +08:00
Code Issues Packages Projects Releases Wiki Activity

Don't take the address of a kernel on CUDA devices that don't support this feature.

Browse Source
This commit is contained in:
Benoit Steiner 2016-04-19 14:35:11 -07:00
parent 884c075058
commit b9ea40c30d
1 changed files with 6 additions and 3 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

9
unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h
View File

@ -291,14 +291,17 @@ struct GpuDevice {
int max_blocks_; int max_blocks_;
}; };
#ifndef __CUDA_ARCH__ #if !defined(__CUDA_ARCH__)
#define LAUNCH_CUDA_KERNEL(kernel, gridsize, blocksize, sharedmem, device, ...) \ #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); assert(cudaGetLastError() == cudaSuccess);
#else #elif __CUDA_ARCH__ >= 350
#define LAUNCH_CUDA_KERNEL(kernel, ...) \ #define LAUNCH_CUDA_KERNEL(kernel, ...) \
{ const auto __attribute__((__unused__)) __makeTheKernelInstantiate = &(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 #endif