SIMT 分支发散：warp 遇到 if/else 为什么利用率腰斩

更新时间：2026-09-25。本文是「集成电路大类 · 芯片设计」进阶层（intermediate） 的一篇。GPU 之所以敢堆几千个计算核心，靠的是一种叫 SIMT（Single Instruction, Multiple Threads，单指令多线程） 的执行模型：32 个线程组成一个 warp，像一支绑腿齐步走的队伍，同一拍执行同一条指令。这套模型在数据并行时近乎完美，可一旦队伍里有人想走 if、有人想走 else，齐步走就被迫拆成两轮——这就是 SIMT 分支发散（warp divergence）。为什么 GPU 选型要走吞吐路线见 GPU 与 CPU：为什么 GPU 走向了众核吞吐，warp 调度器的硬件全景见 GPU 架构深入：SIMT、warp 调度与内存层级。

本文要回答的问题

• SIMT 到底「多」在哪？它和 CPU 上的 SIMD 向量指令是一回事吗？
• 一个 warp 里 32 个线程条件判断结果不同，硬件这一拍该听谁的？
• if/else 各走一半时，利用率具体掉多少？最坏情况能掉到什么程度？
• 等内存要等几百个周期，GPU 为什么不怕？warp 切换凭什么说「零开销」？
• 写 CUDA kernel 时，哪些写法必然发散，怎么改才能让 warp 重新齐步走？

一、SIMT：一条指令驱动 32 个线程的锁步执行

GPU 的基本执行单元是 SM（Streaming Multiprocessor，流式多处理器），内部结构见 SM/CU 内部结构。SM 不直接调度单个线程，而是以 warp 为最小调度单位：一个 warp 固定包含 32 个线程（NVIDIA 的口径；AMD 对应的概念叫 wavefront，宽度 64），它们共享同一个程序计数器，在每一个时钟周期执行同一条指令，只是各自作用在自己那份数据上。这种执行方式叫锁步（lock-step）执行。

线程在 warp 里的编号叫 lane，从 0 到 31。一条加法指令发出后，32 个 lane 的加法器同时动作，一拍完成 32 个加法。线程层级（grid/block/thread）如何映射到 warp，见 CUDA 线程层次：grid、block 与 warp。这里先记住一个关键事实：block 只是软件的组织形式，硬件眼里只有 warp，同一个 block 的线程被切成若干 warp，warp 之间的执行顺序没有任何保证。

SIMT 这个名字很容易和 CPU 的 SIMD 混为一谈：两者确实都是「一条指令驱动多份数据」，但抽象层次和硬件代价完全不同。

对比维度	SIMD（AVX/NEON 向量指令）	SIMT（GPU warp 执行）
编程抽象	向量寄存器 + 显式向量指令，靠编译器或手写 intrinsic 把数据打包进宽寄存器	程序员写普通标量线程，硬件自动把 32 个线程组成 warp 锁步执行
每通道寄存器状态	一条向量指令操作一个整体向量寄存器，通道没有独立的线程身份	每个线程有自己的寄存器堆、线程号和（Volta 架构之后）独立 PC 状态
条件分支	没有原生通道级分支，靠掩码寄存器做谓词执行，分支逻辑需手工展开	硬件提供 active mask 和发散/收敛栈，自动串行化 warp 内不同路径
延迟容忍	靠乱序执行和少量硬件线程，停顿通常直接变成流水线气泡	靠几十个 warp 零开销切换，用并行度换延迟
典型代表	x86 AVX-512、ARM SVE/NEON	NVIDIA CUDA warp、AMD wavefront

block 切 warp 的方式也值得记住：block 里的线程按线性线程号连续切分，每 32 个一个 warp，128 线程的 block 正好是 4 个 warp。这解释了两个常见现象：一是 block 大小不是 32 的整数倍时，最后一个 warp 会带若干永不工作的「僵尸 lane」，block 取 96、100 这类数字等于白送一部分发射槽；二是多维线程索引（比如 threadIdx.x、threadIdx.y）在一个 warp 里通常只有最内层连续变化，想让一个 warp 协作处理一段连续数组，数据划分要顺着这个方向做。

