第五课：分支、谓词执行与 Warp 发散

导语：这一课继续讲控制流：同一个 Warp 遇到 if/else、循环和提前退出时，硬件如何处理，以及分支发散为什么会降低效率。

前面已经知道，NVIDIA GPU 通常把 32 个线程组成一个 Warp，并以 Warp 为单位发射指令。

理想情况下，Warp 中 32 个线程执行相同的指令：

C[i] = A[i] + B[i];

只是每个线程处理的数据不同。

但实际程序中经常存在：

if (...) {
    ...
} else {
    ...
}

这时，Warp 中不同线程可能选择不同路径，从而产生 Warp Divergence，线程束分支发散。

---

一、先看一个没有发散的分支

假设程序为：

if (blockIdx.x == 0) {
    C[i] = A[i] + B[i];
} else {
    C[i] = A[i] - B[i];
}

同一个 Block 中所有线程的 blockIdx.x 相同。

因此，一个 Warp 中的线程通常会一起进入同一个分支：

Block 0 中的 Warp：
Thread 0～31 全部执行加法

Block 1 中的 Warp：
Thread 0～31 全部执行减法

不同 Warp 执行不同路径通常没有问题，因为 Warp 本来就是独立调度的。

真正的问题是：

同一个 Warp 内的线程执行不同路径。

---

二、典型的分支发散

考虑：

if (threadIdx.x % 2 == 0) {
    C[i] = A[i] + B[i];
} else {
    C[i] = A[i] - B[i];
}

一个 Warp 中：

Thread 0：加法
Thread 1：减法
Thread 2：加法
Thread 3：减法
……

Warp 中线程需要执行不同指令。

但是 Warp 在某一时刻通常只能发射同一条指令，因此硬件会把两个路径分别执行。

第一次执行 if 路径：

Thread 0：活跃
Thread 1：屏蔽
Thread 2：活跃
Thread 3：屏蔽
……

第二次执行 else 路径：

Thread 0：屏蔽
Thread 1：活跃
Thread 2：屏蔽
Thread 3：活跃
……

完整过程可以理解为：

图示：二、典型的分支发散

因此，两条路径并不是同时完成的。

---

三、Active Mask 是什么

Warp 中通常会维护一个活跃线程掩码。

假设一个 Warp 有 8 个线程，用 8 位表示：

11111111

表示所有线程活跃。

如果偶数线程执行 if 路径：

01010101

如果奇数线程执行 else 路径：

10101010

硬件执行加法指令时，只有对应掩码为 1 的线程真正写入结果。

例如：

ADD 指令
Active Mask = 01010101

那么部分执行通道虽然仍然存在，但被屏蔽的线程不会更新自己的寄存器或内存。

可以把 Active Mask 理解为：

当前指令允许哪些线程参与执行。

---

四、分支发散为什么降低性能

假设没有分支时，Warp 只需要执行一次加法：

32 个线程同时执行 ADD

发生发散后：

16 个线程执行 ADD
16 个线程空闲

然后：

16 个线程执行 SUB
16 个线程空闲

原本可能只需要一段执行时间，现在需要依次执行两条路径。

如果两个分支长度相近，实际效率可能接近：

$$\frac{\text{useful thread work}} {\text{all issued lane slots}}$$

例如一半线程走 A，一半线程走 B，且两条路径长度相同，那么很多时候平均执行通道利用率只有约 50%。

不过这只是简化理解，实际性能还受到：

• 分支指令数量；
• 指令流水线；
• 内存访问；
• 编译器优化；
• 路径长度；
• 架构实现；

等因素影响。

---

五、分支发散的严重程度取决于路径长度

考虑：

if (condition) {
    x = x + 1;
} else {
    x = x - 1;
}

两个路径都只有一条简单指令。

虽然存在发散，但开销可能不算很大。

再看：

if (condition) {
    // 100 条复杂计算指令
} else {
    // 200 条复杂计算指令
}

如果同一 Warp 中两个分支都有线程执行，那么 Warp 可能需要先后执行约 300 条指令路径。

因此，分支发散是否严重主要取决于：

同一 Warp 是否分裂
×
各分支路径有多长
×
各路径执行频率

---

六、重新汇合 Reconvergence

