---

一、先给一个总览

Warp divergence，中文通常叫线程束发散，是 CUDA 性能优化中最基础也最重要的概念之一。

它的核心定义是：

同一个 warp 内的 32 个线程，在执行到某个分支时，走了不同的路径，导致硬件必须串行执行这些路径。

在 SIMT（Single Instruction, Multiple Threads）架构下，一个 warp 内的 32 个线程理论上应该执行同一条指令，只是数据不同。一旦出现分支，硬件无法让 32 个线程同时走两条路，只能：

• 先执行路径 A，此时走路径 B 的线程被禁用；
• 再执行路径 B，此时走路径 A 的线程被禁用。

这样原本一次能做完的事情，变成了两次甚至多次，吞吐直接下降。

一句话总结：warp divergence 让 SIMT 的并行执行退化成串行执行，浪费了 warp 内的并行度。

---

二、Warp divergence 为什么会拖慢 kernel

1、SIMT 执行模型

CUDA 的 GPU 以 warp 为调度单位，一个 warp 包含 32 个线程。

在理想情况下：

• 32 个线程执行同一条指令；
• 每个线程操作自己的数据；
• 一条指令完成 32 个线程的工作。

这就是 SIMT 的并行来源。

2、出现 divergence 时

假设代码是：

if (condition) {
    // 路径 A
} else {
    // 路径 B
}

如果同一个 warp 内：

• 前 16 个线程 condition 为真；
• 后 16 个线程 condition 为假。

那么硬件会：

• 1、执行路径 A，只激活前 16 个线程，后 16 个线程 masked off；
• 2、执行路径 B，只激活后 16 个线程，前 16 个线程 masked off。

原本一条路径能完成的工作，变成了两条路径串行执行。

3、性能损失的量化

假设一个 warp 内分支有 $k$ 条不同路径，且各路径线程数不均：

• 理想情况：1 次执行；
• 实际：$k$ 次执行；
• 利用率可能降到 $1/k$。

如果分支嵌套，路径数指数增长，性能下降更严重。

4、不只是 if-else

以下情况也会导致 divergence：

• switch-case 多分支；
• 循环次数不一致，比如 for (i = 0; i < n; i++) 中每个线程的 $n$ 不同；
• 提前 return 或 break；
• 短循环，warp 内线程循环次数不同；
• 谓词执行不完全等价于无 divergence，但通常比分支好。

5、为什么短分支也贵

即使分支体很短，只要出现 divergence，硬件仍要：

• 分别发出指令；
• 管理 active mask；
• 可能导致流水线停顿。

所以频繁的小分支累积起来也很可观。

---

三、Transformer 自定义算子中容易出现的 warp divergence

1、Attention 中的变长序列和 mask

这是最典型的场景。

在 batch 内，不同序列长度不同，需要 padding 或 mask：

• 短序列后面是 padding；
• padding 位置的 attention 要 mask 掉；
• 同一个 warp 内，有的线程处理有效 token，有的处理 padding。

如果实现不当，就会出现 divergence。

具体表现：

• if (mask == 0) continue;
• if (col_idx >= seq_len) break;
• causal mask 下三角判断。

2、Causal Mask 的下三角判断

自回归 attention 中，每个 query 只能看前面的 key：

$$ \text{mask}(i, j) = \begin{cases} 0 & j \le i \ -\infty & j > i \end{cases} $$

在 tiling 实现中，如果 tile 跨越对角线，同一个 warp 内不同线程可能走不同分支。

3、MoE 中的路由

MoE（Mixture of Experts）中：

• 每个 token 被路由到不同的 expert；
• 不同 expert 计算量不同；
• 同一个 warp 内 token 可能去不同 expert。

这会导致严重的 divergence，也是 MoE kernel 优化的难点。

4、变长序列的预处理

处理变长序列时：

• 每个序列长度不同；
• 循环边界不同；
• 同一个 warp 内线程循环次数不一致。

比如 packing、unpadding、位置编码等。

5、条件归一化和激活

某些算子中有条件逻辑：

• 动态缩放；
• 条件 dropout；
• 自适应计算。

如果按元素判断，容易 divergence。

6、量化中的异常值处理

量化 kernel 中：

• 检测异常值；
• 对异常值单独处理；
• 同一个 warp 内有的线程触发，有的不触发。

7、稀疏化和剪枝

稀疏 kernel 中：

• 每个线程处理的非零元素数量不同；
• 循环次数不一致；
• 容易 divergence。

---

四、如何避免和缓解 warp divergence

1、数据重排，让同 warp 内线程行为一致

核心思路：

让同一个 warp 内的线程尽量走同一条路径。

方法：

• 按长度排序，把长度相近的序列放一起；
• 用 bucket 把长度接近的样本分到同一 batch；
• 变长序列 packing，减少 padding；
• MoE 中按 expert 分组处理，而不是按 token 顺序。

2、用谓词执行代替分支

对于短分支，可以用谓词执行：

