3.7. 什么是线程束组 ？

Warpgroup（线程束组）是指四个连续 线程束 (warps) 的集合，且该组中第一个线程束的秩（warp-rank）必须是 4 的倍数。

在分发 Warpgroup 级别的指令时，系统会协同 128 个 线程 —— 即每个 Warpgroup 包含 4 个线程束，每个线程束包含 32 个线程。以更大的粒度进行操作，消除了显式线程束间同步的需求，并允许单条指令处理更大规模的问题，特别是更大的矩阵乘法。更大规模的矩阵乘法能够更容易地使近期数据中心 GPU 中 Tensor Core (张量核心) 的海量 算术带宽 达到饱和。

Warpgroup 引入于 NVIDIA 的 Hopper 流式多处理器（SM）架构 中，用于支持 Warpgroup 级别的矩阵乘法，例如 wgmma.mma_async 指令。欲深入了解，请参阅 Colfax 的这篇博客文章。Warpgroup 在高性能 Hopper 和 Blackwell 内核（kernels）（如 Flash Attention 4）的流水线组件组织中占据重要地位。

在 并行线程执行 (PTX) 中间表示（IR）中，一个线程束的秩（warp-rank）定义为：

int linearIdx = (%tid.x + %tid.y * %ntid.x  + %tid.z * %ntid.x * %ntid.y);
int warpRank = linearIdx / 32;

其中 tid 是线程索引，通过特殊的 PTX 寄存器 进行访问。

因此，对于包含 8 个线程束的调度，有效的 Warpgroup 分别为：

• Warpgroup 0: 秩为 0, 1, 2 和 3 的线程束

• Warpgroup 1: 秩为 4, 5, 6 和 7 的线程束

据我们需要，关于这种线程束秩（warp-rank）对齐限制的目的尚无官方文档说明。但近期数据中心 GPU 的 流式多处理器（SM） 似乎包含四个（未命名的）子单元，每个子单元都拥有独立的 线程束调度器 和 Tensor Core。