CUDA 中线程层次结构的示例用例

Question

我正在学习 CUDA，并了解到线程层次结构的要点是允许具有不同功能的 GPU 之间的可扩展性和兼容性。 （我不知道这是不是唯一的好处）。

在 NVIDIA 的文档（第 1.3 节）中，有一张小图片突出显示了 8 块“多线程”CUDA 程序如何在具有 2 个 SM 的 GPU 和具有 4 个 SM 的 GPU 上运行。但是，我无法理解为什么不能仅通过线程而不是网格和块抽象来实现这一点。

请有人提供线程层次结构的示例用例，其中只有线程是不够的？

Answer 1

线程层次结构，就像内存层次结构一样，都与引用的局部性有关。

扭曲
线程以 32 为一组进行分组：经纱。扭曲中的线程可以使用

shfl_\[up_/down_/xor_/_\]sync

、扭曲减少、扭曲匹配和（非常有用）扭曲投票指令来共享数据。

经纱中的线程具有廉价的同步：

__syncwarp

螺纹块
下一步是线程块，最多 32 个线程束（1024 个线程）在同一个多处理器 (SM) 中协作。块内通信没有专门的指令，但是块内的线程可以使用

__shared__

内存来交换数据。这几乎和

__shfl_sync

一样有效。

网格块
接下来我们有多处理器 (SM)，一个 GPU 可以有 20 或 100 个。这些（在布莱克威尔之前）只能通过全局内存进行通信。 SM

a

中的线程将数据存储在全局内存中，SM

b

中的线程读取该数据。 两个网格块通常在不同的 SM 上运行（如果它们碰巧在同一个 SM 上，则它们不能共享

__shared__

内存）。 （在 Blackwell 中，SM

a

中的线程可以将数据从其共享内存移动到另一个 SM 中的共享内存）。

不同的GPU
层次结构的下一个级别是在同一台计算机上的多个 GPU 上运行您的程序。 它们只能通过 GPU 互连进行通信（我忘记了它的名字）。
这显然更慢、更复杂。

不同机器 最终级别是在通过网络（本地或互联网）进行通信的不同机器上运行。

线程金字塔上面的每一步都比下面的一步执行得更快，因为芯片上线程之间的距离更小，允许制造商设计快速互连。 这种层次结构还简化了编程模型，以便不同版本的 GPU 可以用相同的方式进行编程。