分支执行结束后，Warp 中的线程通常会在某个位置重新汇合。

例如：

if (condition) {
    x = x + 1;
} else {
    x = x - 1;
}

y = x * 2;

执行过程：

图示：六、重新汇合 Reconvergence

这称为：

Reconvergence，重新汇合。

编译器和硬件通常会确定一个重新汇合点。

对于简单的 if/else，重新汇合点通常位于整个条件结构结束之后。

---

七、什么是后支配点

从控制流图角度，多个分支最终都必须经过的节点，可以称为后支配点。

例如：

          condition
          /       \
       Path A    Path B
          \       /
           Node C

无论走 Path A 还是 Path B，最后都会到达 Node C。

Node C 就可以成为重新汇合位置。

对应代码：

if (condition) {
    path_a();
} else {
    path_b();
}

common_code();  // 可能的重新汇合点

现代 GPU 的具体控制流实现比这个模型复杂，但入门时可以这样理解。

---

八、谓词执行 Predication

对于非常短的分支，编译器可能不真正生成跳转，而是采用谓词执行。

例如：

if (x > 0) {
    y = x;
}

可能被编译成类似：

P0 = (x > 0)
@P0 MOV y, x

其中 P0 是谓词寄存器。

@P0 表示：

如果 P0 为真，执行并写回
如果 P0 为假，不产生有效写回

对于一个 Warp，不同线程都有自己的谓词结果：

Thread 0：P0 = 1
Thread 1：P0 = 0
Thread 2：P0 = 1
Thread 3：P0 = 0

指令仍然被发射，但只有谓词为真的线程更新结果。

---

谓词执行和分支的区别

真正的分支：

根据条件改变程序计数器
分别执行不同代码路径

谓词执行：

指令仍然执行或进入流水线
但部分线程的结果被屏蔽

谓词执行适合短小分支，因为可以避免复杂的跳转和重新汇合。

但如果分支内部有大量指令，谓词执行会让很多线程执行大量无效指令，因此编译器通常更倾向于真正分支。

---

九、一个 if 不一定产生实际分支指令

下面的代码：

if (x > 0) {
    y = x + 1;
}

编译后可能使用谓词。

而下面的代码：

if (x > 0) {
    for (int k = 0; k < 100; ++k) {
        complicated_calculation();
    }
}

更可能生成真正的跳转和控制流。

所以源代码中有 if，不代表硬件一定执行了明显的分支跳转。

需要检查编译后的中间表示或机器指令，才能确定具体实现。

---

十、循环也会产生发散

考虑：

for (int k = 0; k < count[i]; ++k) {
    process();
}

每个线程的 count[i] 不同。

例如：

Thread 0：循环 2 次
Thread 1：循环 10 次
Thread 2：循环 5 次
Thread 3：循环 20 次

执行到第 3 次循环时：

Thread 0：已结束
Thread 1：活跃
Thread 2：活跃
Thread 3：活跃

执行到第 11 次循环时：

Thread 0：已结束
Thread 1：已结束
Thread 2：已结束
Thread 3：活跃

最后可能只剩一个线程工作，其余执行通道全部闲置。

因此，不规则循环通常是 GPU 上比较严重的发散来源。

---

十一、循环发散的实际效率

假设一个 Warp 中：

• 31 个线程只循环 1 次；
• 1 个线程循环 100 次。

整个 Warp 仍然需要运行到那个最慢线程完成。

从第 2 次到第 100 次循环：

只有 1 个线程有效
其余 31 个线程被屏蔽

因此 Warp 的完成时间通常由执行次数最多的线程决定。

这可以概括为：

$$T_{\text{warp loop}} \approx \max_i L_i$$

而不是平均循环次数。

---

十二、提前退出也可能产生发散

例如搜索问题：

for (int k = 0; k < N; ++k) {
    if (data[k] == target[i]) {
        result[i] = k;
        break;
    }
}

不同线程可能在不同位置找到目标：

Thread 0：第 2 次找到
Thread 1：第 50 次找到
Thread 2：没有找到
Thread 3：第 10 次找到

提前结束的线程会被屏蔽，但 Warp 通常仍需等待最慢的线程完成循环。

所以：

单个线程提前退出，不一定意味着对应物理执行资源立刻被其他 Warp 中的线程重新填补。

