如何在 CUDA 中自动计算 2D 图像的块和网格大小？

Question

我已经知道 cuda 中块和网格的想法，我想知道是否有任何编写良好的辅助函数可以帮助我确定任何给定 2D 图像的最佳块和网格大小。

例如，对于此线程中提到的 512x512 图像。网格为 64x64，块为 8x8。

然而有时我的输入图像可能不是 2 的幂，它可能是 317x217 或类似的东西。在这种情况下，也许网格应该是 317x1 而块应该是 1x217。

因此，如果我有一个应用程序接受来自用户的图像，并使用 cuda 对其进行处理，它如何自动确定块和网格的大小和尺寸，用户可以在其中输入任何大小的图像。

是否有任何现有的帮助函数或类来处理这个问题？

Answer 1

通常，您希望根据您的 GPU 架构选择块的大小，目标是在流式多处理器 (SM) 上保持 100% 的占用率。比如我学校的GPU每个SM可以运行1536个线程，每个SM最多可以运行8个block，但是每个block每个维度最多只能有1024个线程。因此，如果我要在 GPU 上启动一个 1d 内核，我可以将一个具有 1024 个线程的块最大化，但随后只有 1 个块位于 SM 上（66% 的占用率）。如果我改为选择较小的数字，例如每块 192 个线程或 256 个线程，那么我可以在 SM 上分别使用 6 个和 8 个块获得 100% 的占用率。

要考虑的另一件事是必须访问的内存量与要完成的计算量。在许多成像应用中，您不仅需要单个像素的值，还需要周围的像素。Cuda 将其线程分组为 warp，它们同时执行每条指令（目前，warp 有 32 个线程，尽管这可能会改变）。使块方形通常可以最大限度地减少需要加载的内存量与可以完成的计算量，从而提高 GPU 的效率。同样，由于 Cuda 一次加载内存行而不是单个值，因此 2 的幂的块会更有效地加载内存（如果与内存地址正确对齐）。

因此，对于您的示例，即使拥有 317x1 的网格和 1x217 的块似乎更有效，但如果您在 20x14 的网格上启动 16x16 的块，您的代码可能会更高效，因为它会导致更好的计算/内存比率和 SM 占用率。但是，这确实意味着您必须在内核中进行检查以确保线程在尝试访问内存之前没有超出图片，例如

const int thread_id_x = blockIdx.x*blockDim.x+threadIdx.x;
const int thread_id_y = blockIdx.y*blockDim.y+threadIdx.y;
if(thread_id_x < pic_width && thread_id_y < pic_height)
{
  //Do stuff
}

最后，您可以使用 (N+M-1)/M 确定每个网格维度中所需的最少块数，其中 N 是该维度中的总线程数，并且每个块有 M 个线程那个维度。