title: cudaMemcpyFromSymbol on a __device__ variable
body: <p>I am trying to apply a kernel function on a <code>__device__</code> variable, which, according to the specs, resides "in global memory"</p>
  
  <pre><code>#include &lt;stdio.h&gt;
  #include "sys_data.h"
  #include "my_helper.cuh"
  #include "helper_cuda.h"
  #include &lt;cuda_runtime.h&gt;
  
  
  double X[10] = {1,-2,3,-4,5,-6,7,-8,9,-10};
  double Y[10] = {0};
  __device__ double DEV_X[10];
  
  
  int main(void) {
      checkCudaErrors(cudaMemcpyToSymbol(DEV_X, X,10*sizeof(double)));
      vector_projection&lt;double&gt;&lt;&lt;&lt;1,10&gt;&gt;&gt;(DEV_X, 10);
      getLastCudaError("oops");
      checkCudaErrors(cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double)));
      return 0;
  }
  </code></pre>
  
  <p>The kernel function <code>vector_projection</code> is defined in <code>my_helper.cuh</code> as follows:</p>
  
  <pre><code>template&lt;typename T&gt; __global__ void vector_projection(T *dx, int n) {
      int tid;
      tid = threadIdx.x + blockIdx.x * blockDim.x;
      if (tid &lt; n) {
          if (dx[tid] &lt; 0)
              dx[tid] = (T) 0;
      }
  }
  </code></pre>
  
  <p>As you can see, I use <code>cudaMemcpyToSymbol</code> and <code>cudaMemcpyFromSymbol</code> to transfer data to and from the device. However, I'm getting the following error:</p>
  
  <pre><code>CUDA error at ../src/vectorAdd.cu:19 code=4(cudaErrorLaunchFailure) 
    "cudaMemcpyFromSymbol(Y, DEV_X, 10*sizeof(double))" 
  </code></pre>
  
  <p><strong>Footnote:</strong> I can of course avoid to use <code>__device__</code> variables and go for something <a href="http://pastebin.com/rx9nUnGX" rel="nofollow">like this</a> which works fine; I just want to see how to do the same thing (if possible) with <code>__device__</code> variables. </p>
  
  <p><strong>Update:</strong> The output of <code>cuda-memcheck</code> can be found at <a href="http://pastebin.com/AW9vmjFs" rel="nofollow">http://pastebin.com/AW9vmjFs</a>. The error messages I get are as follows:</p>
  
  <pre><code>========= Invalid __global__ read of size 8
  =========     at 0x000000c8 in /home/ubuntu/Test0001/Debug/../src/my_helper.cuh:75:void vector_projection&lt;double&gt;(double*, int)
  =========     by thread (9,0,0) in block (0,0,0)
  =========     Address 0x000370e8 is out of bounds
  </code></pre>