一个 Warp 的执行通道不会临时混入另一个 Warp 的线程。

---

十三、Warp 中的空闲通道不能被其他 Warp 补上

假设 Warp 0 只有 8 个线程活跃：

Warp 0：
8 个线程工作
24 个线程被屏蔽

硬件一般不会把 Warp 1 的 24 个线程填入剩余通道，组成一个新的 32 线程组合。

原因是每个 Warp 有自己的：

• 程序计数器；
• 指令； -寄存器状态；
• Active Mask；
• 控制流状态。

因此 Warp 0 发射指令时，空闲执行通道通常就只能空闲。

这是分支发散造成利用率下降的根本原因之一。

---

十四、嵌套分支

考虑：

if (condition_a) {
    if (condition_b) {
        path_1();
    } else {
        path_2();
    }
} else {
    path_3();
}

一个 Warp 可能被进一步划分为多个活跃线程集合：

线程集合 1 → path_1
线程集合 2 → path_2
线程集合 3 → path_3

硬件需要分别执行这些路径。

嵌套分支越多、线程条件越不规则，控制流管理越复杂，有效执行通道利用率也可能越低。

---

十五、分支发散只发生在同一 Warp 内

假设 Block 有 256 个线程，即 8 个 Warp。

代码为：

if (threadIdx.x < 128) {
    path_a();
} else {
    path_b();
}

线程分布为：

图示：十五、分支发散只发生在同一 Warp 内

虽然整个 Block 中线程走了不同分支，但每个 Warp 内部路径一致，所以没有明显 Warp 发散。

再看：

if (threadIdx.x % 2 == 0) {
    path_a();
} else {
    path_b();
}

每个 Warp 内一半线程走 A，一半线程走 B，因此所有 Warp 都发散。

这说明，关键不是整个 Block 是否走不同路径，而是线程条件与 Warp 边界之间的关系。

---

十六、数据排列会影响分支发散

假设有 64 个任务：

• 32 个任务属于类型 A；
• 32 个任务属于类型 B。

如果任务顺序为：

A A A A ... A
B B B B ... B

那么：

Warp 0：全部处理 A
Warp 1：全部处理 B

基本不会发散。

如果任务排列为：

A B A B A B ...

那么每个 Warp 都会同时包含 A 和 B：

Warp 0：一半 A，一半 B
Warp 1：一半 A，一半 B

会发生明显发散。

因此，GPU 优化有时不仅要改代码，还要重新组织输入数据，使同一 Warp 中的线程尽量处理相同类型的任务。

---

十七、如何减少分支发散

1. 让分支条件与 Warp 对齐

较差：

if (threadIdx.x % 2 == 0) {
    path_a();
} else {
    path_b();
}

较好：

if ((threadIdx.x / 32) % 2 == 0) {
    path_a();
} else {
    path_b();
}

后一种情况下，一个 Warp 内的线程更可能走相同路径。

这只是说明原则，并不意味着实际程序一定要这样写。

---

2. 按任务类型重新排序数据

原始顺序：

A B A B B A A B ...

重排为：

A A A A ... B B B B ...

让相邻线程处理相同类型的数据。

代价是可能需要：

• 数据预处理；
• 排序；
• 额外索引；
• 结果恢复顺序。

因此需要比较重排开销和发散损失。

---

3. 把不同任务拆成不同 Kernel

原来：

if (type[i] == A) {
    process_a();
} else {
    process_b();
}

可以先把任务分类，然后分别启动：

Kernel A：只处理 A 类任务
Kernel B：只处理 B 类任务

这样每个 Kernel 内控制流更加一致。

但会增加：

• Kernel 启动开销；
• 数据分类开销；
• 额外内存访问。

---

4. 用无分支表达式代替短分支

例如：

if (condition) {
    x = a;
} else {
    x = b;
}

有时可以写成：

x = condition ? a : b;

或者使用数学表达式：

x = condition * a + (1 - condition) * b;

但后者未必总是更快，也可能增加额外计算。

编译器通常已经能对简单条件做谓词化，因此不应机械地消除所有 if。

---

5. 减少不规则循环

如果不同线程循环次数差异很大，可以考虑：

