假设您要编写一个内核，该内核对大小为 400x900 像素的图像进行操作。您还想为每个像素分配一个 GPU 线程。您的线程块是方形的，您希望在设备上使用每个块可能的最大线程数。每个块的最大线程数为 1024。如何选择内核的网格尺寸和块尺寸？
我对其工作原理的理解是，为每个像素分配一个线程，我需要 360,000 (400x900) 个线程。数据层次结构是网格 -> 块 -> 线程。我认为该公式最终将是 360,000 = (# of blocks)*(# of threads per block)，其中# of blocks 必须是完美的平方数和 32 的倍数。
我尝试了从 2 到 4096 的数字，但在除以 360,000 时，它们都没有给我一个偶商。这是否意味着线程可以是十进制数？
【问题讨论】：
标签： cuda
使用 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）。
我们可以看到，每个维度的线程总数都小于对应的图像维度。我们需要做的是四舍五入块的数量，以便如果图像尺寸不是块尺寸的倍数，我们创建一个额外的块来覆盖剩余的像素。 以下是两种方法。
我们不使用整数除法来计算块数，而是使用浮点除法和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
}
}
像这样创建网格和块：
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
}
}
【讨论】：
通常，您将尺寸设置为所需尺寸的下一个倍数，然后在内核中进行边界检查。
这里有一个简单的例子： 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];
}
【讨论】：