extern __shared__ uint16_t shared[];
__global__ void ComputeDiff(
    uint32_t* data1,
    uint32_t* data2,
    uint32_t mx,
    uint32_t my,
    uint32_t x0,
    uint32_t y0,
    uint32_t block_w,
    uint32_t block_h,
    uint32_t search_w,
    uint32_t search_h,
    size_t pitch_data1,
    size_t pitch_data2,
    float* rmse)
{
    int i, j;
    uint32_t* mrow;
    uint32_t* vrow;
    uint32_t delta;
    uint32_t sum = 0;
    uint32_t shared_w = search_w+block_w;

    /*
    // version 1 start--
    {
        for (i = threadIdx.y; i < search_h+block_h; i += blockDim.y)
        {
            row2 = (uint16_t*)((char*)data2 + (y0+i)*pitch_data2);

            for (j = threadIdx.x; j < shared_w; j += blockDim.x)
            {
                shared[threadIdx.y*blockDim.x+threadIdx.x] = row2[x0+j];
                //shared[i*shared_w+j] = row2[x0+j];
            }
        }
    }
    __syncthreads();
    // --version 1 end
     */

    uint32_t idx = threadIdx.x;
    uint32_t bidx = blockIdx.x * blockDim.x + idx;

    uint32_t sx = bidx % search_w;
    uint32_t sy = bidx / search_w;

    if (sx < search_w && sy < search_h)
    {
        for (i = 0; i < block_h; i++)
        {
            mrow = (uint32_t*)((char*)data1 + ((my+i)*pitch_data1));
            vrow = (uint32_t*)((char*)data2 + ((y0+sy+i)*pitch_data2));

            for (j = 0; j < block_w; j++)
            {
                // version 2:
                delta = mrow[mx+j] - vrow[x0+sx+j];
                // version 1:
                //delta = row1[mx+j] - shared[(sy+i)*shared_w + (sx+j)];
                //delta = row1[mx+j] - shared[(sy+i)*shared_w + (sx+j)];
                sum += delta*delta;
            }
        }
        rmse[sy*search_w+sx] = sqrtf((float) sum / (block_w*block_h));
    }

}