一句话概括：SIMD 是程序员手动打包的向量，SIMT 是硬件帮你把标量线程自动锁步成向量。后者编程简单，代价是硬件必须回答一个 SIMD 不用面对的问题——线程们要是想执行不同指令怎么办。

二、active mask：硬件用一个 32 位掩码管理「谁在干活」

答案是谓词执行（predicated execution）加活动掩码。每个 warp 附带一个 32 位的 active mask（活动掩码），某一位为 1，表示对应 lane 此刻是活动的、指令的结果对它生效；为 0 的 lane 不执行，寄存器保持原样。这一拍不是 32 个线程「投票决定」执行哪条指令，而是 warp 按路径逐条发射指令，每条指令只对掩码允许的 lane 生效。

遇到条件分支时，硬件先求出 32 个 lane 各自的条件结果，形成一个掩码，然后按路径组织执行：

• 所有 32 个 lane 条件一致（uniform branch）：要么全员走 if、要么全员走 else，另一条路径整体跳过，零损失，等效于普通分支。
• 条件不一致（non-uniform）：warp 发散，先走 if 路径——此时掩码里只有条件为真的 lane，其余 lane 空闲等待；走完后在分支汇合点（post-dominator）切换掩码，再走 else 路径，最后全员重新收敛。

上图就是 warp 发散掩码执行的完整时序。读图抓一个点：两轮执行期间发射槽都没有闲着，但每轮只有一半的 lane 在做有用功——指令照发、功耗照耗，有效产出减半。分支路径的保存和恢复由硬件的发散栈（reconvergence stack）自动管理，软件感知不到；2017 年 Volta 架构之后每个线程有了独立 PC，线程间同步更灵活，但「锁步发射 + 掩码执行 + 发散串行」的基本规则和性能代价依旧。

三、利用率账：五五对开腰斩，最坏只剩 1/32

把发散区（从分支到汇合之间的代码）的 lane 利用率算一算，数字非常直观。

warp 内分流情况	串行轮数	发散区平均 lane 利用率
全员同向（uniform）	1 轮	100%
if/else 各 16 个	2 轮	50%
24 个走 if、8 个走 else	2 轮	50%（每轮都按 32 槽计）
三路分支各约 11 个	3 轮	约 33%
32 个 lane 走 32 条不同路径	最多 32 轮	最低约 3%

几个容易误判的细节值得点明。其一，利用率只和「路径数」挂钩，不和分流比例成正比：24/8 的分流同样要两轮，平均利用率恰好也是 50%——人少的那一轮 24 个 lane 全在空等。其二，利用率腰斩只发生在发散区，分支前后的无分支代码仍是 100%，所以 kernel 整体损失要按发散区耗时占比折算：一段只占 10% 时间的发散代码，全局损失只有 5 个百分点左右。其三，嵌套分支是乘法关系：外层五五开、内层再五五开，最坏路径数累乘，优化器里最隐蔽的性能黑洞往往就是这种多层嵌套。

用一个具体例子把账算实。假设某 kernel 每个 block 128 线程、共 4 个 warp，发散区每轮约 100 条指令，发散前与汇合后各有 200 条无分支指令。无发散时整个 block 完成约 400 条指令；发散后每个 warp 多走一轮，耗时变成 500 条，block 整体慢了 25%。若把分流改成 24/8，轮数不变，耗时一模一样——这就是「分流再悬殊也救不回轮数」的含义。反过来，如果发散区只有 20 条指令、其余 380 条不发散，总耗时 420 条，损失只剩 5%，这种发散完全可以不处理。

还有一种非 if 形式的发散：数据相关的循环次数。while (x[tid] > 0) x[tid] >>= 1; 这类代码，每个 lane 的循环次数由数据决定，warp 必须等循环次数最多的那个 lane 走完全部轮次，提前结束的 lane 全程掩码等待。它的利用率等于平均迭代次数除以最大迭代次数，数据越参差越难看。

四、warp 调度：用零开销切换把数百周期延迟藏起来

