想象一个工厂车间，有一种特殊的“三十二人班组”：32名工人并排站在流水线上，每人面前有一套相同的工位。班长站在高台上，拿着大喇叭喊口号，口号一响，32个人必须同时做完全相同的动作——不能有人快一步，也不能有人慢一步。这个“一声口号、全员同步”的班组，就是 GPU 里的 warp。GPU 中有成千上万个这样的班组，它们靠“人多力量大”把任务扛下来。可一旦某个班组里出现了“分岔”——有人要往左、有人要往右，班长就只能先带一拨人走完全程，再回来带另一拨，另一拨人只能原地干等。这个“干等”的代价，就是今天要讲的 warp divergence（分支发散）。这篇文章会带你理解发散为什么会发生、会造成多大损失，并学会用数据重排、predication（预测执行）和循环展开三种手段，让班组重新整齐划一。

岗位关联

对 AI 从业者而言，warp divergence 是 GPU 性能优化的第一道门槛。搞懂它，你就能解释为什么同一个 kernel 在某批数据上快、在另一批数据上慢；学会消除它，你的推理、训练代码才能真正吃满 GPU 算力。

目录
  1. 可视化演示
  2. Warp执行模型
  3. 发散成因
  4. 性能影响
  5. 数据重排
  6. 预测执行
  7. 循环展开
  8. 调优工具
  9. AI 应用场景
  10. Python 演示
  11. 常见误区
  12. 面试题

可视化演示

本节配有交互动效（进入视口后自动加载，约 3KB）。

Warp执行模型

车间里，班组长的喇叭每次只广播一个口号，32名工人必须在这一拍内做出同一个动作。这个“32人共进退”的班组就是 warp，工人就是线程。GPU 并不是一个线程一个线程地独立调度的，而是以 warp 为单位，把一条指令同时广播给32个线程。只要这32个线程中还有一个在忙，整个班组就要等它一起进入下一拍。这就是 SIMT（单指令多线程）的“锁步”气质：像齐步走，谁也不能抢拍。

在 CUDA 中，当你启动一个 kernel，线程会被组织成 block，而 block 内部又被自动切分成一个个 warp。绝大多数时候你不需要手动创建 warp，但你必须清楚每个线程属于哪个班组、站在哪个位置。可以通过线程编号 threadIdx 计算：

回到车间：知道了班组编号和站位，班长才能精确指挥。但班组真正的挑战不在站位，而在分岔——同一个口号，有人要往东，有人要往西。

warp_id=⌊threadIdx.x32⌋\text{warp\_id} = \left\lfloor \frac{\text{threadIdx.x}}{32} \right\rfloor
lane_id=threadIdx.x mod 32\text{lane\_id} = \text{threadIdx.x} \bmod 32
warp_id 告诉这个线程属于第几个班组，lane_id 告诉它在班组里的站位号。GPU 的很多原语（如 __shfl 洗牌指令）都依赖 lane_id 来交换数据。
  1. 整数除法定位班组threadIdx.x 除以32，向下取整，得到该线程属于第几个 warp。比如线程编号65，65÷32=2，它就在第2个班组。
    warp_id=⌊threadIdx.x32⌋\text{warp\_id} = \left\lfloor \frac{\text{threadIdx.x}}{32} \right\rfloor
  2. 取余得到站位号线程编号对32取余，得到该线程在 warp 内的 lane 编号，用于区分“我是组里第几个人”。
    lane_id=threadIdx.x mod 32\text{lane\_id} = \text{threadIdx.x} \bmod 32

发散成因

订单来了，班长喊：“奇数订单向左，偶数订单向右！”这一嗓子下去，32个人里有人迈左脚、有人迈右脚。班长为难了：喇叭只有一张嘴，不可能同时广播两条相反的口令。他只能先带奇数的人走完左边流程，再回来带偶数的人走右边流程。这个“一嗓子喊出两拨人”的局面，就是 warp divergence。

