如何计算每个网格的块?

首先定义栅格尺寸。这取决于你在做什么。假设您有一张尺寸为1024x768的图像,然后进行逐像素计算。那么你的网格自然会 dim3(1024, 768, 1) 。如果你有几帧视频,它可能是 dim3(width, height, frames) 。

您的块尺寸由内核决定。如果您没有特定的需要,我建议您不要使用非常大的块,如1024。这可能导致低效利用,因为诸如 __syncthreads() 或者内核的结束导致相当多的线程等待几个掉队者。当有疑问时,坚持使用128x1x1或16x16x1之类的小东西。

考虑到这一点,您可以计算块。

dim3 grid(1024, 768, 1);
dim3 blockdim(16, 16, 1);
dim3 blocks((grid.x + blockdim.x - 1) / blockdim.x,
            (grid.y + blockdim.y - 1) / blockdim.y,
            (grid.z + blockdim.z - 1) / blockdim.z);

这个计算是四舍五入的除法。这样可以确保在网格尺寸不能被块尺寸整除时启动足够的块。不利的一面是,现在您可能会启动比所需更多的线程。实际上,你的网格被填充到块尺寸的倍数。有多种方法可以解决这一问题,例如确保阵列维度始终是16的倍数。然而,最简单的方法是简单地检查是否超出范围。

__global__ void kernel(int xdim, int ydim, int zdim)
{
    int x_idx = blockIdx.x * blockDim.x + threadIdx.x;
    int y_idx = blockIdx.y * blockDim.y + threadIdx.y;
    int z_idx = blockIdx.z * blockDim.z + threadIdx.z;
    if(x_idx < xdim && y_idx < ydim && z_idx < zdim) {
        do_something();
    }
}

通常,如果沿Y和Z轴的块尺寸为1,则可以跳过Y和Z尺寸的检查。

另一种方法是只启动GPU上可以同时运行的线程,然后在内核内进行循环。与大型网格的实际工作相比,这具有降低启动开销比例的好处,因为所有索引计算只需要完成一次,并且通过将部分计算从循环体中提取出来,可以进行更多优化。

__global__ void kernel(int xdim, int ydim, int zdim)
{
    int x_start = blockIdx.x * blockDim.x + threadIdx.x;
    int y_start = blockIdx.y * blockDim.y + threadIdx.y;
    int z_start = blockIdx.z * blockDim.z + threadIdx.z;
    int x_stride = gridDim.x * blockDim.x;
    int y_stride = gridDim.y * blockDim.y;
    int z_stride = gridDim.z * blockDim.z;
    for(int z_idx = z_start; z_idx < zdim; z_idx += z_stride)
        for(int y_idx = y_start; y_idx < ydim; y_idx += y_stride)
            for(int x_idx = x_start; x_idx < xdim; x_idx += x_stride)
                do_something();
}