Hardware
例如,如果 GPU 设备有 4 个多处理单元,并且每个单元可以运行 768 个线程:那么在给定时刻,真正并行运行的线程不会超过 4*768 个(如果您计划了更多线程,它们将等待轮到他们了)。
Software
线程以块的形式组织。块由多处理单元执行。
块的线程可以使用 1Dimensions(x)、2Dimensions (x,y) 或 3Dim 索引 (x,y,z) 来标识(索引),但在任何情况下 xy对于我们的示例,z
显然,如果您需要 4*768 以上的线程,则需要 4 个以上的块。
块也可以按 1D、2D 或 3D 索引。有一个块队列等待进入
GPU(因为在我们的示例中,GPU 有 4 个多处理器,并且只有 4 个块
同时执行)。
现在是一个简单的案例:处理 512x512 图像
假设我们希望一个线程处理一个像素 (i,j)。
我们可以使用每个 64 个线程的块。那么我们需要 512*512/64 = 4096 个块
(所以有 512x512 线程 = 4096*64)
通常将线程组织在 blockDim = 8 x 8(每个块 64 个线程)的 2D 块中(以便更轻松地索引图像)。我更喜欢将其称为“threadsPerBlock”。
dim3 threadsPerBlock(8, 8); // 64 threads
2D gridDim = 64 x 64 块(需要 4096 个块)。我更喜欢称之为 numBlocks。
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;