title: cudaMemcpy error from Device to Host
body: <p>I am returning a two-dimensional structure after computation on a kernel, from device to host.</p>
  
  <pre><code>HANDLE_ERROR(cudaMemcpy(Pixel,Pixel_gpu,img_wd*img_ht*sizeof(pixel),cudaMemcpyDeviceToHost));
  </code></pre>
  
  <p>Pixel is declared on host, Pixel_gpu is allocated on device, as below:</p>
  
  <pre><code>**Pixel_gpu;
  HANDLE_ERROR(cudaMalloc(&amp;Pixel_gpu,img_wd*img_ht*sizeof(pixel)));
  
  pixel **Pixel = (pixel**)malloc((img_ht)*sizeof(pixel*));
  for(int i=0;i&lt;(img_ht);i++)
      Pixel[i]=(pixel*)malloc((img_wd)*sizeof(pixel));
  </code></pre>
  
  <p>Using this I end up getting illegal memory access error. </p>
  
  <p>Trying a similar memory alignment for result, doesn't help either.</p>
  
  <pre><code> pixel *Pixel_res = (pixel*)malloc(img_wd*img_ht*sizeof(pixel));
  
  
  
  HANDLE_ERROR(cudaMemcpy(Pixel_res,Pixel_gpu,img_wd*img_ht*sizeof(pixel),cudaMemcpyDeviceToHost));
  </code></pre>
  
  <p>Kernel launching:</p>
  
  <pre><code>cudaDeviceProp prop;
  HANDLE_ERROR(cudaGetDeviceProperties(&amp;prop, 0));
  
  
  int thread_block=sqrt(prop.maxThreadsPerBlock);
  dim3 DimGrid(ceil(img_wd/thread_block),ceil(img_ht/thread_block),1);
  dim3 DimBlock(sqrt(prop.maxThreadsPerBlock),sqrt(prop.maxThreadsPerBlock),1);
  
  //allocating gpu memory
  
  
  pixel **Pixel_tmp_gpu, **Pixel_gpu;
  
  
  HANDLE_ERROR(cudaMalloc(&amp;Pixel_tmp_gpu,img_wd*img_ht*sizeof(pixel)));
  HANDLE_ERROR(cudaMalloc(&amp;Pixel_gpu,img_wd*img_ht*sizeof(pixel)));
  
  
  float **kernel0_gpu, **kernel1_gpu;
  
  HANDLE_ERROR(cudaMalloc(&amp;kernel0_gpu,k*1*sizeof(float)));
  HANDLE_ERROR(cudaMalloc(&amp;kernel1_gpu,1*k*sizeof(float)));
  
  cout&lt;&lt;"memory allocated"&lt;&lt;endl;
  
  //copying needed data
  
  HANDLE_ERROR(cudaMemcpy(Pixel_tmp_gpu,Pixel_tmp,img_wd*img_ht*sizeof(pixel),cudaMemcpyHostToDevice));
  HANDLE_ERROR(cudaMemcpy(Pixel_gpu,Pixel,img_wd*img_ht*sizeof(pixel),cudaMemcpyHostToDevice));
  HANDLE_ERROR(cudaMemcpy(kernel0_gpu,kernel0,k*1*sizeof(float),cudaMemcpyHostToDevice));
  HANDLE_ERROR(cudaMemcpy(kernel1_gpu,kernel1,1*k*sizeof(float),cudaMemcpyHostToDevice));
  
  cout&lt;&lt;"memory transfers done"&lt;&lt;endl;
  
  vertical_conv&lt;&lt;&lt;DimGrid,DimBlock&gt;&gt;&gt;(Pixel_gpu, Pixel_tmp_gpu,img_wd, img_ht,kernel0_gpu,k);
  time_t vertical_convolution=time(NULL);
  
  cout&lt;&lt;" vertical_convolution time: "&lt;&lt;double(vertical_convolution - reading_file)&lt;&lt;"sec"&lt;&lt;endl;
  
  
  horizontal_conv&lt;&lt;&lt;DimGrid,DimBlock&gt;&gt;&gt;(Pixel_tmp_gpu, Pixel_gpu, img_wd, img_ht, kernel1_gpu, k);
  time_t horizontal_convolution=time(NULL);
  
  cout&lt;&lt;" horizontal convolution time:" &lt;&lt;double(horizontal_convolution-vertical_convolution)&lt;&lt;" sec"&lt;&lt;endl;
  
  pixel *Pixel_res = (pixel*)malloc(img_wd*img_ht*sizeof(pixel));
  
  HANDLE_ERROR(cudaMemcpy(Pixel_res,Pixel_gpu,img_wd*img_ht*sizeof(pixel),cudaMemcpyDeviceToHost));
  </code></pre>
  
  <p>The functions used:</p>
  
  <pre><code>struct  pixel //to store RGB values
  {
      unsigned char r;
      unsigned char g;
      unsigned char b;
  };
  
  static void HandleError( cudaError_t err, const char *file, int line ) {
      if (err != cudaSuccess) {
          cout&lt;&lt;cudaGetErrorString(err)&lt;&lt;" in "&lt;&lt; file &lt;&lt;" at line "&lt;&lt; line&lt;&lt;endl;
      }
  }
  
  #define HANDLE_ERROR( err ) (HandleError( err, __FILE__, __LINE__ ))
  
  __device__ void padding(pixel** Pixel_val, int x_coord, int y_coord, int img_width, int img_height, pixel Px) //padding the image,depending on pixel coordinates, can be replaced by reflect for better result //currently zero padding
  {
      if(x_coord&lt;img_width &amp;&amp; y_coord&lt;img_height &amp;&amp; x_coord&gt;=0 &amp;&amp; y_coord&gt;=0) 
          Px=Pixel_val[y_coord][x_coord];
  }
  </code></pre>
  
  <p>The vertical convolution:</p>
  
  <pre><code>__global__ void vertical_conv(pixel** Pixel_in, pixel** Pixel_out,int img_wd, int img_ht, float** kernel, int k)
  {
      float tmp_r, tmp_g, tmp_b;
      pixel pix_val;
      pix_val.r=0;pix_val.g=0;pix_val.b=0;
      int row=blockIdx.y*blockDim.y + threadIdx.y;
  
      int col = blockIdx.x*blockDim.x + threadIdx.x;
      if(row&lt;img_ht &amp;&amp; col&lt;img_wd){
          tmp_r=0, tmp_g=0, tmp_b=0;
          for(int l=0;l&lt;k;l++)
          {   
  
              padding(Pixel_in, col, row+l-(k-1)/2, img_wd, img_ht, pix_val);
              tmp_r+=pix_val.r * kernel[l][0];
              tmp_b+=pix_val.b * kernel[l][0];
              tmp_g+=pix_val.g * kernel[l][0];
          }
  
          Pixel_out[row][col].r=tmp_r;
          Pixel_out[row][col].g=tmp_g;
          Pixel_out[row][col].b=tmp_b;
      }
  }
  </code></pre>
  
  <p>The horizontal convolution:</p>
  
  <pre><code>   __global__ void horizontal_conv(pixel** Pixel_in, pixel** Pixel_out, int img_wd, int img_ht, float** kernel, int k)
  {
      float tmp_r, tmp_b, tmp_g;
      pixel pix_val;
      pix_val.r=0;pix_val.g=0;pix_val.b=0;
  
      //horizontal convolution
      int row=blockIdx.y*blockDim.y + threadIdx.y;
  
      int col = blockIdx.x*blockDim.x + threadIdx.x;
      tmp_r=0, tmp_g=0, tmp_b=0;
      if(row&lt;img_ht &amp;&amp; col&lt;img_wd)
      {
          for(int l=0; l&lt;k;l++)
          {
              padding(Pixel_in, col+l-(k-1)/2, row, img_wd, img_ht, pix_val);
              tmp_r+=pix_val.r * kernel[0][l];
              tmp_g+=pix_val.g * kernel[0][l];
              tmp_b+=pix_val.b * kernel[0][l];
          }
          Pixel_out[row][col].r=tmp_r;
          Pixel_out[row][col].g=tmp_g;
          Pixel_out[row][col].b=tmp_b;
      }
  }
  </code></pre>
  
  <p>Can someone help me know what could be wrong here?</p>