title: Extracting a sub-matrix from a larger one with cudaMemcpy2D not copying correct data
body: <p>I have two square matrices: d_img and d_template. I am trying to copy a region of d_img (in this case from the top left corner) into d_template using cudaMemcpy2D(). But it's not copying the correct ranges. Most questions I've found on this are because an argument is wrong. But I'm pretty sure I have that right. I've checked and there're no cudaErrors occurring.</p>
  
  <pre><code>int const TEMPLATE_DIM = 10;
  int const OFFSET_DIM = 1;
  int const IMG_DIM = 2 * OFFSET_DIM + TEMPLATE_DIM; //12
  size_t const TEMPLATE_DIM_BYTES = TEMPLATE_DIM * sizeof(int);
  size_t const IMG_DIM_BYTES = IMG_DIM * sizeof(int);
  
  int main(){
  
      //Larger matrix
      int h_img[IMG_DIM][IMG_DIM];
      int* d_img;
      size_t imgPitch;
      cudaMallocPitch(&amp;d_img, &amp;imgPitch, IMG_DIM_BYTES, IMG_DIM);
  
      //Subset matrix
      int h_template[TEMPLATE_DIM][TEMPLATE_DIM];
      int* d_template;
      size_t templatePitch;   
      cudaMallocPitch(&amp;d_template, &amp;templatePitch, TEMPLATE_DIM_BYTES, TEMPLATE_DIM);
  
      //populate h_img, copy to d_img
      srand(time(NULL)+1);
      for (int y = 0; y &lt; IMG_DIM; ++y)
          for (int x = 0; x &lt; IMG_DIM; ++x)   
              h_img[y][x] = y*IMG_DIM+x;
      cout &lt;&lt; "h_img: \n"; printTemplateImg(h_img);
      cudaMemcpy(d_img, h_img, IMG_DIM_BYTES*IMG_DIM, cudaMemcpyHostToDevice);
  
      //copy subset of d_img to d_template
      cudaMemcpy2D(d_template, templatePitch, d_img, imgPitch, TEMPLATE_DIM_BYTES, TEMPLATE_DIM, cudaMemcpyDeviceToDevice);
      //copy d_template to h_template to view it.
      cudaMemcpy(h_template, d_template, TEMPLATE_DIM_BYTES*TEMPLATE_DIM, cudaMemcpyDeviceToHost);
      cudaDeviceSynchronize();
      cout &lt;&lt; "h_template: \n"; printTemplate(h_template);
  }
  </code></pre>
  
  <p>And this is the output</p>
  
  <pre><code>h_img: 
  {
  {0,1,2,3,4,5,6,7,8,9,10,11,}
  {12,13,14,15,16,17,18,19,20,21,22,23,}
  {24,25,26,27,28,29,30,31,32,33,34,35,}
  {36,37,38,39,40,41,42,43,44,45,46,47,}
  {48,49,50,51,52,53,54,55,56,57,58,59,}
  {60,61,62,63,64,65,66,67,68,69,70,71,}
  {72,73,74,75,76,77,78,79,80,81,82,83,}
  {84,85,86,87,88,89,90,91,92,93,94,95,}
  {96,97,98,99,100,101,102,103,104,105,106,107,}
  {108,109,110,111,112,113,114,115,116,117,118,119,}
  {120,121,122,123,124,125,126,127,128,129,130,131,}
  {132,133,134,135,136,137,138,139,140,141,142,143,}
  }
  h_template: 
  {
  {0,1,2,3,4,5,6,7,8,9,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  {0,0,0,0,0,0,0,0,0,0,}
  }
  </code></pre>
  
  <p>Why does it only do the first row?
  Also, if you change TEMPLATE_DIM to 32 or 96, it gives weird row skipping patterns, in case that helps you see a pattern I'm missing.</p>