• 固定循环次数；
• 分阶段处理；
• 把长任务与短任务分开；
• 使用工作队列；
• 按任务长度排序；
• 采用 persistent thread 等更复杂方案。

这些方法会增加实现复杂度，需要根据实际性能瓶颈选择。

---

十八、分支发散不一定是首要问题

不能看到 if 就认为程序一定很慢。

例如：

if (i < N) {
    C[i] = A[i] + B[i];
}

只有最后一个 Warp 中可能有少量线程越界。

如果总共有数万个 Warp，而只有最后一个 Warp 不完整，那么影响几乎可以忽略。

同样，边界处理：

if (x < width && y < height) {
    ...
}

通常只影响图像边缘的一小部分线程。

所以优化时应关注：

• 发散发生得是否频繁；
• 影响多少 Warp；
• 分支路径是否很长；
• 是否确实成为性能瓶颈。

---

十九、现代 NVIDIA GPU 的独立线程调度

早期可以把 Warp 简单理解为：

Warp 中所有线程共享一个程序计数器，严格锁步执行。

从 Volta 架构开始，NVIDIA 引入了更灵活的独立线程调度机制。线程可以维护更细粒度的执行状态，硬件可以在更复杂的条件下组织线程执行。

但这并不意味着：

• Warp 消失了；
• 32 个线程完全像 CPU 线程一样独立；
• 分支发散没有代价；
• Warp 级同步不再需要注意。

Warp 仍然是重要的发射和执行组织单位，发散仍然会导致执行路径被分别处理。

入门阶段可以继续使用下面的模型：

Warp 统一发射指令
不同路径通过 Active Mask 分别执行
之后重新汇合

只是要知道，现代硬件的真实实现比这个模型更灵活。

---

二十、独立线程调度带来的同步问题

过去有些程序默认：

同一个 Warp 中的线程天然始终同步。

例如：

shared_data[lane] = value;
value = shared_data[lane - 1];

程序员可能认为 Warp 内线程锁步执行，因此不需要同步。

在现代架构中，这种假设可能不安全。

当 Warp 内线程通过 Shared Memory 交换数据时，必要时应使用明确的同步机制，例如：

__syncwarp();

示意：

shared_data[lane] = value;

__syncwarp();

if (lane > 0) {
    value = shared_data[lane - 1];
}

__syncwarp() 用于同步指定 Warp 线程，并处理相关的内存可见性要求。

---

二十一、__syncthreads() 和 __syncwarp() 的区别

__syncthreads()：

同步整个 Block

同一个 Block 中所有参与线程都需要到达同步点。

__syncwarp()：

同步一个 Warp 中指定的线程

作用范围更小，通常开销也更低。

可以简单区分为：

同步方式	同步范围
__syncwarp()	Warp 内线程
__syncthreads()	整个 Block

如果数据需要在不同 Warp 之间共享，仅使用 __syncwarp() 不够。

---

二十二、分支与 __syncthreads() 的危险组合

考虑：

if (threadIdx.x < 128) {
    __syncthreads();
}

假设一个 Block 有 256 个线程。

只有前 128 个线程到达屏障，后 128 个线程没有到达。

这可能导致：

前 128 个线程一直等待
后 128 个线程不会进入屏障
Block 无法继续

可能造成死锁或未定义行为。

因此，通常要求：

Block 中所有需要参与同步的线程，以一致的控制流到达 __syncthreads()。

安全形式：

if (threadIdx.x < 128) {
    do_something();
}

__syncthreads();

这里所有线程最终都会执行同步指令。

---

二十三、Warp Vote 指令

GPU 提供一些 Warp 级操作，用于在 Warp 内汇总线程条件。

常见操作包括：

__all_sync(mask, predicate)
__any_sync(mask, predicate)
__ballot_sync(mask, predicate)

含义如下。

__all_sync：

指定线程是否全部满足条件。

__any_sync：

指定线程中是否至少一个满足条件。

__ballot_sync：

返回每个线程条件结果组成的位掩码。

例如：

unsigned int result =
    __ballot_sync(0xffffffff, x > 0);

如果 Thread 0、2、3 的 x > 0，可能得到类似：

0000...1101

每一位对应一个线程的判断结果。

这些操作常用于：

