From dfa991cbae98cde7db5aef5ff1bb4b3d51cc362b Mon Sep 17 00:00:00 2001 From: Benoit Steiner Date: Tue, 21 Apr 2015 16:15:45 -0700 Subject: [PATCH] Make sure that the copy constructor of the evaluator is always called before launching the evaluation of a tensor expression on a cuda device. --- unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h | 12 ++++++++++-- 1 file changed, 10 insertions(+), 2 deletions(-) diff --git a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h index bb2f8b977..02e1667b9 100644 --- a/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h +++ b/unsupported/Eigen/CXX11/src/Tensor/TensorExecutor.h @@ -157,7 +157,11 @@ class TensorExecutor template __global__ void __launch_bounds__(1024) -EigenMetaKernel_NonVectorizable(Evaluator eval, Index size) { +EigenMetaKernel_NonVectorizable(Evaluator memcopied_eval, Index size) { + // Cuda memcopies the kernel arguments. That's fine for POD, but for more + // complex types such as evaluators we should really conform to the C++ + // standard and call a proper copy constructor. + Evaluator eval(memcopied_eval); const Index first_index = blockIdx.x * blockDim.x + threadIdx.x; const Index step_size = blockDim.x * gridDim.x; @@ -171,7 +175,11 @@ EigenMetaKernel_NonVectorizable(Evaluator eval, Index size) { template __global__ void __launch_bounds__(1024) -EigenMetaKernel_Vectorizable(Evaluator eval, Index size) { +EigenMetaKernel_Vectorizable(Evaluator memcopied_eval, Index size) { + // Cuda memcopies the kernel arguments. That's fine for POD, but for more + // complex types such as evaluators we should really conform to the C++ + // standard and call a proper copy constructor. + Evaluator eval(memcopied_eval); const Index first_index = blockIdx.x * blockDim.x + threadIdx.x; const Index step_size = blockDim.x * gridDim.x;