技术上的触发点是一条 if-else 条件分支。当 warp 内的32个线程执行到 if (cond) { A } else { B } 时，只要有一个线程的 cond 为真、另一个为假，硬件就让所有走 A 的线程先执行，走 B 的线程被屏蔽（inactive）；然后反过来，让走 B 的线程执行，走 A 的被屏蔽。屏蔽并不意味着不花时间，它只是让那些线程“站在工位上发呆”，整个班组的时间仍然被每一条路径串行消耗掉。理解发散的关键在于：并不是代码里有 if 就发散，而是同一个 warp 里出现了多种不同的分支方向。用数学语言描述：

其中 $b_i$ 是第 i 个 lane 的分支方向。当 $N_{\text{path}} = 1$ 时，全员同路，称为 uniform 分支，无发散；当 $N_{\text{path}} > 1$ 时，发散发生，路径们等待串行登上指令流水线。

班长后来学会了一个窍门：如果订单全都来自同一个客户，那所有订单都是奇数或者都是偶数，一嗓子就能带走全班。这说明，发散与否不仅取决于代码结构，更取决于 warp 里的数据是不是“同路人”。

Npath=∣{b0,b1,…,b31}∣N_{\text{path}} = \left| \{ b_0, b_1, \ldots, b_{31} \} \right|
  1. 收集每个 lane 的分支结果每个线程根据自己的条件产生一个分支方向，把这些方向放进一个集合。
    {b0,b1,…,b31}\{ b_0, b_1, \ldots, b_{31} \}
  2. 统计不同方向的数量集合自动去掉重复项，集合的大小就是 warp 实际需要走的路径数目。
    Npath=∣{b0,b1,…,b31}∣N_{\text{path}} = |\{ b_0, b_1, \ldots, b_{31} \}|

性能影响

发散最直接的代价是“时间账”。假设一个班组里有 N 条互斥的路径，班长必须一条一条地带。每带一条路径，全班的时间单元都在流逝，哪怕只有一小半人在真正干活。发散的耗时不是“最慢路径的时间”，而是所有路径耗时之和——因为每一条路径都真实地占据了指令发射周期，另一拨人只是被屏蔽了，并不是消失了。

如果 N 条路径耗时大致相同，都是 $T_0$，那么总耗时变为 $N \times T_0$，而理想情况下（全员同路）只需 $T_0$。也就是说，吞吐量骤降到原来的 $1/N$。例如一半线程走 A、一半走 B，A 和 B 各耗时一个周期，那么原来一个周期能完成的活，现在要两个周期，吞吐直接减半。更隐蔽的是调度器的浪费：每个 warp 耗尽执行槽之后，调度器还要承担额外的分支跳转和屏蔽检查，这些开销会进一步挤压其他 warp 的上场机会，让 GPU 整体的延迟掩盖能力变弱。

账算清楚了，班长决定改变策略。与其每次都眼睁睁看着班组分成两拨发呆，不如在开工之前就把订单分好类——这就是数据重排。

Twarp=∑i=1NTiT_{\text{warp}} = \sum_{i=1}^{N} T_i
TuniformTdiverged=T0∑i=1NTi≈1N\frac{T_{\text{uniform}}}{T_{\text{diverged}}} = \frac{T_0}{\sum_{i=1}^{N} T_i} \approx \frac{1}{N}
  1. 对每条路径求和warp 串行执行 N 条不同的路径，总时间等于每条路径耗时的累加。
    Twarp=∑i=1NTiT_{\text{warp}} = \sum_{i=1}^{N} T_i
  2. 等耗时特例如果每条路径耗时都相同 $T_0$，则总时间为 $N \times T_0$，与不分发相比效率降为 $1/N$。
    Twarp=N×T0T_{\text{warp}} = N \times T_0

数据重排

车间主任提出一个大胆的方案：以后订单进车间之前，先过一道“分拣线”。所有奇数订单进左筐，偶数订单进右筐，然后再把整筐订单交给班组。这样一来，一个班组拿到的32份订单几乎全是同一类，班长喊“奇数向左”时，32个人齐刷刷向左走，没有任何人发呆。这个思路在 GPU 里叫数据重排（data rearrangement）。

