title: 什么是线程束发散？

线程束发散 (warp divergence) 发生在线程束 (warp) 内的线程由于控制流语句而采取不同执行路径时。

例如，考虑以下内核 (kernel)：

__global__ void divergent_kernel(float* data, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        if (data[idx] > 0.5f) {
            // A
            data[idx] = data[idx] * 4.0f;
        } else {
            // B
            data[idx] = data[idx] + 2.0f;
        }
        data[idx] = data[idx] * data[idx];
    }
}

当线程束 (warp) 内的线程 (threads) 遇到数据相关的条件语句时，根据 data[idx] 的值，一些线程 (threads) 必须执行代码块 A，而其他线程必须执行代码块 B。由于这种数据依赖性以及CUDA 编程模型 (CUDA programming model) 及其在PTX 机器模型 (PTX machine model) 中实现的结构性约束，程序员或编译器无法避免线程束 (warp) 内部控制流的这种分裂。

相反，线程束调度器 (warp scheduler) 必须处理这些发散代码路径的并发执行，它通过"屏蔽"一些线程 (threads) 使其不执行指令来实现这一点。这是通过使用谓词寄存器 (registers) 实现的。

让我们检查生成的SASS (Streaming Assembler) (Godbolt 链接) 来理解执行流程：

LDG.E.SYS R4, [R2]                       // L1 加载 data[idx]
FSETP.GT.AND P0, PT, R4.reuse, 0.5, PT   // L2 设置 P0 为 data[idx] > 0.5
FADD R0, R4, 2                           // L3 将 2 + data[idx] 存储到 R0
@P0 FMUL R0, R4, 4                       // L4 在某些线程中，将 4 * data[idx] 存储到 R0
FMUL R5, R0, R0                          // L5 将 R0 * R0 存储到 R5
STG.E.SYS [R2], R5                       // L6 将 R5 存储到 data[idx]

将数据加载到 R4 (L1) 后，线程束 (warp) 中的所有 32 个线程 (threads) 并发执行 FSETP.GT.AND (L2)，每个线程 (thread) 根据 R4 中的 data 值获得自己的 P0 值。然后，我们看到一点编译器 (nvcc) 的巧妙之处：在 L3 中，所有线程 (threads) 都执行代码块 A 的代码，写入 R0。只有那些 P0 为真的线程随后执行代码块 B 的代码 (L4)，覆盖在 L3 中写入 R0 的值。在这条指令上，线程束 (warp) 被称为"发散的"。在 L5 上，所有线程 (threads) 都回到执行相同的代码。一旦线程束调度器 (warp scheduler) 通过在同一时钟周期发出相同指令使它们重新对齐，线程束就"收敛"了。

这可能比将分支天真地编码到SASS (Streaming Assembler) 中更高效，后者会对 L3 和 L4 两行都进行谓词化 — 说"可能"是因为我们可以信任编译器 (nvcc)，并且启发式地看，我们是在用廉价、充足的CUDA 核心 (CUDA Core) 计算来换取更昂贵的流控制。正如在 GPU 编程中常见的那样，即使只是简单的谓词化，浪费计算（每次执行 L4 时进行一次不必要的 FADD）也比增加复杂性要好！

编译器可能积极避免发散的一个原因是，在早期（Volta 架构之前）的 GPU 中，发散的线程束 (warps) 总是完全串行化的。虽然线程束发散仍然会降低效率，但具有独立线程调度的现代 GPU 不一定会经历完全的串行化惩罚。