发散让单个 warp 的 lane 利用率下降，而 GPU 对抗另一种浪费——停顿——的办法，是让 SM 同时驻留大量 warp，靠调度隐藏延迟。SM 内部一般有 4 个 warp 调度器，最多驻留 48 到 64 个 warp（如 A100 的 SM 可驻留 64 个 warp、共 2048 个线程）。一次全局显存访问的延迟约 400 到 800 个周期，任何一个 warp 发起访存后都会停住等待，调度器立刻在就绪 warp 里挑一个发射，执行槽因此始终被填满。

为什么这种切换敢称零开销（zero-overhead）？CPU 超线程切换的代价再小也有状态共享的摩擦，而 GPU 每个驻留 warp 的寄存器一直实实在在地占在 SM 的大寄存器堆里（A100 每个 SM 的寄存器堆约 256KB），warp 停住时寄存器原封不动，换 warp 只是调度器改选一个程序计数器，没有任何保存、恢复动作。

需要多少 warp 才够藏住延迟，可以粗算一笔：一次访存延迟按 600 周期计，一个 warp 在两轮访存之间平均执行约 20 条指令，那么单个 warp 每 20 拍就要「请假」600 拍；想让发射槽始终有就绪指令，理论上需要约 600/20 = 30 个 warp 轮流站岗。这正是 SM 支持几十个 warp 驻留的原因，也是「访存密集型 kernel 需要高 occupancy、纯计算 kernel 低 occupancy 也能跑满」的根据——计算密集时 warp 很少请假，少量战友就够轮转。此外，warp 内部连续两条独立指令也能靠调度交错执行，称为指令级并行，让每个 warp 多挂一些独立的进行中操作，等价于减少需要的 warp 数。

这里引出 occupancy（占用率） 的概念：SM 实际驻留 warp 数与上限的比值。寄存器用得太多、每个 block 线程数不合理，都会压低驻留 warp 数，occupancy 不足时没有就绪 warp 可换，显存延迟就裸露成气泡。但 occupancy 不是越高越好：实测中 50% 左右往往已足以藏住延迟，剩下的应让给寄存器复用和单 warp 指令级并行。注意区分两个概念——发散伤的是 warp 内 32 个 lane 的利用率，occupancy 伤的是 warp 之间填满发射槽的能力，一个 kernel 完全可能 occupancy 很高却被发散拖垮，反之亦然。

五、典型发散场景：先能认出它，才谈得上规避

常见的发散写法可以归成四类。

按线程号取奇偶。 if (tid % 2 == 0) 让 lane0、2、4… 与 lane1、3、5… 在 warp 内完美交错，是最教科书式的五五发散。树形归约（reduction）的朴素写法就长这样，应改用 warp 级原语或跨步访存的归约结构。

按数据值分支。 if (a[tid] > threshold)、分类、裁剪（clipping）等，走向完全取决于输入数据，数据随机时 warp 内必然参差不齐。

数据相关的循环与递归式展开。 前面算过的变长 while、稀疏数据结构（哈希表、图遍历、稀疏矩阵）里的空元素跳过，邻居数不同的图算法是重灾区。

以 block/grid 级条件误写成线程级。 条件其实对整个 block 或整个 grid 都一致，却写在 kernel 里按线程判断——这类是「假发散」，本该零成本，改成按 blockIdx 判断或直接由主机端发起不同 kernel 即可。一个典型例子是边界处理：本可以只让边界上的 block 走边界分支，却用全局线程号逐线程判断，结果连内部 block 的 warp 也平白多一轮。

辨认发散不用靠肉眼读码猜数据，profiler 提供两类直接证据：一是发散分支指令的计数和占比，二是每个 SM 的「warp 发射效率」（实际发射的活跃 lane 数除以理论槽位数），这个比例显著低于 1 且对得上分支位置，就是发散在出血。

六、规避手段：让 warp 内的线程尽量同构

核心思路一句话：把分支的一致性边界对齐到 warp，让同一个 warp 的 32 个线程拿到同构的数据、做出相同的判断。

按数据重排（最根本）。 先按分支关键字排序或分桶（bucket），让取值相近的元素在数组里连续分布，再启动 kernel。连续 32 个元素大概率同走一条路，发散从「每个 warp 都发散」变成「只有边界处少数 warp 发散」。这是图计算、稀疏算子的标准做法。举个数字感受收益：原始数据随机正负各半时，几乎每个 warp 都五五发散，发散区两轮、约 50% 利用率；分桶后 1024 个元素里只有两类交界处的一两个 warp 发散，其余 warp 单轮通过，发散区平均利用率从 50% 抬到接近 100%，代价是一次性的排序时间——kernel 只跑一遍未必划算，数据被反复处理多轮时几乎稳赚。