实现时，算法通常会根据某个特征值把数据映射到有限个“桶”里，桶的编号对应着分支方向。然后对数据按桶号排序，让同一个桶的数据在内存中连续排布。这样，当一个 warp 加载数据时，它大概率拿到的都是同一个桶里的数据，分支自然 uniform 化。分桶函数可以写成：

其中 key(x) 是数据 x 的分类特征，B 是桶的总数。排序后，内存中数据的排列大致为“桶0、桶1、桶2……”，warp 按顺序取数时，内部数据类别的一致性大幅提高。

不过主任也提醒：分拣线本身要花钱——排序和拷贝也在消耗 GPU 时间。如果每个数据每帧都在变，排序成本可能超过发散省下的时间。所以数据重排适合“分类稳定、类别少、发散代价大”的场景。

bucket(x)=hash(key(x)) mod B\text{bucket}(x) = \text{hash}(\text{key}(x)) \bmod B
  1. 计算特征哈希对每个数据的 key 计算哈希值，把复杂的分类信息压缩成一个整数。
    hash(key(x))\text{hash}(\text{key}(x))
  2. 取模分桶哈希值对桶数 B 取模，得到 0 到 B-1 之间的桶号。
    bucket(x)=hash(key(x)) mod B\text{bucket}(x) = \text{hash}(\text{key}(x)) \bmod B
  3. 排序归拢所有数据按桶号排序，同桶数据在内存中连续，warp 取数时更可能整组落入同一桶。

预测执行

主任又带来一种新设备：电子工牌。每个工人的工牌上有一个开关，开关由订单决定。班长不再喊“向左走还是向右走”，而是喊“向前三步！”——开关打开的工人真走三步，开关关闭的工人保持立正，假装走了三步。所有人同时做完整个口令，不存在先带一拨再带一拨。这个“人人带开关、同做一套动作”的方案，就是预测执行（predication）。

在 CUDA 中，编译器遇到很短的 if-else 时，可能不会生成真正的分支跳转，而是把两条路径的指令都保留下来，每条指令前挂一个谓词（predicate）。指令发射给整个 warp 时，每个 lane 根据自己手上的谓词决定“是否把结果写回”。不满足条件的 lane 会丢弃计算结果，但它依然参与了取指和发射，所以两条路径的指令都被执行了一遍。predication 的优势是省去了跳转和路径切换的开销；代价是一条路径的线程在另一条路径上“空转弯”。代价模型可以表示如下：

当 $T_{\text{pred}} < T_{\text{div}}$，即跳转开销 $O_{\text{jump}}$ 足够大、且两条路径都很短时，predication 更划算。

工牌方案不是万能的——如果左边是一条十米长廊，右边是一条三步小道，让所有人都走完十米长廊显然太亏。所以 CUDA 编译器只对指令条数很少的分支自动启用 predication，长分支仍要靠数据重排让路径 uniform。

Tpred=TA+TBT_{\text{pred}} = T_A + T_B
Tdiv=TA+TB+OjumpT_{\text{div}} = T_A + T_B + O_{\text{jump}}
  1. predication 的总耗时两条路径的指令全部执行一次，时间为 A 路径耗时加 B 路径耗时。
    Tpred=TA+TBT_{\text{pred}} = T_A + T_B
  2. 发散的总耗时串行执行两条路径之外，还要加上分支跳转和路径切换的额外开销。
    Tdiv=TA+TB+OjumpT_{\text{div}} = T_A + T_B + O_{\text{jump}}
  3. 选择条件当 predication 的耗时小于发散耗时，编译器就倾向用谓词指令；否则保留真实分支。

循环展开

班长发现，喊口号本身也有成本。每喊一句“抬起左脚”，都要先清一下嗓子、确认没人掉队。如果一个班组要连续做100次相同的短动作，班长就要喊100次口号，检查100次。能不能把“第1步到第10步”合成一句长口号，一次喊完？能，这就是循环展开（loop unrolling）。

在 GPU kernel 中，循环会产生两类开销：循环体本身的计算指令，以及循环控制指令——判断 i 是否到达边界、跳回循环头、维护计数器。控制指令虽然不直接产生业务结果，却同样占据 warp 的指令发射周期。把循环体复制 K 份后，循环控制执行的次数大约降为原来的 $1/K$。CUDA 提供了 #pragma unroll 指令，让编译器展开短循环；也可以手工展开，把循环体重复写几遍。

