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

Question

我了解了CUDA中块和网格的概念，想知道是否有写得很好的帮助函数可以帮助我确定任何给定2D图像的最佳块和网格大小。例如，在this thread中提到的512x512图像中，网格为64x64，块为8x8。然而，有时我的输入图像可能不是2的幂，可能是317x217或类似的尺寸。在这种情况下，网格应该是317x1，块应该是1x217。因此，如果我有一个接受用户图像并使用CUDA处理它的应用程序，如何自动确定块和网格的大小和维度，其中用户可以输入任何大小的图像。是否存在任何已经存在的帮助函数或类来处理这个问题？

- Philip Zhang

Answer 1

通常，您应该根据GPU架构选择块的大小，以保持流处理器（SM）的100％占用率。例如，我学校的GPU可以每个SM运行1536个线程，每个SM最多可以有8个块，但每个块在每个维度上只能有1024个线程。因此，如果我在GPU上启动一个1d内核，我可以使用1024个线程使块达到最大值，但然后只有1个块会在SM上（占用率为66％）。如果我选择较小的数字，如192个线程或256个线程每个块，则我可以在SM上具有6个和8个块的100％占用率。

另一件要考虑的事情是必须访问的内存量与要执行的计算量。在许多图像应用程序中，您不仅需要单个像素的值，而且还需要周围的像素。 Cuda将其线程分组成warp，这些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 个线程每个块。

- user3288829

Answer 2

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

所以没有像你描述的函数那样。线程和块数取决于你如何处理数据，它不是普遍适用的。

- Kill Console