• 判断 Warp 是否可以整体退出；
• 压缩有效线程；
• Warp 级归约；
• 构造任务掩码；
• 优化不规则控制流。

---

二十四、Warp Shuffle

Warp 中线程还可以直接交换寄存器数据，例如：

__shfl_sync(mask, value, source_lane);

它允许一个线程读取同一 Warp 中另一个线程的寄存器值。

例如：

int other =
    __shfl_sync(0xffffffff, value, 0);

可以让所有线程读取 lane 0 的 value。

这类操作不需要经过 Shared Memory，常用于：

• Warp 归约；
• 前缀和；
• 数据广播；
• 矩阵计算；
• 高效线程通信。

不过执行 Shuffle 时，必须正确指定参与线程掩码，并保证参与线程的控制流符合要求。

---

二十五、一个 Warp 归约示例

假设 Warp 中每个线程有一个数，希望求和。

可以通过 Shuffle 逐步相加：

for (int offset = 16; offset > 0; offset /= 2) {
    value += __shfl_down_sync(
        0xffffffff,
        value,
        offset
    );
}

过程大致为：

第一次：lane 0 加 lane 16
第二次：lane 0 加 lane 8 的部分和
第三次：继续合并
……

最终 lane 0 得到整个 Warp 的总和。

这体现了 Warp 内线程虽然共享一条执行指令，但每个线程仍有自己的寄存器数据，并可以通过专用网络交换。

---

二十六、分支发散和数据冒险不是一回事

分支发散：

同一个 Warp 内线程走不同控制路径

数据依赖：

后一条指令需要前一条指令的结果

例如：

a = b + c;
d = a * e;

这里有数据依赖，但没有分支发散。

例如：

if (threadIdx.x % 2 == 0) {
    x++;
}

这里可能有分支发散，但不一定存在复杂的数据依赖。

两者都可能使流水线停顿或利用率下降，但原因不同。

---

二十七、从硬件角度总结控制流过程

假设 Warp 执行：

if (condition) {
    A();
} else {
    B();
}

C();

可以简化为：

1. 每个线程计算 condition
2. 生成线程条件掩码
3. 判断是否所有线程结果相同
4. 如果相同：
      只执行对应路径
5. 如果不同：
      保存另一条路径状态
      用 Mask A 执行路径 A
      用 Mask B 执行路径 B
6. 到达重新汇合点
7. 恢复完整活跃线程集合
8. 继续执行 C

需要注意，具体硬件实现可能采用：

• 分支栈；
• 重汇合栈；
• 每线程状态；
• 编译器插入的控制指令；
• 独立线程调度；

但功能目标基本相同。

---

二十八、本节核心结论

第一，分支发散只关心同一个 Warp 内的线程是否走不同路径。

第二，发生发散后，多个路径通常需要分别执行，而不是同时执行。

第三，Active Mask 用于指定当前指令中哪些线程有效。

第四，被屏蔽线程通常不会产生有效结果，但仍会占据 Warp 的执行位置。

第五，Warp 中的空闲执行通道不能临时由其他 Warp 的线程填充。

第六，路径越长、分裂越严重，分支发散带来的性能损失通常越大。

第七，循环次数不一致也会导致发散，Warp 的执行时间通常受最慢线程影响。

第八，短小条件可能被编译为谓词执行，而不是真正的分支跳转。

第九，不同 Warp 走不同分支通常没有明显发散问题。

第十，数据重排、任务分类和控制流重构可以减少发散。

第十一，从 Volta 开始，现代 NVIDIA GPU 支持更灵活的独立线程调度，但 Warp 和分支发散仍然存在。

第十二，不能依赖隐式 Warp 锁步；Warp 内通信必要时应使用 __syncwarp()。

第十三，__syncthreads() 需要谨慎放置，不能让部分 Block 线程永久无法到达屏障。

下一课适合继续学习：

第六课：GPU 中的数据依赖、流水线停顿与指令调度

主要包括：

RAW、WAR、WAW 数据依赖
GPU 是否存在数据冒险
Scoreboard 如何解决 RAW 依赖
为什么 GPU 通常不需要 CPU 式乱序执行
同一个 Warp 和不同 Warp 的依赖关系
访存依赖与内存一致性
原子操作
Memory Fence