其中 $T_{\text{body}}$ 是循环体耗时，$C$ 是每轮循环的控制开销，$K$ 是展开因子。展开后控制开销从 $N \times C$ 降到约 $N/K \times C$。与此同时，更少的控制流指令也意味着更少的潜在分支点，warp 遇到分支的概率随之下降。

回看车间：合并口号省去了无数次清嗓子，工人动作不会变快，但“检查”的时间被压缩了。这是循环展开的哲学——不减少干活量，只减少“指挥”的开销。

Tloop=N×Tbody+N×CT_{\text{loop}} = N \times T_{\text{body}} + N \times C
Tunrolled=N×Tbody+⌈NK⌉×CT_{\text{unrolled}} = N \times T_{\text{body}} + \left\lceil \frac{N}{K} \right\rceil \times C
  1. 原循环总开销N 次迭代的业务指令加上 N 次循环控制，控制开销按次累加。
    Tloop=N×Tbody+N×CT_{\text{loop}} = N \times T_{\text{body}} + N \times C
  2. 展开后控制开销循环体被复制 K 份，循环控制次数降为 $\lceil N/K \rceil$ 次。
    Tunrolled=N×Tbody+⌈NK⌉×CT_{\text{unrolled}} = N \times T_{\text{body}} + \left\lceil \frac{N}{K} \right\rceil \times C
  3. 展开因子 KK 越大，检查次数越少，但代码体积增大、寄存器压力上升，过犹不及。

调优工具

前面讲了这么多理论，落到实际开发中，班长的第一件事不是盲改代码，而是用“摄像机”复盘车间。NVIDIA 的 Nsight Compute（ncu）就是这台摄像机：它能逐 kernel 分析 warp 的完整状态，告诉我某个 kernel 有多少 warp 因为分支停滞、多少周期花在了等待不同路径上，还能展示分支覆盖率、每条指令的活跃线程占比等关键指标。有了这些数据，优化不再是玄学，而是“哪里发散打哪里”。

优化手段的选择也有章法可循。先定位热点 kernel——用 profiler 找出耗时占比最高的函数；再看它的发散指标，如果 branch divergence 相关的占比很高，接着分析数据分布：类别是否稳定、分桶成本是否可接受。分支短则尝试 predication，分支长则数据重排，涉及循环边界发散则循环展开。每一种方案都应以实验数据验证，最终的判断标准是真实的加速比：

班长的摄像机让他看清了每一分钟浪费在哪里。发散、重排、predication、循环展开，这四样工具组合起来，才是一个成熟的 GPU 工程师的完整工具箱。

speedup=TbeforeTafter\text{speedup} = \frac{T_{\text{before}}}{T_{\text{after}}}
  1. 优化前耗时用 profiler 测出 kernel 修改前的执行时间。
    TbeforeT_{\text{before}}
  2. 优化后耗时应用数据重排、predication 或循环展开后重新测量。
    TafterT_{\text{after}}
  3. 计算加速比两数相除，比值大于1说明优化有效，越大越好；小于1说明优化得不偿失，应回退方案。
    speedup=TbeforeTafter\text{speedup} = \frac{T_{\text{before}}}{T_{\text{after}}}

AI 应用场景

LLM推理请求的分桶填充

阶段：模型推理引擎优化（inference serving）

解决的问题：大语言模型推理时，不同请求的序列长度、注意力掩码各不相同。GPU 批处理中，掩码判断和 padding 逻辑会产生大量 if 分支，同一个 warp 内既有有效 token 又有 padding token，warp divergence 导致计算单元一半空闲。AI 工程师需要让 SIMT 流水线满负荷运转。

具体案例：使用 vLLM 一类推理框架时，连续批处理会把多个请求拼成一个 batch 喂给 GPU。长度差异大的请求混在一起，attention 计算中遇到 if (mask > 0) 这样的条件，warp 内有效与无效线程分道扬镳。工程上按序列长度把请求分桶，同一桶内的请求 padding 到相近长度再进 GPU，从源头保证一个 warp 内数据的“同路性”。

