Nvidia GPU 中，一个 Warp 包含 32 个线程，
这里为了说明的简便，假定一个 Warp 包含 4 个线程。
不管 Warp 中包含多少个线程，它们都并行执行相同的指令，因此可以认为它们有一个共同的 CURR_PC，
但由于有分支指令的存在，它们即将执行的下一条指令可能各不相同，因此，它们都有各自的 NEXT_PC。
Warp 中的 4 个线程分别记为 T0、T1、T2 和 T3，它们即将执行的下一条指令的 PC 记为 T~N~_NEXT_PC，它们正在执行的指令的 PC 记为 CURR_PC。
给定一个 Kernel 程序，如下，
__kernel void
test(__global int* OUT) {
int tid = (int)(get_global_id(0)); // ← PC 0 // 获取当前线程的 ID
if (tid < 2) { // ← PC 1 // 如果当前线程的 ID 小于 2
OUT[tid] += 10; // ← PC 2
}
OUT[tid] += 100; // ← PC 3
}
执行过程，如下，
执行到 PC-1 的时候，由于 4 个线程的 NEXT_PC 不一样，因此在这里发生了线程分叉（Diverge）。
发生 Diverge 时，需要选择 NEXT_PC 最小的线程来继续执行，其他线程陪跑。
本例中，选择 T0 和 T1 执行，T2 和 T3 陪跑。
执行到 PC-2 的时候，由于执行到了分支代码的最后一行，因此在这里需要做线程合并（Reconverge）。
Reconverge 的时机，硬件并不能自动判断得到。需要编译器在指令中标记好。
发生 Reconverge 时，需要选择 NEXT_PC 最小的线程来继续执行，其他线程陪跑。
本例中，选择 T0、T1、T2 和 T3 执行。
分支指令会非常影响 GPU 的性能，这是因为 Diverge 时会有线程空跑，出现了执行“气泡”，降低了执行单元的利用率。
为了加深印象，提供一个带嵌套分支的 Kernel 程序，读者可以自行推理一下何时会发生 Diverge 或者 Reconverge。
__kernel void
test(__global int* OUT) {
int tid = (int)(get_global_id(0)); // ← PC 0
if (tid < 2) { // ← PC 1
OUT[tid] += 10; // ← PC 2
if (tid < 1) // ← PC 3
OUT[tid] -= 10; // ← PC 4
else
OUT[tid] -= 5; // ← PC 5
} else
OUT[tid] += 20; // ← PC 6
OUT[tid] += 100; // ← PC 7
}
参考的执行过程，如下，
线程发散会带来 GPU 执行性能的降低。减少分支发散的方法一般有：
-
减少一个 Warp 包含的线程的数量。例如，假设程序逻辑为：当线程 ID 小于 2 时执行 if 语句内容，且一个 Warp 是 4 个线程。那么此时会发生线程分散。如果设置一个 Warp 只包含 2 个线程，那么就不会发生线程分散了。该方法降低了线程发散的概率
英伟达一个 Warp 包含 32 个线程
ARM 一个 Warp 包含 4/8/16 个线程
-
AMD 每个 Warp 包含 32 或者 64 个线程
- 官方推荐图形渲染时使用 64 的配置，通用计算时使用 32 的配置，因为通用计算通常包含更加复杂的控制逻辑，会出现更多的分支指令
一个 Warp 中的线程并不是越少越好，这是硬件架构设计过程中综合多个因素而定的
分支合并
参看：fetch.php (cmu.edu)
