Warp States
GPU 通过在多个 Warp 之间切换来隐藏指令和访存延迟。理解 Warp States 的关键不是记住一串状态名，而是理解：Warp Scheduler 每个周期都会从已经驻留的 Warp 中，寻找当前能够发射指令的 Warp。
本文讨论的是一个 SM 内部、以 Warp Scheduler（通常对应一个 SMSP）为观察范围的调度状态，而不是整个 GPU 的全局状态。
1. Warp Slot
一个线程块被分配到 SM 后，其中的 Warp 会驻留在该 SM 上，直到执行结束。每个 Warp 会被分配到某个 SMSP 的 Warp Scheduler，并占用一个 Warp Slot。在这个 Warp 的生命周期内，它不会在不同 SMSP 之间迁移。
Warp Slot 有两种基本情况：
- Unused：该 Slot 没有驻留 Warp。
- Active：该 Slot 中已经驻留一个尚未结束的 Warp。
这里的 Active 只表示 Warp 还活着并占用硬件资源，不表示它正在执行指令。Active Warp 可能正在等待数据、等待同步，也可能已经准备好发射下一条指令。
图 1：Warp Scheduler 的简化调度模型。图中的 8 个 Warp Slot 只是示意，具体数量与架构有关。图源：NVIDIA GTC 2019：CUDA Kernel Profiling Using NVIDIA Nsight Compute。
2. 状态关系
Active Warp 在某个周期内可以进一步分为：
- Stalled Warp：下一条指令当前不能发射。例如输入数据还没准备好、访存结果还没返回、正在等待同步，或者目标执行流水线暂时不可用。
- Eligible Warp：下一条指令已经具备发射条件，可以参与本周期的调度选择。
- Selected / Issued Warp：Warp Scheduler 从 Eligible Warp 中选中了该 Warp，并从中发射一条或多条指令。
它们不是几个相互独立的并列集合，关系可以写成：
Warp Slot
├── Unused
└── Active
├── Stalled
└── Eligible
├── Not Selected
└── Selected / Issued对于同一个调度器、同一个周期，有：
因此，Selected 是 Eligible 的子集。它表示本周期的调度结果，而不是 Warp 长时间保持的一种状态。Nsight Compute 的 Stall Reason 统计会把 selected 和 not_selected 也列在同一组指标中，便于让所有 Active Warp 在每个采样周期都有唯一归类；其中 selected 实际上不是停顿。
3. 每个周期如何调度
可以把 Warp Scheduler 的工作简化为以下过程：
- 检查调度池中的所有 Active Warp。
- 判断每个 Warp 的下一条指令是否可以发射。
- 不能发射的 Warp 归为 Stalled，能够发射的归为 Eligible。
- 从 Eligible Warp 中选择一个 Warp。
- 从被选中的 Warp 发射指令，然后在下个周期重新判断其状态。
例如，一个调度器中驻留了 4 个 Warp：
| 周期 | Warp 0 | Warp 1 | Warp 2 | Warp 3 | 发射结果 |
|---|---|---|---|---|---|
| Selected | Stalled | Eligible | Stalled | Warp 0 发射 | |
| Stalled | Stalled | Selected | Stalled | Warp 2 发射 | |
| Stalled | Stalled | Stalled | Stalled | 无指令发射 |
在周期 Not Selected。到周期
如果某个周期没有 Eligible Warp，调度器就无法使用该周期的 Issue Slot。偶尔出现这种情况很正常；如果长期如此，才说明可用的 Warp 并行度不足以隐藏延迟。
4. 常见 Stall Reason
Stalled 只说明“现在不能发射”，还需要结合具体 Stall Reason 判断原因。Nsight Compute 中常见的原因包括：
| Stall Reason | 含义 | 常见来源 |
|---|---|---|
long_scoreboard | 等待 L1TEX 相关操作产生的数据 | Global、Local、Texture 或 Surface Memory 的访问延迟 |
short_scoreboard | 等待 MIO 相关操作产生的数据 | Shared Memory、特殊函数或动态分支等 |
barrier | 等待同一线程块中的其他 Warp 到达 CTA Barrier | __syncthreads() 前工作量不均衡 |
wait | 等待固定延迟指令的数据依赖 | 连续存在依赖关系的算术指令 |
math_pipe_throttle | 目标数学流水线暂时不可用 | 指令集中压在同一种执行流水线上 |
mio_throttle | MIO 指令队列已满 | Shared Memory、特殊数学指令等过于密集 |
no_instructions | 等待取到下一条指令 | 指令缓存未命中或大范围跳转 |
not_selected | Warp 可以发射，但本周期选择了另一个 Warp | Eligible Warp 数量多于当前可用的 Issue Slot |
不同 GPU 架构和 Nsight Compute 版本展示的 Stall Reason 可能不同。尤其要注意，not_selected 通常不是坏事：它往往表示调度器手里还有多个可以执行的 Warp，具备较好的延迟隐藏能力。
5. 如何分析性能
不要看到 Stalled Warp 就直接判定性能有问题。几乎所有 Kernel 都会产生停顿，真正需要关注的是：停顿是否导致调度器无法持续发射指令。
可以按下面的顺序分析：
- 在 Nsight Compute 的 Scheduler Statistics 中观察调度器是否经常无法发射指令。
- 如果 Issue Slot 经常空闲，再查看 Warp State Statistics 和 Source Counters，判断哪些 Stall Reason 占主导。
- 回到对应的 CUDA 源码或 SASS，判断是访存延迟、数据依赖、同步等待，还是执行流水线压力。
- 结合吞吐率和 Occupancy 决定优化方向，而不是孤立地追求更少的 Stall。
几个典型现象可以这样理解：
- Unused Slot 很多：实际驻留的 Warp 少，可能受到线程块数量、寄存器或 Shared Memory 用量限制。
- Active Warp 很多，但 Eligible Warp 经常为 0：大量 Warp 同时等待，延迟没有被有效隐藏。
not_selected较高：通常说明有足够多的 Eligible Warp，并不意味着调度失败。long_scoreboard较高且 Issue Slot 经常空闲：更可能是访存延迟问题，应继续检查访问合并、缓存命中率和数据局部性。short_scoreboard较高且涉及 Shared Memory：应继续检查 Shared Memory 依赖和 Bank Conflict。
Nsight Compute 中以 smsp__ 开头的指标通常是在 SMSP 粒度上统计。分析这些指标时，应把它们理解为“每个 Warp Scheduler 的调度池是否有足够的 Eligible Warp”，而不是把多个 SM 或整个 GPU 的 Warp 混在一起理解。