计算内核的网格和块尺寸

Question

假设您要编写一个对大小为 400x900 像素的图像进行操作的内核。您还希望为每个像素分配一个 GPU 线程。您的线程块是方形的，并且您希望在设备上使用每个块可能的最大线程数。每个块的最大线程数是 1024。您将如何选择内核的网格尺寸和块尺寸？

我对其工作原理的理解是，将一个线程分配给每个像素，我需要 360,000 (400x900) 个线程。数据层次结构为网格 -> 块 -> 线程。我认为公式最终会是 360,000 = (块数)*(每个块的线程数)，块数必须是完全平方数且是 32 的倍数。

我尝试过 2 到 4096 之间的数字，但除 360,000 时没有一个给出偶数商。这是否意味着线程可以是十进制数？

Answer 1

当使用 CUDA 处理 2D 图像时，自然的直觉是使用 2D 块和网格形状。如果我们想设置最大可能的块大小，我们必须确保其尺寸的乘积不超过块大小限制。请记住块大小的限制 (1024)，以下是有效块大小的一些示例。

dim3 block(32,32); //32 x 32 = 1024
or
dim3 block(64,16); //64 x 16 = 1024
or
dim3 block(16,64); //16 x 64 = 1024 ... Duh

接下来是二维网格尺寸的计算。如果我们想为每个像素映射一个线程，那么应该创建网格，使得每个维度中的线程总数至少等于相应的图像维度。请记住，网格大小意味着每个维度中的块数。这意味着某个维度中的线程总数等于该维度中的网格大小和块大小的乘积。对于 2D 网格，X 维度上的线程数等于block.x * grid.x，Y 维度上的线程数等于block.y * grid.y。

假设您有一个尺寸为 400 x 900 的图像，那么相应维度中的线程总数也应该至少相同。

假设您选择了一个大小为 (32,32) 的块。那么图像的 x 和 y 维度的块数应为400/32和900/32。但是两个图像尺寸都不是相应块尺寸的整数倍，因此由于整数除法，我们最终将创建大小为12 x 28的网格，这将导致线程总数等于384 x 896。（因为 32 x 12 = 384 和 32 x 28 = 896）。

我们可以看到，每个维度的线程总数都小于相应的图像维度。我们需要做的是对块的数量进行四舍五入，以便如果图像尺寸不是块尺寸的倍数，我们会创建一个额外的块来覆盖剩余的像素。以下是执行此操作的 2 种方法。

我们不使用整数除法来计算块数，而是使用浮点除法和ceil结果。

int image_width = 400;
int image_height = 900;
dim3 block(32,32);
dim3 grid;
grid.x = ceil( float(image_width)/block.x );
grid.y = ceil( float(image_height)/block.y );

另一种聪明的方法是使用以下公式

int image_width = 400;
int image_height = 900;
dim3 block(32,32);
dim3 grid;
grid.x = (image_width + block.x - 1 )/block.x;
grid.y = (image_height + block.y - 1 )/block.y;

当以上述方式创建网格时，您最终将创建一个大小为13 x 29的网格，这将导致线程总数等于416 x 928。

现在在这种情况下，我们每个维度的线程总数都大于相应的图像维度。这将导致某些线程访问图像边界之外的内存，从而导致未定义的行为。这个问题的解决方案是我们在内核内部执行边界检查，并仅对落在图像边界内的线程进行处理。当然，要做到这一点，我们需要将图像尺寸作为参数传递给内核。以下示例内核显示了此过程。

__global__ void kernel(unsigned char* image, int width, int height)
{
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x; //image x index or column number
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y; //image y index of row number

    if(xIndex < width && yIndex < height)
    {
       //Do processing only here
    }
}

总长DR

像这样创建网格和块：

dim3 block(32,32);
dim3 grid;
grid.x = (image_width + block.x - 1)/block.x;
grid.y = (image_height + block.y - 1)/block.y;

调用内核并将图像尺寸作为参数传递，如下所示：

kernel<<<grid, block>>>(...., image_width, image_height);

像这样在内核内部执行绑定检查：

__global__ void kernel(unsigned char* image, int width, int height)
{
    int xIndex = blockIdx.x * blockDim.x + threadIdx.x; //image x index or column number
    int yIndex = blockIdx.y * blockDim.y + threadIdx.y; //image y index of row number

    if(xIndex < width && yIndex < height)
    {
       //Do processing only here
    }
}