图像阈值分类的直方图分桶

阶段：计算机视觉预处理（CV pipeline）

解决的问题：Canny 边缘检测等算法会对像素梯度做多级阈值判定——强边缘、弱边缘、非边缘三条路径。GPU 上直接实现会产生三路分支发散，图像特征分布越均匀，发散越严重。

具体案例：某自动驾驶感知团队发现边缘检测 kernel 耗时异常。先用 ncu 定位到 warp divergence 指标偏高，随后在预处理阶段计算梯度直方图，将像素按梯度幅值分桶，使每个 warp 处理的像素尽量落在同一阈值区间，发散率下降 70%，流水线整体延迟降低 18%。

粒子模拟的类型排序法

阶段：科学计算与物理仿真（HPC）

解决的问题：粒子系统中，刚体、流体、布料粒子对应不同的物理方程。未排序时 warp 内粒子类型混杂，每个分支路径都要串行执行，模拟步进时间被成倍拉长。

具体案例：在 CUDA 分子动力学模拟中，按粒子类型排序后每个 warp 处理同一类粒子，分支几乎全部 uniform。排序代价通过基数排序控制在建树阶段一次完成，后续每步模拟的分支发散率从 35% 降到 2%，整体模拟速度提升 1.6 倍。

Python 演示

python36 行
import random WARP = 32  def simulate(data, labels, costs, predication=False):    n_warp = len(data) // WARP    total = 0    for w in range(n_warp):        chunk = labels[w * WARP:(w + 1) * WARP]        if predication:            total += sum(costs.values())        else:            total += sum(costs[p] for p in set(chunk))    return total  random.seed(42)n = 65536data = list(range(n))labels = [random.randint(0, 1) for _ in range(n)]costs = {0: 1, 1: 2} before = simulate(data, labels, costs) pairs = sorted(zip(data, labels), key=lambda x: x[1])data_ok = [d for d, _ in pairs]labels_ok = [l for _, l in pairs]after = simulate(data_ok, labels_ok, costs) pred = simulate(data, labels, costs, predication=True) print(f"divergence before reorder: {before}")print(f"after data reorder:        {after}")print(f"predication always:        {pred}")print(f"speedup by reorder:        {before / after:.2f}x")
目的：模拟一个包含 65536 个元素的 GPU 任务，每个元素只有 0/1 两种分类，处理耗时分别为 1 和 2 个周期，对比数据重排前后的 warp 发散总耗时。
  1. chunk = labels[w * WARP:(w + 1) * WARP]每 32 个元素形成一个 warp，取当前 warp 内元素的分类标签
  2. sum(costs[p] for p in set(chunk))用集合去掉重复路径后累加各路径耗时，模拟发散串行执行
  3. labels = [random.randint(0, 1) for _ in range(n)]生成随机分布的分类，保证初始发散严重
  4. pairs = sorted(zip(data, labels), key=lambda x: x[1])按标签排序完成数据重排，让同类数据相邻
  5. after = simulate(data_ok, labels_ok, costs)重排后每个 warp 内几乎只剩一种标签，耗时为单路径
  6. pred = simulate(data, labels, costs, predication=True)模拟 predication 总是执行 0 和 1 两条路径
  7. print(f"speedup by reorder: {before / after:.2f}x")输出重排优化带来的真实加速比

常见误区

误区我觉得 GPU kernel 里只要写了 if-else，性能就一定完蛋，所有条件分支都应该尽早删掉。

正解发散与否取决于同一个 warp 内 32 个线程的分支方向是否一致。如果条件对所有线程相同，属于 uniform 分支，硬件像执行普通指令一样直接通过，没有任何额外开销。比如 if (blockIdx.x == 0) 这种按 block 划分的条件，整个 warp 天然一致。真正的敌人是“数据相关”的分支——同一 warp 里不同线程拿着不同数据，才可能走向不同路径。

为何常见：把 CPU 上分支预测的经验直接套到 GPU 上。CPU 关心跳转是否预测准确，GPU 关心 warp 内步调是否一致，两者机制完全不同。

误区既然 GPU 有几千个核心，那 warp 分组后各走各的路径，不是正好并行吗？