让条件对 warp 整体一致。 能用 blockIdx、tile 编号判断的，不要用 threadIdx；把不同路径拆成两个独立 kernel 各自处理全体数据，用两次 kernel 启动换两轮锁步满速执行，通常比一个发散 kernel 更快。

用 select 代替短分支。 x = cond ? a : b 这类两侧都只是简单表达式的写法，编译器会生成无分支的谓词指令（如 CUDA 的 selp），两个值都算但不串行化，两轮变一轮。注意两侧若有函数调用或访存，谓词化反而两边都付代价，不划算。

规整循环。 把变长循环改成固定迭代次数、循环体内用掩码跳过无效工作；或先做一遍压缩（compact）把有效元素挤紧，再跑满速的定长 kernel。前者的代价是所有 lane 都按最长迭代次数计费，只有无效 lane 占比不高时才划算；后者多一遍扫描，却让主 kernel 的每个 warp 全程满员，稀疏度越高、主计算越重，压缩越值得。

使用 warp 级原语。 __shfl_sync（束内 shuffle）、__ballot_sync（收集全 warp 条件位）、__activemask 等让开发者显式利用锁步性质，warp 级归约、投票、广播都可以在不发散的前提下完成。以 warp 内 32 元素求和为例，朴素写法要按奇偶分五轮、每轮都伴随掩码发散；用 shuffle 做五层蝶式相加，没有分支、没有共享内存读写，全部 lane 始终活动，是实测中又快又稳的写法。

还要澄清一个流传甚广的误解：发散的损失不等于「分支指令本身变慢」。分支指令的发射开销极小，真正的成本来自后续路径上每条指令都按 32 个 lane 的槽位计费、却只有部分 lane 产出。所以一段分支即使内部有上百条指令，只要 warp 全员同向，硬件把它当普通跳转处理，价格与 CPU 分支无异；反过来，哪怕只有一个 lane 落单，整个 warp 也要为它单独发射一轮——这也是「一个掉队者拖累全队」说法的由来。

最后留一条工程判断：不要见到分支就优化。分支结果高度偏向一边（比如 31 比 1）时硬件代价很小，重排数据带来的预处理开销可能远超收益；先用 profiler 看发散造成的实际停顿，再决定是否动手——性能优化始终要回到测量，而不是凭直觉。

这一篇在知识体系里的位置

本篇属于芯片设计进阶层，回答「SIMT 锁步执行的红利在分支面前如何折损，以及硬件和软件各自怎么补救」。向前补基础，GPU 为什么选择众核吞吐路线见 GPU 与 CPU：为什么 GPU 走向了众核吞吐；横向配套，线程如何组织成 grid/block/warp 见 CUDA 线程层次：grid、block 与 warp，warp 跑在什么样的硬件上见 SM/CU 内部结构，warp 调度器与内存合并的全景见 GPU 架构深入：SIMT、warp 调度与内存层级。再往后，内存合并（coalescing）、bank conflict 与 occupancy 调优、roofline 性能建模是这条 GPU 线的后续专题。

总结

一句话总结：SIMT 让一个 warp 的 32 个线程锁步执行同一条指令，用 32 位 active mask 管理谁在干活——条件一致时分文不取，条件不一致时 if/else 各发一轮、lane 在另一轮中空等，五五对开利用率直接腰斩到 50%，32 条路径各走各的最坏只剩约 3%，且损失只按发散区耗时占比折算；warp 等显存的数百周期延迟则靠 SM 驻留几十个 warp、调度器零开销切换来隐藏（寄存器常驻大寄存器堆、无需保存恢复），其健康度由 occupancy 衡量；软件侧的总原则是让 warp 内 32 线程同构——按数据排序分桶重排、条件对齐 block/warp、短分支改 select 谓词执行、变长循环规整化或先压缩、善用 __shfl_sync/__ballot_sync 原语，先测量真实停顿再动手优化。