// 不推荐：分支
if (cond) {
    y = a;
} else {
    y = b;
}

// 推荐：谓词
y = cond ? a : b;

编译器有时会自动做，但显式写更可控。

谓词执行不是完全没有代价，但比分支轻。

3、用数学技巧消除分支

很多分支可以用数学表达式替代：

• 用 max、min 代替 if-else；
• 用 mask 乘法代替条件赋值；
• 用 step、select 等函数。

比如 causal mask 可以用：

$$ \text{score} = \text{score} + \text{mask} \cdot (-\infty) $$

其中 mask 是 0 或 1，避免显式分支。

4、循环边界对齐

对于循环次数不一致：

• 让所有线程循环相同次数；
• 用边界判断保护越界访问；
• 用 __syncthreads() 对齐。

比如：

for (int i = 0; i < max_len; i++) {
    if (i < len) {
        // 有效计算
    }
}

虽然仍有谓词，但避免了不同线程循环次数不同导致的 divergence。

5、用 warp-level 原语

CUDA 提供 warp-level 原语：

• __ballot_sync：投票；
• __shfl_sync：warp 内 shuffle；
• __any_sync、__all_sync：聚合判断。

这些可以帮助在 warp 内做统一决策，减少 divergence。

6、分离 kernel

如果不同路径差异很大：

• 拆成多个 kernel；
• 每个 kernel 处理一类数据；
• 避免在一个 kernel 内混合。

比如 MoE 中：

• 先做路由；
• 再按 expert 分组；
• 每个 expert 单独 kernel。

7、使用 CUTLASS、Triton 等成熟库

这些库已经做了大量优化：

• 减少 divergence；
• 优化内存访问；
• 自动调优。

如果不是必须手写，优先用成熟库。

8、使用 shared memory 做数据重排

在 warp 内或 block 内：

• 先把数据写到 shared memory；
• 再按统一模式读回；
• 让后续计算对齐。

9、合理设置 block 和 warp 大小

• 让一个 warp 处理同类数据；
• 避免一个 warp 跨越不同长度的序列；
• 调整 block 维度，让线程映射更规整。

10、Profiling 定位

用 Nsight Compute 看：

• branch efficiency；
• warp stall reason；
• divergence 相关指标。

不要凭感觉，要用数据定位。

---

五、一个具体例子：Causal Attention

1、朴素实现

for (int j = 0; j < seq_len; j++) {
    if (j <= i) {
        score += q[i] * k[j];
    }
}

同一个 warp 内，不同线程的 $i$ 不同，导致 j <= i 判断结果不同，出现 divergence。

2、优化实现

方法一：用 mask 替代分支

float mask = (j <= i) ? 0.0f : -INFINITY;
score += q[i] * k[j] + mask;

方法二：tiling + 对角线特殊处理

• 非对角线 tile 全部有效，无分支；
• 对角线 tile 单独处理，用 mask；
• 减少 divergence 范围。

方法三：按 query 分块，让 warp 内 query 连续

• 同一个 warp 处理相近的 query；
• mask 模式接近；
• divergence 减少。

---

六、Warp divergence 和分支效率

Nsight Compute 中有个指标叫 branch efficiency：

$$ \text{Branch Efficiency} = \frac{\text{实际执行的分支数}}{\text{理想无 divergence 的分支数}} $$

如果 branch efficiency 低，说明 divergence 严重。

但要注意：

• 不是所有 divergence 都值得优化；
• 如果分支体很小，优化收益有限；
• 优先优化热路径上的 divergence。

---

七、简单评价这道面试题

1、这道题的价值

• 这是 CUDA 性能优化中最基础的一道题；
• 一道题能考察 SIMT 执行模型、divergence 原理、Transformer 中的实际场景和优化手段；
• 区分度好，初级可能只知道“分支会慢”，中级能说清 masked execution，高级能结合 Attention、MoE 讲实际优化；
• 贴近实战，写自定义算子必然遇到。

2、这道题的不足

• 题目偏基础，对高级 CUDA 候选人来说不够深；
• 没有涉及更复杂的性能问题，比如 bank conflict、occupancy、指令级并行；
• 如果候选人只背概念，容易答得泛泛；
• 对 profiling 和量化分析的考察不够。

3、建议的追问方向

• 谓词执行和分支在硬件层面有什么区别？
• 如何用 Nsight Compute 定位 divergence？
• MoE 中怎么避免路由带来的 divergence？
• causal mask 有哪些无分支实现？
• divergence 和 occupancy 之间有什么关系？
• 如果分支不可避免，怎么减少它的影响？

4、总体评价

这是一道很好的 CUDA 基础面试题，适合考察候选人对 SIMT 执行模型和 divergence 的理解。它既能考察基础原理，也能延伸到 Transformer 自定义算子的实际优化。如果候选人能把 divergence 的原理、性能损失、Transformer 中的典型场景以及多种避免手段都讲清楚，并能结合 profiling 数据做分析，基本可以判断其具备 CUDA 算子优化的扎实基础。

---

---