理解CUDA网格维度、块维度和线程组织（简单解释）

Question

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

- cibercitizen1

---

CUDA编程指南应该是一个很好的起点。我还建议从这里查看CUDA介绍。 - Tom

Answer 1

硬件

假设一个GPU设备有4个多处理器单元，并且它们可以每个运行768个线程：那么在某一时刻，最多只有4*768个线程会真正并行运行（如果您计划更多的线程，则它们将等待它们的轮到）。

软件

线程被组织在块中。一个块由一个多处理器单元执行。块的线程可以使用1维（x）、2维（x、y）或3维（x、y、z）索引来识别（编号），但无论如何，对于我们的示例，xyz<=768（对x、y、z还有其他限制，请参见指南和设备能力）。

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

现在考虑一个简单的情况：处理512x512像素的图像

假设我们想要一个线程来处理一个像素（i,j）。

我们可以使用每个块64个线程的块。然后我们需要4096个块（这样就有512x512个线程=4096x64）。

通常为了使图像索引更容易，将线程组织在2D块中，每个块的blockDim为8 x 8（每个块64个线程）。我更喜欢称其为threadsPerBlock。

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

2D的网格大小为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;

- cibercitizen1

---

如果每个区块可以运行768个线程，为什么只使用64个呢？如果使用最大限制的768个线程，您将拥有更少的区块，从而获得更好的性能。 - Aliza

@Aliza：块是“逻辑上的”，768线程的限制是针对每个“物理处理单元”的。您根据问题的规格使用块，以将工作分配给线程。您可能不能总是针对每个问题使用768线程的块。想象一下，您需要处理一个64x64像素（4096像素）的图像。4096/768 = 5.333333块？ - cibercitizen1

@cibercitizen1 - 我认为Aliza的观点很好：如果可能的话，一个块应该使用尽可能多的线程。如果有约束需要使用更少的线程，最好在第二个示例中解释为什么会出现这种情况（但仍然要首先解释更简单和更理想的情况）。 - user227667

@thouis 是的，也许是这样。但问题在于每个线程需要的内存量取决于应用程序。例如，在我的最后一个程序中，每个线程都调用了一个最小二乘优化函数，需要“很多”内存。所需的内存如此之多，以至于块的大小不能超过4x4个线程。即使如此，与顺序版本相比，获得的加速比是惊人的。 - cibercitizen1

@MySchizoBuddy 你应该填充图像，添加像素（或修剪它，删除像素），使其适合2的幂维度。 - cibercitizen1

Answer 2

假设有一块9800GT显卡：

• 它有14个多处理器（SM）
• 每个SM有8个线程处理器（也称为流处理器、SP或核心）
• 每个块最多允许512个线程
• Warp大小为32（这意味着14x8=112个线程处理器中的每个处理器可以调度最多32个线程）

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

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

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

如果内核必须运行两次，执行顺序如下：

1. 首先为前512个线程执行func1
2. 然后为前512个线程执行func2
3. 接着为剩余的线程执行func1
4. 最后为剩余的线程执行func2

注意：

__syncthreads是块级操作，它不会同步所有线程。

---

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

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

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

- Bizhan

---

如果func2()依赖于func1()的结果，会发生什么？我认为这是错误的。 - Chris

抱歉，我认为这是错误的。此外，该GPU只能同时运行112个线程。 - Steven Lu

@StevenLu 请阅读第40页中间段落：http://developer.download.nvidia.com/compute/cuda/3_1/toolkit/docs/NVIDIA_CUDA_C_BestPracticesGuide_3.1.pdf - Bizhan

......因为整个10年前的GPU上只有N=112个FP32执行单元（CUDA核心）！ 我试图指出，您正在使用这些非常具体的数字，例如可以分配给SM的最大线程数，意味着同时运行那么多个warp，但事实并非如此，尽管存在某种思考方式可以扭曲这些话的含义。 实际上，对于可能被分配给给定SM的所有24个warp，它们都将逐渐被处理，直到完成才能被驱逐。 - Steven Lu

@StevenLu 最大线程数不是问题，__syncthreads 是一个块级操作，它实际上并没有同步所有线程，这对于 CUDA 学习者来说是一个麻烦。所以我根据你给我的信息更新了我的答案。非常感谢。 - Bizhan