arrays - 二维数组上的CUDA网格跨度循环

Question

我试图在CUDA上高效地遍历二维数组。在主机代码中，我有

double **h_matrix; // Matrix on host of size Nx by Ny
double tmp;
...
for(i = 0; i < Nx; i++) {
    for(j = 0; j < Ny; j++) {
        tmp = h_matrix[i][j];
        ... // Perform some operation on tmp
        h_matrix[i][j] = tmp;
    }
}

为了在CUDA中有效执行相似的任务，我了解到必须使用 cudaMallocPitch()为2D数组分配内存，如 CUDA Programming guide所示(例如，滚动一下)。该示例并没有多大帮助，因为即使内核以 <<<100, 512>>>的形式启动，该内核也不会使用任何有关网格，执行块或执行线程的信息。

NVidia'a并行的所有博客 suggests使用网格跨度循环来编写灵活且可伸缩的内核，但是，它们的示例仅使用一维数组。如何为使用 cudaMallocPitch()分配的2D数组编写网格步幅循环以并行化上面显示的代码？我应该使用2D dimGrid和dimBlock吗？

Answer 1

网格步幅循环概念在2D情况下与cudaMallocPitch分配的2D矩阵有关的扩展可能看起来像:

#define N 11
#define M 3

__global__ void kernel(float * d_matrix, size_t pitch) {

    int idx = blockIdx.x*blockDim.x + threadIdx.x;
    int idy = blockIdx.y*blockDim.y + threadIdx.y;

    for (int j = blockIdx.y * blockDim.y + threadIdx.y; j < N; j += blockDim.y * gridDim.y) 
    {
        float* row_d_matrix = (float*)((char*)d_matrix + idy*pitch);
        for (int i = blockIdx.x * blockDim.x + threadIdx.x; i < M; i += blockDim.x * gridDim.x) {
            row_d_matrix[i] = ....
       } 

    }

}

int main()
{

    float *d_matrix;

    size_t pitch;
    cudaMallocPitch(&d_matrix,&pitch,M*sizeof(float),N);

    kernel<<<GridSize,BlockSize>>>(d_matrix,pitch);

    // Other stuff

}