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

Posted 2023-04-15

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

【英文标题】：How to automatically calculate the block and grid size of a 2D image in CUDA?

【发布时间】：2014-07-24 13:56:26

【问题描述】：

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

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

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

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

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

【参考方案1】：

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

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

因此，对于您的示例，即使拥有 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 个线程该维度中的每个块。

【参考方案2】：

这取决于您如何处理图像。如果您的线程仅单独处理每个像素，例如，将每个像素值加 3，则您可以将一个维度分配给您的块大小，将另一个维度分配给您的网格大小（只是不要超出范围）。但是如果要进行过滤或腐蚀之类的操作，这种操作通常需要访问中心像素附近的像素，例如 9*9 的 3*3。然后块应该是你提到的 8*8 或其他值。而且你最好使用纹理内存。因为当线程访问全局内存时，总是有32个线程被一个块包裹一次。

所以没有你描述的功能。线程和块的数量取决于你如何处理数据，它不是通用的。