二维多维数组传递给内核，CUDA答案

【问题标题】：2D multidimensional array passing to kernel, CUDA二维多维数组传递给内核，CUDA
【发布时间】：2016-04-08 01:35:29
【问题描述】：

我一直在努力寻找一种允许我将预定义的二维数组传递给内核的实现。

int values[2][3];

我需要将所有数据保存在正确的列和行中。现在我知道 CUDA 接受线性形式的 2D 数组，但是如何传递已经构建的数组？

【问题讨论】：

int [2][3] 的布局与int[2*3] 相同

标签： c++ cuda

【解决方案1】：

正如@jarod42 所指出的，对于您所展示的“自动”、“非可变长度”C 样式数组：

int values[2][3];

这种数组的存储格式等同于：

int values[2*3];

这意味着我们可以将该数组视为线性单下标数组（即使它不是）：

为了从主机传输到设备：

#define W 3
#define H 2
int values[H][W];
int *d_values;
cudaMalloc(&d_values, H*W*sizeof(int));
cudaMemcpy(d_values, values, H*W*sizeof(int), cudaMemcpyHostToDevice);

为了访问设备代码，使用“模拟”2D 访问：

__global__ void kernel(int *values, int width, ...){
  int col = threadIdx.x+blockDim.x*blockIdx.x;
  int row = threadIdx.y+blockDim.y*blockIdx.y;
  int my_value = values[row*width+col];
  ...
}  

int main(){
  ...
  kernel<<<...>>>(d_values, W, ...);
  ...
  }

但根据您问题中的措辞：

现在我知道 CUDA 接受线性形式的 2D 数组，但是如何传递已经构建的数组？

您似乎知道上述方法，我通常将其称为“展平”二维数组以以线性方式处理它（可能使用“模拟”二维访问）。

一般来说，处理一个宽度在编译时未知的二维数组，同时仍然允许设备代码中的双下标访问，是rather involved，我不推荐它，尤其是对于 CUDA 初学者。但实际上您所提出的情况并非如此：

内核的预定义二维数组。

int values[2][3];
              ^
             the "width"

我认为这是指数组的“宽度”（即第二个的范围，即最后一个下标）在编译时是已知的。在这种情况下，我们可以利用编译器为我们生成必要的数组索引，使传输和使用过程仅比“扁平化”情况稍微复杂一些，同时仍然允许内核中的双下标访问：

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

#define W 3
#define H 2
#define BSIZE 8

typedef int arrtype[W];

__global__ void kernel(arrtype *values, int width, int height){

  int col=threadIdx.x+blockDim.x*blockIdx.x;
  int row=threadIdx.y+blockDim.y*blockIdx.y;

  if ((row < height)&&(col < width)){

    int my_val = values[row][col]; //doubly-subscripted access
    printf("row: %d, col: %d, value: %d\n", row, col, my_val);
    }
}

int main(){

  int values[H][W];
  for (int i = 0; i < H; i++)
    for (int j = 0; j < W; j++)
      values[i][j] = i+j;
  arrtype *d_values;
  cudaMalloc(&d_values, H*W*sizeof(int));
  cudaMemcpy(d_values, values, H*W*sizeof(int), cudaMemcpyHostToDevice);
  dim3 block(BSIZE,BSIZE);
  dim3 grid((W+block.x-1)/block.x, (H+block.y-1)/block.y);
  kernel<<<grid,block>>>(d_values, W, H);
  cudaDeviceSynchronize();
  return 0;
}
$ nvcc -o t1023 t1023.cu
$ ./t1023
row: 0, col: 0, value: 0
row: 0, col: 1, value: 1
row: 0, col: 2, value: 2
row: 1, col: 0, value: 1
row: 1, col: 1, value: 2
row: 1, col: 2, value: 3
$

有关完整工作的 3D（即三重下标）示例，请参阅 here

【讨论】：