title: CUDA, "illegal memory access was encountered" in Memcpy
body: <p>I have this cuda file:</p>
  
  <pre><code>#include "cuda.h"
  #include "../../HandleError.h"
  #include "Sphere.hpp"
  #include &lt;stdlib.h&gt;
  #include &lt;CImg.h&gt;
  
  #define WIDTH 1280
  #define HEIGHT 720
  #define rnd(x) (x*rand()/RAND_MAX)
  #define SPHERES_COUNT 5
  
  using namespace cimg_library;
  
  __global__
  void kernel(unsigned char* bitmap, Sphere* s)
  {
     // Map threadIdx/blockIdx to pixel position
     int x = threadIdx.x + blockIdx.x * blockDim.x;
     int y = threadIdx.y + blockIdx.y * blockDim.y;
     int offset = x + y * blockDim.x * gridDim.x;
     float ox = x - blockDim.x * gridDim.x / 2;
     float oy = y - blockDim.y * gridDim.y / 2;
     float r = 0.2, g = 0.2, b = 0.5;
     float maxz = -INF;
     for (int i = 0; i &lt; SPHERES_COUNT; i++) {
         float n, t = s[i].hit(ox, oy, &amp;n);
         if (t &gt; maxz) {
             float fscale = n;
             r = s[i].r * fscale;
             g = s[i].g * fscale;
             b = s[i].b * fscale;
             maxz = t;
         }
     }
  
     bitmap[offset*3] = (int)(r * 255);
     bitmap[offset*3 + 1] = (int)(g * 255);
     bitmap[offset*3 + 2] = (int)(b * 255);
  }
  
  __constant__ Sphere s[SPHERES_COUNT];
  
  int main ()
  {
      //Capture start time
      cudaEvent_t start, stop;
      HANDLE_ERROR(cudaEventCreate(&amp;start));
      HANDLE_ERROR(cudaEventCreate(&amp;stop));
      HANDLE_ERROR(cudaEventRecord(start, 0));
  
      //Create host bitmap
      CImg&lt;unsigned char&gt; image(WIDTH, HEIGHT, 1, 3);
      image.permute_axes("cxyz");
  
      //Allocate device bitmap data
      unsigned char* dev_bitmap;
      HANDLE_ERROR(cudaMalloc((void**)&amp;dev_bitmap, image.size()*sizeof(unsigned char)));
  
      //Generate spheres and copy them on the GPU one by one
      Sphere* temp_s = (Sphere*)malloc(SPHERES_COUNT*sizeof(Sphere));
      for (int i=0; i &lt;SPHERES_COUNT; i++) {
          temp_s[i].r = rnd(1.0f);
          temp_s[i].g = rnd(1.0f);
          temp_s[i].b = rnd(1.0f);
          temp_s[i].x = rnd(1000.0f) - 500;
          temp_s[i].y = rnd(1000.0f) - 500;
          temp_s[i].z = rnd(1000.0f) - 500;
          temp_s[i].radius = rnd(100.0f) + 20;
      }
  
      HANDLE_ERROR(cudaMemcpyToSymbol(s, temp_s, sizeof(Sphere)*SPHERES_COUNT));
      free(temp_s);
  
      //Generate a bitmap from spere data
      dim3 grids(WIDTH/16, HEIGHT/16);
      dim3 threads(16, 16);
      kernel&lt;&lt;&lt;grids, threads&gt;&gt;&gt;(dev_bitmap, s);
  
      //Copy the bitmap back from the GPU for display
      HANDLE_ERROR(cudaMemcpy(image.data(), dev_bitmap,
                              image.size()*sizeof(unsigned char),
                              cudaMemcpyDeviceToHost));
  
      cudaFree(dev_bitmap);
  
      image.permute_axes("yzcx");
      image.save("render.bmp");
  }
  </code></pre>
  
  <p>It compiles fine, but when executed I get this error:</p>
  
  <pre><code>an illegal memory access was encountered in main.cu at line 82
  </code></pre>
  
  <p>that is, here:</p>
  
  <pre><code>    //Copy the bitmap back from the GPU for display
      HANDLE_ERROR(cudaMemcpy(image.data(), dev_bitmap,
                              image.size()*sizeof(unsigned char),
                              cudaMemcpyDeviceToHost));
  </code></pre>
  
  <p>I cannot understand why...
  I know that If remove this:</p>
  
  <pre><code>  bitmap[offset*3] = (int)(r * 255);
    bitmap[offset*3 + 1] = (int)(g * 255);
    bitmap[offset*3 + 2] = (int)(b * 255);
  </code></pre>
  
  <p>The error is not reported, so I thought It may be an out of index error, reported later, but I have An identical version of this program that makes no use of constant memory, and it works fine with the very same version of the <em>kernel</em> function...</p>