title: Using Eigen 3.3 in a CUDA kernel
body: <p>Since Nov. 2016 it's possible to compile CUDA code which references Eigen3.3 - see <a href="https://stackoverflow.com/a/39745481/4069571">this answer</a></p>
  
  <p><a href="https://stackoverflow.com/a/23804859/4069571">This answer</a> is not what I'm looking for and may now be "outdated" in the sense that there might now be an easier way, since the following is written in the <a href="http://eigen.tuxfamily.org/index.php?title=3.3#Experimental_CUDA_support" rel="noreferrer">docs</a></p>
  
  <blockquote>
    <p>Starting from Eigen 3.3, it is now possible to use Eigen's objects and
    algorithms within CUDA kernels. However, only a subset of features are
    supported to make sure that no dynamic allocation is triggered within
    a CUDA kernel.</p>
  </blockquote>
  
  <p>See also <a href="http://eigen.tuxfamily.org/dox/TopicCUDA.html" rel="noreferrer">here</a>. Unfortunately, I was not able to find any example of how this might look like.</p>
  
  <h2>My Question</h2>
  
  <p><strong>Is it now possible to write a kernel such as the following</strong>, which should simply calculate a bunch of dot products?</p>
  
  <pre><code>__global__ void cu_dot(Eigen::Vector3d *v1, Eigen::Vector3d *v2, double *out, size_t N)
  {
      int idx = blockIdx.x * blockDim.x + threadIdx.x;
      if(idx &lt; N)
      {
          out[idx] = v1[idx].dot(v2[idx]);
      }
      return;
  }
  </code></pre>
  
  <p>I can compile this, but it does not seem to work. When I try to copy the data to host, I get <code>illegal memory access</code>. Note that I originally store the Vector3d's as `std::vector and then respectively use </p>
  
  <pre><code>cudaMalloc((void **)&amp;p_device_v1, sizeof(Eigen::Vector3d)*n);
  cudaMemcpy(p_v1_device, v1.data(), sizeof(Eigen::Vector3d)*n, cudaMemcpyHostToDevice);
  </code></pre>
  
  <p>I have set up an <strong>MWE project</strong> using CMake at <a href="https://github.com/GPMueller/eigen-cuda" rel="noreferrer">https://github.com/GPMueller/eigen-cuda</a></p>