title: cudaMemcpy segmentation fault
body: <p>I've been haunted by this error for quite a while so I decided to post it here.</p>
  
  <p>This segmentation fault happened when a cudaMemcpy is called:</p>
  
  <pre><code>CurrentGrid-&gt;cdata[i] = new float[size];
  cudaMemcpy(CurrentGrid-&gt;cdata[i], Grid_dev-&gt;cdata[i], size*sizeof(float),\
                  cudaMemcpyDeviceToHost);
  </code></pre>
  
  <p><code>CurrentGrid</code> and <code>Grid_dev</code> are pointer to a <code>grid</code> class object on host and device respectively and i=0 in this context. Class member <code>cdata</code> is a float type pointer array. For debugging, right before this cudaMemcpy call I printed out the value of each element of <code>Grid_Dev-&gt;cdata[i]</code>, the address of <code>CurrentGrid-&gt;cdata[i]</code> and <code>Grid_dev-&gt;cdata[i]</code> and the value of <code>size</code>, which all looks good. But it still ends up with "Segmentation fault (core dumped)", which is the only error message. cuda-memcheck only gave "process didn't terminate successfully". I'm not able to use cuda-gdb at the moment. Any suggestion about where to go? </p>
  
  <p><strong>UPDATE</strong>: It seems now I have solved this problem by cudaMalloc another float pointer A on device and cudaMemcpy the value of Grid_dev->cdata[i] to A, and then cudaMemcpy A to host.
  So the segment of code written above becomes:</p>
  
  <pre><code>float * A;
  cudaMalloc((void**)&amp;A, sizeof(float));
  ...
  ...
  cudaMemcpy(&amp;A, &amp;(Grid_dev-&gt;cdata[i]), sizeof(float *), cudaMemcpyDeviceToHost);    
  CurrentGrid-&gt;cdata[i] = new float[size];
  cudaMemcpy(CurrentGrid-&gt;cdata[i], A, size*sizeof(float), cudaMemcpyDeviceToHost);            
  </code></pre>
  
  <p>I did this because valgrind popped up "invalid read of size 8", which I thought referring to <code>Grid_dev-&gt;cdata[i]</code>. I checked it again with gdb, printing out the value of <code>Grid_dev-&gt;cdata[i]</code> being NULL. So I guess I cannot directly dereference the device pointer even in this cudaMemcpy call.  But why ?  According to the comment at the bottom of this <a href="https://stackoverflow.com/questions/14790999/how-to-pass-a-c-class-with-array-of-pointers-to-cuda">thread</a> , we should be able to dereference device pointer in cudaMemcpy function.</p>
  
  <p>Also, I don't know the  the underlying mechanism of how cudaMalloc and cudaMemcpy work but I think by cudaMalloc a pointer, say A here, we actually assign this pointer to point to a certain address on the device. And by cudaMemcpy the <code>Grid_dev-&gt;cdata[i]</code> to A as in the modified code above, we re-assign the pointer A to point to the array. Then don't we lose the track of the previous address that A pointed to when it is cudaMalloced? Could this cause memory leak or something? If yes, how should I work around this situation properly?
  Thanks!</p>
  
  <p>For reference I put the code of the complete function in which this error happened  below.</p>
  
  <p>Many thanks!</p>
  
  <pre><code>__global__ void Print(grid *, int);
  __global__ void Printcell(grid *, int);
  void CopyDataToHost(param_t p, grid * CurrentGrid, grid * Grid_dev){
  
      cudaMemcpy(CurrentGrid, Grid_dev, sizeof(grid), cudaMemcpyDeviceToHost);
  #if DEBUG_DEV
      cudaCheckErrors("cudaMemcpy1 error");
  #endif
      printf("\nBefore copy cell data\n");
      Print&lt;&lt;&lt;1,1&gt;&gt;&gt;(Grid_dev, 0);            //Print out some Grid_dev information for 
      cudaDeviceSynchronize();                //debug 
      int NumberOfBaryonFields = CurrentGrid-&gt;ReturnNumberOfBaryonFields();
      int size = CurrentGrid-&gt;ReturnSize();
      int vsize = CurrentGrid-&gt;ReturnVSize();
      CurrentGrid-&gt;FieldType = NULL;
      CurrentGrid-&gt;FieldType = new int[NumberOfBaryonFields];
      printf("CurrentGrid size is %d\n", size);
      for( int i = 0; i &lt; p.NumberOfFields; i++){
          CurrentGrid-&gt;cdata[i] = NULL;
          CurrentGrid-&gt;vdata[i] = NULL;
          CurrentGrid-&gt;cdata[i] = new float[size];
          CurrentGrid-&gt;vdata[i] = new float[vsize];
  
          Printcell&lt;&lt;&lt;1,1&gt;&gt;&gt;(Grid_dev, i);//Print out element value of Grid_dev-&gt;cdata[i]
          cudaDeviceSynchronize();        
  
          cudaMemcpy(CurrentGrid-&gt;cdata[i], Grid_dev-&gt;cdata[i], size*sizeof(float),\
                  cudaMemcpyDeviceToHost);               //where error occurs
  #if DEBUG_DEV
          cudaCheckErrors("cudaMemcpy2 error");
  #endif
          printf("\nAfter copy cell data\n");
          Print&lt;&lt;&lt;1,1&gt;&gt;&gt;(Grid_dev, i);
          cudaDeviceSynchronize();
          cudaMemcpy(CurrentGrid-&gt;vdata[i], Grid_dev-&gt;vdata[i], vsize*sizeof(float),\
                  cudaMemcpyDeviceToHost);
  #if DEBUG_DEV
          cudaCheckErrors("cudaMemcpy3 error");
  #endif
      }
      cudaMemcpy(CurrentGrid-&gt;FieldType, Grid_dev-&gt;FieldType,\
              NumberOfBaryonFields*sizeof(int), cudaMemcpyDeviceToHost);
  #if DEBUG_DEV
      cudaCheckErrors("cudaMemcpy4 error");
  #endif
  }
  </code></pre>
  
  <p>EDIT: here is the information from valgrind, from which I'm trying to track down where the memory leak happened.</p>
  
  <pre><code>==19340== Warning: set address range perms: large range [0x800000000, 0xd00000000) (noaccess)
  ==19340== Warning: set address range perms: large range [0x200000000, 0x400000000) (noaccess)
  ==19340== Invalid read of size 8
  ==19340==    at 0x402C79: CopyDataToHost(param_t, grid*, grid*) (CheckDevice.cu:48)
  ==19340==    by 0x403646: CheckDevice(param_t, grid*, grid*) (CheckDevice.cu:186)
  ==19340==    by 0x40A6CD: main (Transport.cu:81)
  ==19340==  Address 0x2003000c0 is not stack'd, malloc'd or (recently) free'd
  ==19340== 
  ==19340== 
  ==19340== Process terminating with default action of signal 11 (SIGSEGV)
  ==19340==  Bad permissions for mapped region at address 0x2003000C0
  ==19340==    at 0x402C79: CopyDataToHost(param_t, grid*, grid*) (CheckDevice.cu:48)
  ==19340==    by 0x403646: CheckDevice(param_t, grid*, grid*) (CheckDevice.cu:186)
  ==19340==    by 0x40A6CD: main (Transport.cu:81)
  ==19340== 
  ==19340== HEAP SUMMARY:
  ==19340==     in use at exit: 2,611,365 bytes in 5,017 blocks
  ==19340==   total heap usage: 5,879 allocs, 862 frees, 4,332,278 bytes allocated
  ==19340== 
  ==19340== LEAK SUMMARY:
  ==19340==    definitely lost: 0 bytes in 0 blocks
  ==19340==    indirectly lost: 0 bytes in 0 blocks
  ==19340==      possibly lost: 37,416 bytes in 274 blocks
  ==19340==    still reachable: 2,573,949 bytes in 4,743 blocks
  ==19340==         suppressed: 0 bytes in 0 blocks
  ==19340== Rerun with --leak-check=full to see details of leaked memory
  ==19340== 
  ==19340== For counts of detected and suppressed errors, rerun with: -v
  ==19340== ERROR SUMMARY: 1 errors from 1 contexts (suppressed: 2 from 2)
  </code></pre>