如何在 CUDA 中的 CPU / GPU 之间分配内存和复制 2D 数组而不使它们变平？

Question

所以我想在CUDA中分配2D数组并在CPU和GPU之间复制它们，但是我是一个完全初学者，其他在线资料对我来说很难理解或不完整。重要的是我能够在内核代码中以二维数组的形式访问它们，如下所示。

请注意，数组的高度 != 宽度，如果可能的话，这会让我更加困惑，因为我总是难以选择网格大小。

我考虑过将它们展平，但我真的想让它以这种方式工作。

这是我自己研究的结果。

__global__ void myKernel(int *firstArray, int *secondArray, int rows, int columns) {
    int row = blockIdx.x * blockDim.x + threadIdx.x;
    int column = blockIdx.y * blockDim.y + threadIdx.y;

    if (row >= rows || column >= columns)
        return;

    // Do something with the arrays like you would on a CPU, like:
    firstArray[row][column] = row * 2;
    secondArray[row[column] = row * 3;  
}


int main() {
    int rows = 300, columns = 200;
    int h_firstArray[rows][columns], h_secondArray[rows][columns];
    int *d_firstArray[rows][columns], *d_secondArray[rows][columns];

    // populate h_ arrays (Can do this bit myself)

    // Allocate memory on device, no idea how to do for 2D arrays.
    // Do memcopies to GPU, no idea how to do for 2D arrays.

    dim3 block(rows,columns);
    dim3 grid (1,1);
    myKernel<<<grid,block>>>(d_firstArray, d_secondArray, rows, columns);

    // Do memcopies back to host, no idea how to do for 2D arrays.

    cudaFree(d_firstArray);
    cudaFree(d_secondArray);

    return 0;
}

编辑：有人问我在我试图解决的问题中是否会在编译时知道数组宽度。您可以假设我目前主要对这种特殊情况感兴趣。

Answer 1

在一般情况下（直到运行时才知道数组维度），在 CUDA 设备代码中处理双下标访问需要一个指针数组，就像在主机代码中一样。 C 和 C++ 将每个下标作为指针解引用处理，以便到达“二维数组”中的最终位置。

从cuda tag info page 链接的canonical answer 已经涵盖了一般情况下设备代码中的双指针/双下标访问。这样做有几个缺点，该答案已涵盖，因此我不会在此重复。

但是，如果数组 width 在编译时是已知的（数组高度可以是动态的 - 即在运行时确定），那么我们可以利用编译器和语言类型机制来规避大多数缺点。您的代码演示了 CUDA 和/或 C/C++ 使用的其他几种错误模式：

1. 不能使用像int *firstarray 这样的简单单指针类型来传递对 C 或 C++ 函数进行双下标访问的项

2. 通过基于堆栈的机制分配大型主机阵列：

   int h_firstArray[rows][columns], h_secondArray[rows][columns];
   

   在 C 和 C++ 中经常出现问题。这些是基于堆栈的变量，如果足够大，通常会遇到堆栈限制。

3. CUDA 线程块总数限制为 1024 个线程。因此这样的线程块维度：

   dim3 block(rows,columns);
   

   除非rows 和columns 的尺寸非常小（乘积必须小于或等于1024），否则将无法使用。

4. 在 CUDA 中为设备数组声明指针变量时，创建指针数组几乎是不正确的：

   int *d_firstArray[rows][columns], *d_secondArray[rows][columns];
   

   我们也不在主机上分配空间，然后“重新分配”这些指针以供设备使用。

以下是一个工作示例，其中解决了上述项目并演示了上述方法，其中数组宽度在运行时是已知的：

$ cat t50.cu
#include <stdio.h>

const int array_width = 200;

typedef int my_arr[array_width];

__global__ void myKernel(my_arr *firstArray, my_arr *secondArray, int rows, int columns) {
    int column = blockIdx.x * blockDim.x + threadIdx.x;
    int row = blockIdx.y * blockDim.y + threadIdx.y;

    if (row >= rows || column >= columns)
        return;

    // Do something with the arrays like you would on a CPU, like:
    firstArray[row][column] = row * 2;
    secondArray[row][column] = row * 3;
}


int main() {
    int rows = 300, columns = array_width;
    my_arr *h_firstArray, *h_secondArray;
    my_arr *d_firstArray, *d_secondArray;
    size_t dsize = rows*columns*sizeof(int);
    h_firstArray = (my_arr *)malloc(dsize);
    h_secondArray = (my_arr *)malloc(dsize);
    // populate h_ arrays
    memset(h_firstArray, 0, dsize);
    memset(h_secondArray, 0, dsize);

    // Allocate memory on device
    cudaMalloc(&d_firstArray, dsize);
    cudaMalloc(&d_secondArray, dsize);
    // Do memcopies to GPU
    cudaMemcpy(d_firstArray, h_firstArray, dsize, cudaMemcpyHostToDevice);
    cudaMemcpy(d_secondArray, h_secondArray, dsize, cudaMemcpyHostToDevice);

    dim3 block(32,32);
    dim3 grid ((columns+block.x-1)/block.x,(rows+block.y-1)/block.y);
    myKernel<<<grid,block>>>(d_firstArray, d_secondArray, rows, columns);

    // Do memcopies back to host
    cudaMemcpy(h_firstArray, d_firstArray, dsize, cudaMemcpyDeviceToHost);
    cudaMemcpy(h_secondArray, d_secondArray, dsize, cudaMemcpyDeviceToHost);
    // validate
    if (cudaGetLastError() != cudaSuccess) {printf("cuda error\n"); return -1;}
    for (int i = 0; i < rows; i++)
      for (int j = 0; j < columns; j++){
        if (h_firstArray[i][j] != i*2) {printf("first mismatch at %d,%d, was: %d, should be: %d\n", i,j,h_firstArray[i][j], i*2); return -1;}
        if (h_secondArray[i][j] != i*3) {printf("second mismatch at %d,%d, was: %d, should be: %d\n", i,j,h_secondArray[i][j], i*3); return -1;}}

    printf("success!\n");


    cudaFree(d_firstArray);
    cudaFree(d_secondArray);

    return 0;
}
$ nvcc -arch=sm_61 -o t50 t50.cu
$ cuda-memcheck ./t50
========= CUDA-MEMCHECK
success!
========= ERROR SUMMARY: 0 errors
$

我已经颠倒了内核索引 (x,y) 的含义，以帮助合并全局内存访问。我们看到，通过这种类型的创建，我们可以利用编译器和语言特性来最终得到一个允许在主机和设备代码中进行双下标访问的代码，同时允许 CUDA 操作（例如cudaMemcpy）就好像我们正在处理单指针（例如“扁平化”）数组。