了解 CUDA 网格尺寸、块尺寸和线程组织（简单解释）[关闭]

Question

如何组织线程以供 GPU 执行？

Answer 1

硬件

例如，如果一个 GPU 设备有 4 个多处理单元，每个单元可以运行 768 个线程：那么在给定时刻，真正并行运行的线程不会超过 4*768 个（如果您计划更多线程，它们将等待轮到他们）。

软件

线程以块的形式组织。一个块由一个多处理单元执行。 可以使用 1Dimension(x)、2Dimensions (x,y) 或 3Dim 索引 (x,y,z) 来识别（索引）块的线程，但在任何情况下 xyz

显然，如果您需要超过 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;

Answer 2

假设一个 9800GT GPU：

• 它有 14 个多处理器 (SM)
• 每个 SM 有 8 个线程处理器（AKA 流处理器、SP 或内核）
• 每个块最多允许 512 个线程
• warpsize 为 32（这意味着每个 14x8=112 线程处理器最多可以调度 32 个线程）

https://www.tutorialspoint.com/cuda/cuda_threads.htm

一个块不能有超过 512 个活动线程，因此__syncthreads 只能同步有限数量的线程。即如果您使用 600 个线程执行以下操作：

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

那么内核必须运行两次，执行顺序为：

1. func1 为前 512 个线程执行
2. func2 为前 512 个线程执行
3. func1 为剩余线程执行
4. func2 为剩余线程执行

注意：

重点是__syncthreads是块范围的操作，它不会同步所有线程。

---

我不确定__syncthreads 可以同步的确切线程数，因为您可以创建一个具有超过 512 个线程的块并让 warp 处理调度。据我了解，更准确的说法是：func1至少在前 512 个线程中执行。

在我编辑这个答案之前（早在 2010 年），我测量了 14x8x32 个线程是使用 __syncthreads 同步的。

如果有人再次对此进行测试以获得更准确的信息，我将不胜感激。