正解一个 warp 是由一个调度器在一个周期内广播同一条指令的，warp 内的 32 个线程共享指令流，不存在“两条路径同时发射”的硬件机制。发散发生时，路径必须在时间上串行排队：先执行路径 A 的指令，再执行路径 B 的指令。并行发生在不同 warp 之间——多个 warp 可以同时被不同调度器处理，但同一个 warp 内部只能锁步前进。

为何常见：混淆了“多 warp 并行”和“warp 内锁步执行”两个层次的并行模型。前者是 GPU 并行性的来源，后者是发散代价的来源。

误区predication 不跳转、不分流，听上去完美无缺，那我把所有分支都改成 predication 不就彻底没有发散了吗？

正解predication 的本质是“两条路径的指令全部执行一遍，只是符合条件的 lane 才写回结果”。这意味着它要付出路径 A 和路径 B 的完整指令开销。对于很短的 if-else（一两句赋值），predication 非常划算；但对于长达几十条指令的分支，predication 会让所有线程都走完两条长路径，比串行发散更糟。CUDA 编译器只在分支很短时自动启用 predication，长分支请改用数据重排。

为何常见：看到文档里“predication 可以避免 divergence”就当作免费午餐，忽略了它“两条路径都要走”的本质代价。

误区数据重排消除发散效果立竿见影，那无论什么任务，先按分支标签排序总没错。

正解数据重排本身包含排序、拷贝、索引重映射的开销，这些也消耗 GPU 带宽和计算资源。如果数据每帧都在剧烈变化、分类标签不稳定，排序成本可能超过发散省下的时间。实践上要先用 ncu 量化发散损失，再估算重排开销，只有 排序成本 < 发散损失 时才值得做。一般适用于“分类稳定、类别少、发散路径耗时长”的算法。

为何常见：只看重排后分支路径变少，忽略了排序本身的时间和空间开销，犯了“局部最优”的直觉错误。

误区GPU 靠大量 warp 切换隐藏延迟，就算一个 warp 因为发散变慢了，其他 warp 也能顶上，整体性能应该不受影响。

正解warp 切换确实能隐藏延迟，但每个 warp 发散会占据调度器的发射槽位，即使它被切走，它之前的发散周期也真实消耗了指令发射机会。当一个 kernel 中大量 warp 都发散时，调度器反复发射无效的屏蔽指令，全局执行效率被系统性拉低。极端情况下，所有 warp 同时发散会让 SM 的指令吞吐降到理论值的几分之一，没有任何“别的 warp”能补回来。

为何常见：把延迟隐藏想象成无限资源，忽略了调度器和发射带宽本身就是共享的稀缺资源。

面试题

Q1. Warp 是 GPU 中怎样的执行单位？

单选题 · easy

  • A.存储单位，包含 256 字节缓存
  • B.执行基本单位，包含 32 个线程← 正确答案
  • C.调度单位，包含 8 个线程
  • D.网络单位，连接 64 个核心

答案：B。warp 是 GPU 调度的基本执行单位，固定包含 32 个线程，它们在同一指令流上锁步执行。选项 A 说的是缓存，C 和 D 的数字都不正确，只有 B 准确描述了 warp 的核心特征。

Q2. 下列哪种情况一定会发生 warp divergence？

单选题 · easy

  • A.32 个线程全部执行 if 分支
  • B.部分线程执行 if 分支，部分执行 else 分支← 正确答案
  • C.32 个线程全部执行 else 分支
  • D.所有线程执行同一个 for 循环

答案：B。发散的本质是 warp 内出现了不同的分支方向。部分线程走 if、部分走 else，路径集合大小为 2，必然发散。A 和 C 是全员同路，D 是无条件循环，都不会发散。

Q3. 当 warp 内线程走向不同路径时，GPU 硬件如何执行？

单选题 · mid

  • A.并行执行所有路径，互不干扰
  • B.随机选择一条路径执行
  • C.串行执行各路径并屏蔽非活跃线程← 正确答案
  • D.直接终止走少数路径的线程

