在 cuda 中将 3D 数组展平为 1D答案

【问题标题】：Flattening a 3D array to 1D in cuda在 cuda 中将 3D 数组展平为 1D
【发布时间】：2017-10-08 09:53:36
【问题描述】：

我尝试在 cuda 中实现以下代码，但在 cuda 中将 3D 数组展平为 1D 时遇到问题

C++ 代码

for(int i=0; i<w; i++)
  for(int j=0; j<h; j++)
    for(int k=0; k<d; k++)
     arr[h*w*i+ w*j+ k] = (h*w*i+ w*j+ k)*2;

这就是我目前在 Cuda 中所拥有的

  int w = h = d;
  int N = 64;

 __global__ void getIndex(float* A)
{
  int i = blockIdx.x;
  int j = blockIdx.y;
  int k = blockIdx.z;
  A[h*w*i+ w*j+ k] = h*w*i+ w*j+ k;
}


int main(int argc, char **argv)
 {

    float *d_A;
    cudaMalloc((void **)&d_A, w * h * d * sizeof(float) );
    getIndex <<<N,1>>> (d_A);
  }

但我没有得到我期望的结果，我不知道如何获得正确的 i,j 和 k 索引

【问题讨论】：

你应该看看 CUDA 指南，尤其是Chapter 2.2。您可以“在三个维度上调用 CUDA 内核”。我认为这是你想做的，但你没有。还要考虑到每个块绝对应该启动多个线程。

标签： c++ arrays indexing 3d cuda

【解决方案1】：

考虑大小为w x h x d 的3D 问题。（这可能是一个简单的数组，必须像您的问题或任何其他易于并行化的 3D 问题一样设置。）我将使用您的简单设置任务进行演示。

使用 CUDA 内核处理此问题的最简单方法是为每个数组条目启动一个线程，即w*h*d 线程。 This answer 讨论了为什么每个元素一个线程可能并不总是最好的解决方案。

现在让我们看看下面几行代码

dim3 numThreads(w,h,d);
getIndex <<<1, numThreads>>> (d_A, w, h, d);

在这里，我们启动了一个总共有w*h*d 线程的内核。内核可以实现为

__global__ void getIndex(float* A, int w, int h, int d) // we actually do not need w
{
    int i = threadIdx.x;
    int j = threadIdx.y;
    int k = threadIdx.z;
    A[h*d*i+ d*j+ k] = h*d*i+ d*j+ k;
}

但是这个内核和内核调用有一个问题：每个线程块的线程数是有限的（也是“特定方向的线程”的数量是有界的=z方向通常是最有界的）。由于我们只调用一个线程块，我们的问题大小不能超过这些特定限制（例如w*h*d <= 1024）。

这就是线程块的用途。实际上，您可以根据需要启动具有任意多线程的内核。（这不是真的，但是线程块的最大数量的限制不太可能用尽。）

这样调用内核：

dim3 numBlocks(w/8,h/8,d/8);
dim3 numThreads(8,8,8);
getIndex <<<numBlocks, numThreads>>> (d_A, w, h, d);

将为w/8 * h/8 * d/8 线程块启动内核，而每个块都包含8*8*8 线程。所以总共会调用w*h*d 个线程。现在我们必须相应地调整我们的内核：

__global__ void getIndex(float* A, int w, int h, int d) // we actually do not need w
{
    int bx = blockIdx.x;
    int by = blockIdx.y;
    int bz = blockIdx.z;
    int tx = threadIdx.x;
    int ty = threadIdx.y;
    int tz = threadIdx.z;
    A[h*d*(8*bx + tx)+ d*(8*by + ty)+ (8*bz + tz)] = h*d*(8*bx + tx)+ d*(8*by + ty)+ (8*bz + tz);
}

注意：

您可以使用blockDim.x 而不是固定大小的8 和gridDim.x 编写更通用的内核，以通过gridDim.x*blockDim.x 计算w。其他两个维度也同样处理。
在建议的示例中，所有三个维度 w、h 和 d 必须是 8 的倍数。您还可以概括内核以允许每个维度。（然后您必须将所有三个维度解析到内核，以检查计算出的位置是否仍在问题范围内。）
如前所述，每个线程编辑多个数组条目可能更有效。调用内核时必须再次考虑这一点。一个包装函数，它接收问题大小和数据，并使用正确的块和线程配置调用内核。

【讨论】：