title: Efficient zero padding using cudaMemcpy3D
body: <p>I would like to transfer a 3d array stored in linear memory on the host, into a larger (3D) array on the device. As an example (see below), I tried to transfer a (3x3x3) array into a (5x5x3) array.
  I expect that on the host I get 2D slices with the following pattern:</p>
  <pre><code>x x x 0 0
  x x x 0 0
  x x x 0 0
  0 0 0 0 0 
  0 0 0 0 0
  </code></pre>
  <p>where x are the values of my array. However, I get something like that, where <code>y</code> are the values of the next 2D slice:</p>
  <pre><code>x x x 0 0
  x x x 0 0
  x x x 0 0
  y y y 0 0 
  y y y 0 0
  </code></pre>
  <p>According to the <a href="http://developer.download.nvidia.com/compute/cuda/2_3/toolkit/docs/online/group__CUDART__MEMORY_gc1372614eb614f4689fbb82b4692d30a.html#gc1372614eb614f4689fbb82b4692d30a" rel="nofollow noreferrer">cudaMemcpy3D documentation</a> I would have expect that the <code>extent</code> parameter would take into account the padding in the vertical axis but apparently not.
  Am I mistaken in the understanding of the documentation? If yes, is there any other way to perform this operation? The final size of the array to transfer will be <code>60x60x900</code> into an array of size <code>1100x1500x900</code>. I use the zero padding to prepare a Fourier transform.</p>
  <p>Here is the simplified code that I used:</p>
  <pre><code>cudaError_t cuda_status;
  cudaPitchedPtr d_ptr;
  cudaExtent d_extent = make_cudaExtent(sizeof(int)*5,sizeof(int)*5,sizeof(int)*3);
  cudaExtent h_extent = make_cudaExtent(sizeof(int)*3,sizeof(int)*3,sizeof(int)*3);
  
  int* h_array = (int*) malloc(27*sizeof(int));
  int* h_result = (int*) malloc(512*sizeof(int)*5*3);
  for (int i = 0; i&lt;27; i++)
  {
      h_array[i] = i;
  }
  cuda_status = cudaMalloc3D(&amp;d_ptr, d_extent);
  cout &lt;&lt; cudaGetErrorString(cuda_status) &lt;&lt; endl;
  
  cudaMemcpy3DParms myParms = {0};
  myParms.extent = h_extent;
  myParms.srcPtr.ptr = h_array;
  myParms.srcPtr.pitch = 3*sizeof(int);
  myParms.srcPtr.xsize = 3*sizeof(int);
  myParms.srcPtr.ysize = 3*sizeof(int);
  myParms.dstPtr = d_ptr;
  myParms.kind = cudaMemcpyHostToDevice;
  
  cuda_status = cudaMemcpy3D(&amp;myParms);
  cout &lt;&lt; cudaGetErrorString(cuda_status) &lt;&lt; endl;
  
  cout &lt;&lt; &quot;Pitch: &quot; &lt;&lt; d_ptr.pitch &lt;&lt; &quot; / xsize:&quot; &lt;&lt; d_ptr.xsize &lt;&lt; &quot; / ysize:&quot; &lt;&lt; d_ptr.ysize &lt;&lt; endl; // returns Pitch: 512 / xsize:20 / ysize:20 which is as expected
  
  // Copy array to host to be able to print the values - may not be necessary
  cout &lt;&lt; cudaMemcpy(h_result, (int*) d_ptr.ptr, 512*5*3, cudaMemcpyDeviceToHost) &lt;&lt; endl;
  cout &lt;&lt; h_result[128] &lt;&lt; &quot; &quot; &lt;&lt; h_result[3*128] &lt;&lt; &quot; &quot; &lt;&lt; h_result[5*128] &lt;&lt; &quot; &quot; &lt;&lt; endl; // output : 3 9 15 / expected 3 0 9
  </code></pre>