答案：C。硬件按路径逐一发射指令，执行路径 A 时屏蔽走路径 B 的线程，再交换角色执行路径 B。两条路径串行占用指令周期，这就是发散性能损失的来源。A 错误在于同一 warp 无法并行多路，B 和 D 完全不符合 SIMT 模型。

Q4. 若 warp 中一半线程走路径 A、一半走路径 B，且 A 耗时为 B 的 2 倍。与全部走 A 相比，发散后总耗时增加约多少？

单选题 · mid

  • A.不增加
  • B.增加 33%
  • C.增加 50%← 正确答案
  • D.增加 100%

答案：C。设 B 耗时 1 个周期，A 耗时 2 个周期。全部走 A 耗时 2 个周期；发散后总耗时 = A + B = 3 个周期。耗时从 2 变到 3，增加了 50%。注意发散是路径耗时相加，而不是取最大值。

Q5. predication（预测执行）的核心代价是什么？

单选题 · mid

  • A.只执行 if 分支的指令
  • B.两条路径指令都发射，无效 lane 的结果被丢弃← 正确答案
  • C.不执行任何计算，直接跳过
  • D.只对长分支有效，短分支无效

答案：B。predication 把所有路径的指令都发射给整个 warp，利用谓词决定哪些 lane 写回结果。代价是路径 A 和 B 的指令都占用执行周期，因此它只适合很短的分支。A 和 C 不正确，D 说反了——长分支恰恰不适合 predication。

Q6. 以下哪些手段可以消除或缓解 warp divergence？

多选题 · mid

  • A.数据重排，让同一 warp 内数据同路← 正确答案
  • B.对短分支使用 predication← 正确答案
  • C.把线程数从 256 增加到 1024
  • D.循环展开，减少循环分支点← 正确答案

答案：ABD。数据重排从数据侧消除路径分歧；predication 用谓词机制绕过分支；循环展开减少循环控制分支的出现频率。增加线程数只是多了 warp，并不会让已有 warp 内部的分歧消失，所以 C 不选。

Q7. 关于 predication 与分支发散，下列说法正确的有？

多选题 · hard

  • A.predication 会执行所有路径的指令← 正确答案
  • B.predication 通常只适合很短的 if-else← 正确答案
  • C.长分支用 predication 比分支更高效
  • D.分支发散时各路径串行执行← 正确答案

答案：ABD。predication 的本质是“两径都算、按谓词写回”，所以 A 正确；因为要付出全部路径的指令代价，它只适合短分支，B 正确、C 错误；分支发散时路径串行排队执行，D 正确。

Q8. 循环展开的主要收益是什么？

单选题 · easy

  • A.减少循环控制和分支指令数量← 正确答案
  • B.让 warp 包含更多线程
  • C.消除所有条件分支
  • D.提高 GPU 时钟频率

答案：A。循环展开把循环体复制多份，使循环控制从 N 次降为约 N/K 次，也减少了循环边界产生的分支点。它不改变 warp 大小，也不消除所有分支，更不影响时钟频率，所以 A 正确。

Q9. 数据重排优化发散的原理是什么？

单选题 · mid

  • A.增加分支路径的数量，分摊压力
  • B.让同一个 warp 内的数据尽量属于同一分支路径← 正确答案
  • C.减少线程数量，降低调度开销
  • D.把所有分支都改成循环

答案：B。数据重排通过分桶和排序，把相同分类的数据在内存中连续排布，warp 取数时整组落入同一路径，分支从发散变为 uniform。A 南辕北辙，C 和 D 与重排原理无关。

Q10. 使用 NVIDIA Nsight Compute 进行发散调优，合理的顺序是？

单选题 · hard

  • A.直接改代码，逐行猜发散位置
  • B.定位热点 kernel→看发散指标→分析数据分布→选优化手段← 正确答案
  • C.只要把 if 全删掉就能解决问题
  • D.先做数据重排，再确认有没有发散

答案：B。科学的调优流程是先找出耗时最长的 kernel，用 ncu 量化发散贡献率，再结合数据特征选择 predication、数据重排或循环展开，最后用加速比验证。A、C 是凭感觉优化，D 颠倒了“先诊断后治疗”的顺序。