From 53749ff4152191d2f1bd56090a14f6474fe059c2 Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Fri, 8 Jan 2016 13:53:40 -0800 Subject: Prevent nvcc from miscompiling the cuda metakernel. Unfortunately this reintroduces some compulation warnings but it's much better than having to deal with random assertion failures. --- unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h | 6 +----- 1 file changed, 1 insertion(+), 5 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 4d7570077..af140a68b 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorDeviceCuda.h @@ -238,14 +238,10 @@ struct GpuDevice { }; -#ifndef __CUDA_ARCH__ #define LAUNCH_CUDA_KERNEL(kernel, gridsize, blocksize, sharedmem, device, ...) \ (kernel) <<< (gridsize), (blocksize), (sharedmem), (device).stream() >>> (__VA_ARGS__); \ assert(cudaGetLastError() == cudaSuccess); -#else -#define LAUNCH_CUDA_KERNEL(...) \ - eigen_assert(false && "Cannot launch a kernel from another kernel"); -#endif + // FIXME: Should be device and kernel specific. #ifdef __CUDACC__ -- cgit v1.2.3