启动 GPU 内核以对 3D 数据集进行计算的最佳方法是什么？答案

【问题标题】：What is the best way to launch a GPU kernel to do calculation on a 3D data set?启动 GPU 内核以对 3D 数据集进行计算的最佳方法是什么？
【发布时间】：2012-06-12 02:05:51
【问题描述】：

我正在使用 CUDA 对可能很大的 3D 数据集进行计算。我觉得最好先看个短代码sn-p：

void launch_kernel(/*arguments . . . */){
    int bx = xend-xstart, by = yend-ystart, bz = zend-zstart;

    dim3 blocks(/*dimensions*/);
    dim3 threads(/*dimensions*/);
    kernel<<blocks, threads>>();
}

我有一组 3D 单元，我需要启动一个内核来计算每个单元。问题是输入大小可能超出 GPU 的能力，特别是线程。所以代码如下：

void launch_kernel(/*arguments . . . */){
       int bx = xend-xstart, by = yend-ystart, bz = zend-zstart;

       dim3 blocks(bx,by,1);
       dim3 threads(bz);
       kernel<<blocks, threads>>();
   }

... 效果不佳。因为如果尺寸是 1000x1000x1000 怎么办？ - 我无法在每个块中启动 1000 个线程。或者更好的是，如果尺寸是 5x5x1000 怎么办？ - 现在我几乎不启动任何块，但内核需要启动 5x5x512 b/c 的硬件，每个线程将执行 2 次计算。我也不能只是混搭我的所有维度，将一些 z 放在块中，一些放在线程 b/c 中，我需要能够识别内核中的维度。目前：

__global__ void kernel(/*arguments*/){
    int x = xstart + blockIdx.x;
    int y = ystart + blockIdx.y;
    int z = zstart + threadIdx.x;
    if(x < xend && y < yend && z < zend){
        //calculate
    }
}

我需要一种可靠、有效的方法来计算这些变量：

块 x 维度，块 y 维度，线程 x（以及 y? 和 z?），x,y,z 一旦我通过 blockIdx 和 threadIdx 在内核中，如果输入超过硬件，则我在内核计算中的 for 循环中为每个维度采取的“步骤”量。

如果您有任何问题，请提出。这是一个难题，一直困扰着我（尤其是因为我启动的块/线程的数量是性能的主要组成部分）。该代码需要在其针对不同数据集的决策中自动化，我不确定如何有效地做到这一点。提前谢谢你。

【问题讨论】：

您使用的是什么 GPU？如果是 Fermi 或 Kepler 卡（计算能力为 2.x 或 3.x），则硬件中支持 3D 网格，这大大简化了事情。
没有 3D 网格。它需要能够在相当新的 NVIDIA 显卡上运行（我会说是在过去 4 年内发布的）。

标签： c multidimensional-array cuda gpu

【解决方案1】：

我认为你在这里把事情复杂化了。基本问题似乎是您需要在 1000 x 1000 x 1000 计算域上运行内核。因此，您需要 1000000000 个线程，这完全在所有 CUDA 兼容硬件的能力范围内。因此，只需使用标准的 2D CUDA 执行网格，其中至少包含进行计算所需的线程数（如果您不明白如何做到这一点，请发表评论，我会将其添加到答案中），然后在您的内核调用中一个小的设置函数是这样的：

__device__ dim3 thread3d(const int dimx, const int dimxy)
{
    // The dimensions of the logical computational domain are (dimx,dimy,dimz)
    // and dimxy = dimx * dimy
    int tidx = threadIdx.x + blockIdx.x * blockDim.x;
    int tidy = threadIdx.y + blockIdx.y * blockDim.y;
    int tidxy = tidx + gridDim.x * tidy;

    dim3 id3d;
    id3d.z = tidxy / dimxy;
    id3d.y = tidxy / (id3d.z * dimxy);
    id3d.x = tidxy - (id3d.z * dimxy - id3d.y * dimx);

    return id3d;
}

[免责声明：在浏览器中编写，从未编译，从未运行，从未测试。使用风险自负]。

此函数将从 CUDA 2D 执行网格返回 3D 域 (dimx,dimy,dimz) 中的“逻辑”线程坐标。在内核的开头这样称呼它：

__global__ void kernel(arglist, const int dimx, const int dimxy)
{
    dim3 tid = thread3d(dimx, dimxy);

    // tid.{xyx} now contain unique 3D coordinates on the (dimx,dimy,dimz) domain
    .....
}

请注意，设置网格需要大量整数计算开销，因此您可能需要考虑为什么您真的需要 3D 网格。您会惊讶于它实际上没有必要的次数，并且可以避免大部分设置开销。

【讨论】：

很好的答案！正如您所提供的，我希望看到内核的启动，因为我对 dimx 和 dimxy 的起源仍然有些困惑（参数列表中的 dimxy 是错字吗？）。我需要尽可能快的计算，我注意到我启动的线程/块的数量似乎花费了最多的时间，而不是计算。计算部分可能有大约 25-40 FLOP，但它们处理 3D 立方体中与其他单元格相邻的单元格，因此我需要 XYZ 坐标。感谢您的回答。
在上一篇文章中，我提到块的数量似乎提供了很多开销。在运行了更多测试之后，我不确定这是否是真的。它实际上似乎可以很好地扩展，开销方面，只是额外的块需要计算才能发现它们实际上是额外的。但是，正如我认为您的代码所做的那样，让 Z 不再完全依赖线程，这正是我想要的。

【解决方案2】：

我会首先使用cudaGetDeviceProperties() 来查找您的 GPU 的计算能力，这样您就可以确切地知道您的 GPU 允许每个块有多少线程（如果您的程序需要被泛化以便它可以在任何支持 CUDA 的设备上运行） )。

然后，使用该数字，我将创建一个大的嵌套 if 语句来测试您输入的维度。如果所有维度都足够小，您可以拥有一组 (bx,by,bz) 线程（不太可能）。如果这不起作用，则找到可以放入一个块中的最大维度（或两个维度）并据此进行分区。如果这不起作用，那么您必须对最小维度进行分区，以便将其中的一些块放入一个块中 - 例如 (MAX_NUMBER_THREADS_PER_BLOCK,1,1) 线程和 (bx/MAX_NUMBER)THREADS_PER_BLOCK,by,bz) 块，假设 bx<by<bz 和 bx>MAX_NUMBER_THREADS_PER_BLOCK。

每种情况都需要不同的内核，这有点麻烦，但归根结底这是一项可行的工作。

【讨论】：