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

Question

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

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

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

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

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

Answer 1

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

另一件要考虑的事情是必须访问的内存量与要完成的计算量。在许多成像应用程序中，您不仅需要单个像素的值，还需要周围的像素。Cuda将线程分组到warps中，warps同时遍历每条指令(目前，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个线程。

Answer 2

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

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