理解 CUDA 网格尺寸、块尺寸和线程组织(简单说明)

小开

最佳答案

例如，如果一个 GPU 设备有4个多处理单元，并且每个单元可以运行768个线程: 那么在给定的时刻，实际并行运行的线程不会超过4 * 768个(如果你计划更多的线程，它们会等待轮到它们运行)。

线程以块的形式组织。块由多处理单元执行。一个块的线程可以使用1维(x) ，2维(x，y)或3Dim 索引(x，y，z)来标识(索引) ，但是在任何情况下，我们的例子中 x嘿z < = 768(其他限制适用于 x，y，z，参见指南和设备功能)。

显然，如果需要的线程超过4 * 768个，那么需要的块就超过4个。块也可以索引1D，2D 或3D。有一个等待进入的块队列 GPU (因为在我们的示例中，GPU 有4个多处理器，只有4个块是同时执行)。

假设我们希望一个线程处理一个像素(i，j)。

我们可以使用每个64个线程的块，然后我们需要512 * 512/64 = 4096个块 (512x512线程 = 4096 * 64)

组织2D 块中的线程(使图像的索引更容易)是很常见的，块中的线程数为 lockDim = 8 x 8(每个块中有64个线程)。我更喜欢称之为 threadsPerBlock。

dim3 threadsPerBlock(8, 8);  // 64 threads

和2D gridDim = 64 x 64块(需要4096块)。

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
imageHeight/threadsPerBlock.y);

内核是这样启动的:

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );

最后: 将会出现类似于“4096个块的队列”的情况，其中一个块正在等待被分配给 GPU 的一个多处理器来执行它的64个线程。

在内核中，线程处理的像素(i，j)是这样计算的:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;

假设一个9800GT 图形处理器:

块中的活动线程不能多于512个，因此 __syncthreads只能同步有限数量的线程。也就是说。如果使用600个线程执行以下操作:

func1();
__syncthreads();
func2();
__syncthreads();

那么内核必须运行两次，执行顺序是:

注:

主要的一点是 __syncthreads是一个块范围的操作，它不同步所有的线程。

我不确定 __syncthreads可以同步的线程的确切数量，因为您可以创建一个包含超过512个线程的块，并让翘曲处理调度。根据我的理解，更准确的说法是: fun1对前512个线程执行至少。

在我编辑这个答案之前(回到2010年) ，我测量了使用 __syncthreads同步14x8x32个线程。

如果有人能再次测试一下，以获得更准确的信息，我将不胜感激。