如何在 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 标签信息页面链接的规范答案中。这样做有几个缺点，该答案已涵盖这些缺点，因此我不会在这里重复。

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

1. 无法使用简单的单指针类型来传递对 C 或 C++ 函数进行双下标访问的项目，例如int *firstarray

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），就像我们一样正在处理单指针（例如“扁平化”）数组。