计算内核的网格和块尺寸

Posted 2021-04-21

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

我对它如何工作的理解是将一个线程归因于每个像素，我需要360,000（400x900）个线程。数据层次结构为grid - > block - > threads。我认为公式最终会是360,000 =（块数）*（每块的线程数），块数必须是一个完美的平方数和32的倍数。

我已经尝试过从2到4096的数字，当从360,000分开时，它们都没有给我一个平均商数。这是否意味着线程可以是十进制数？

答案

通常，您将维度设置为所需大小的下一个倍数，然后在内核中执行绑定检查。

这里有一个简单的例子：https://devblogs.nvidia.com/parallelforall/easy-introduction-cuda-c-and-c/

这里计算块的数量，因此线程总数等于或高达所需线程数的+256。

  saxpy<<<(N+255)/256, 256>>>(N, 2.0f, d_x, d_y);

在内核中，仅在需要时才执行计算：

__global__
void saxpy(int n, float a, float *x, float *y)
{
  int i = blockIdx.x*blockDim.x + threadIdx.x;
  if (i < n) y[i] = a*x[i] + y[i];
}

另一答案

使用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网格尺寸的计算。如果我们想要为每个像素映射一个线程，那么应该创建网格，使得每个维度中的线程总数至少等于相应的图像维度。请记住，网格大小表示每个维度中的块数。这意味着维度中的线程总数等于该维度中网格大小和块大小的乘积。对于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×28 = 896）。

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

我们使用浮点除法和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
    }
}

TLDR

像这样创建网格和块：

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
    }
}