cuda中的动态3D数组参数

Question

我试图在cuda内核函数中使用动态3D数组参数，但我做得不好。

__global__ void kernel ( 3D array pointer )
{
// do something
}
int main()
{
    const int NUM_OF_ARRAY;
    const int ROW;
    const int CAL;
    int arr[NUM_OF_ARRAY][ROW][CAL]; 
    // Maybe I should use cudaMalloc3D or cudaMalloc3DArray
    dim3 grid( , , ,);
    dim3 block( , , , );
    kernel <<< grid, block >>> ( ? );
 }

我看到罗伯特的答案发送3d数组CUDA内核，但我认为我的情况有点不同。

如果数组的行和调用是在运行时确定的，我怎么能在cuda中分配内存并给出内核函数的指针?

我尝试使用cudaMalloc3D或cudaMalloc3DArray，但我不能很好，因为我从来没有使用过。

谁能展示使用动态3D数组参数的简单示例?

这对我很有帮助。谢谢。

Answer 1

由于前面链接的答案和其他地方建议的所有原因，这不一定是处理3D数组的好方法。更好的方法是将数组平面化，并使用指针算法来模拟3D访问。

但是为了证明前面的示例并不需要硬编码维度，这里修改了这个示例，以显示可变(运行时)维度的使用:

#include <iostream>
inline void GPUassert(cudaError_t code, char * file, int line, bool Abort=true)
{
    if (code != 0) {
        fprintf(stderr, "GPUassert: %s %s %dn", cudaGetErrorString(code),file,line);
        if (Abort) exit(code);
    }
}
#define GPUerrchk(ans) { GPUassert((ans), __FILE__, __LINE__); }

  __global__ void doSmth(int*** a, int sz_x, int sz_y, int sz_z) {
    for(int i=0; i<sz_z; i++)
     for(int j=0; j<sz_y; j++)
      for(int k=0; k<sz_x; k++)
       a[i][j][k]=i-j+k;
 }
 int main() {
  unsigned sx;
  unsigned sy;
  unsigned sz;
  std::cout << std::endl << "Enter x dimension (3rd subscript): " ;
  std::cin >> sx;
  std::cout << std::endl << "Enter y dimension (2nd subscript): " ;
  std::cin >> sy;
  std::cout << std::endl << "Enter z dimension (1st subscript): " ;
  std::cin >> sz;
  int*** h_c = (int***) malloc(sz*sizeof(int**));
  for(int i=0; i<sz; i++) {
   h_c[i] = (int**) malloc(sy*sizeof(int*));
   for(int j=0; j<sy; j++)
    GPUerrchk(cudaMalloc((void**)&h_c[i][j],sx*sizeof(int)));
  }
  int ***h_c1 = (int ***) malloc(sz*sizeof(int **));
  for (int i=0; i<sz; i++){
    GPUerrchk(cudaMalloc((void***)&(h_c1[i]), sy*sizeof(int*)));
    GPUerrchk(cudaMemcpy(h_c1[i], h_c[i], sy*sizeof(int*), cudaMemcpyHostToDevice));
    }
  int*** d_c;
  GPUerrchk(cudaMalloc((void****)&d_c,sz*sizeof(int**)));
  GPUerrchk(cudaMemcpy(d_c,h_c1,sz*sizeof(int**),cudaMemcpyHostToDevice));
  doSmth<<<1,1>>>(d_c, sx, sy, sz);
  GPUerrchk(cudaPeekAtLastError());
  int res[sz][sy][sx];
  for(int i=0; i<sz; i++)
   for(int j=0; j<sy; j++)
    GPUerrchk(cudaMemcpy(&res[i][j][0], h_c[i][j],sx*sizeof(int),cudaMemcpyDeviceToHost));
  std::cout << std::endl;
  for(int i=0; i<sz; i++)
   for(int j=0; j<sy; j++)
    for(int k=0; k<sx; k++)
     printf("[%d][%d][%d]=%dn",i,j,k,res[i][j][k]);
 }

我已经将内核存储的数据修改为i-j+k而不是i+j+k。此外，我还为下标创建了一个[z][y][x]顺序，因为这将建议使用线程索引计算安排，如[threadIdx.z][threadIdx.y][threadIdx.x]，这将最有利于合并访问。然而，这种类型的多重下标数组在内核中仍然是低效的，因为在解析数据的最终位置时需要跟踪指针。