Extracting a sub-matrix from a larger one with cudaMemcpy2D not copying correct data
<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>