深度解析NVIDIA GPU任务调度：从Block分发到Warp发射的完整链路

这是一篇经验总结性质的技术分析。平时我们写 CUDA 程序，习惯把注意力放在 kernel 实现、显存带宽、共享内存 bank conflict 这些“显性”点上，调度这层往往被当成黑盒。但当你开始追性能瓶颈，尤其是 kernel 启动开销、尾效应、SM 占用率上不去、warp 停滞这类问题时，就会发现：没有对任务调度模型的理解，很多优化手段就是死记硬背，换个卡、换个 kernel 就不会用了。

这篇博客是“NVIDIA 调度分析”系列的第二篇，上一篇我们拆了 GPU 的硬件架构和执行流水线，这一篇专门聊调度模型。我会从宏观到微观，把 GPU 上的“任务”从 CPU 端发出到 SM 内部 warp 被发射的完整链路讲清楚，然后落到 CUDA 编程里我们能感知、能调控的部分。

1. 调度模型的整体分层

一上来先把“调度”这个词拆开。GPU 不是一块单一处理器，它内部有好几级执行单元，每一级的“任务”粒度完全不同。NVIDIA 的调度实际上是一个多级、多粒度的队列系统，每一级解决不同的问题。

从用户程序到硬件执行，任务大致经过 4 层：

• CPU/API 层：程序发起 kernel 启动、数据拷贝、同步等操作，通过驱动把这些操作转成 GPU 指令流。
• GPU 前端：GPU 上的 Work Distributor（工作分配器）接收指令流，把 kernel 的线程块（block）分发到各个 GPC（Graphics Processing Cluster）上的 SM。
• SM 内部线程块调度：每个 SM 维护一组 block 槽位，决定哪些 block 可以常驻、哪些排队等资源。
• Warp 调度器：SM 内真正执行指令的单元，从多个 warp 中挑选可发射的 warp，把指令发到执行流水线。

举个例子帮助理解：想象一个大型餐饮厨房。CPU 是老板，负责接外面的订单（启动 kernel）。GPU 前端是前台主管，订单进来后，他把每一桌的菜（线程块）分给不同的灶台（SM）。每个灶台同时能做几道菜有限（block 槽位），菜太多就排队。而灶台上具体哪个灶眼现在炒什么菜（warp 发射），取决于炒菜师傅（warp scheduler）的实时判断——有的菜在等食材（访存），有的菜在等锅（计算单元），师傅就先去炒那些能动的菜。

这个类比虽然粗，但抓住了核心：GPU 调度的高度并行性来自多个独立决策点的流水线配合，而不是中央调度器集中控制。

在 NVIDIA 的不同架构上，这几层的实现细节有差异，但整体框架从 Fermi 开始就基本定型，后续各代（Kepler、Maxwell、Pascal、Volta、Turing、Ampere、Hopper、Ada）只是在此框架上调整硬件宽度、队列深度、调度策略和资源管理方式。

2. 宏观调度：从 CPU 到 GPU 的网格分发

2.1 command queue 与 GPU 前端队列

我们写 CUDA 程序时调用kernel<<<grid, block>>>，这不会立即在 GPU 上执行，而是被放入一个 command queue（命令队列）。这个队列一般按 stream（流）组织。同一个 stream 里的命令按顺序执行，不同 stream 的命令可以并发或交错执行。

驱动把队列里的命令打包成 GPU 能理解的指令，通过前端总线发送到 GPU。GPU 收到后，由front-end解析命令类型，然后交给对应的硬件单元。对于 kernel launch 命令，真正干活的是Work Distributor。

这里有个常见误区：很多初学者以为kernel<<<...>>>后面的 CPU 代码会等 kernel 跑完才继续。实际上 kernel 启动是异步的，CPU 只是把命令丢进队列就返回了。同步行为（如cudaDeviceSynchronize、cudaMemcpy隐含的同步）才会阻塞 CPU。

从架构演进来看，Kepler 时代引入的Hyper-Q是一个关键节点。此前 Fermi 只有单一硬件工作队列，多个 CPU 线程同时提交任务时，会互相阻塞、串行化，GPU 利用率很差。Hyper-Q 把工作队列拆成 32 个硬件队列，对应 32 个 stream，多个 CPU 线程各用各的队列，就可以并发提交、并发调度了。

好消息是，我们用 CUDA 编程时不需要直接操作 Hyper-Q 队列，但理解这一层对提升多流并发性能很重要。实际编码中我经常看到有人创建了一堆 stream，结果根本没有多线程提交，或者 buffer 依赖没有处理好，导致 stream 串行，白费了硬件并发能力。

2.2 GigaThread Engine 与 Thread Block 分发策略

真正在 GPU 内部负责把线程块分发到 SM 的，是GigaThread Engine。它的职责不止是分发，还管理着线程块的生老病死：分配 block ID、为 block 分配 SM 资源（寄存器文件、共享内存）、记录 block 完成、处理 block 之间的依赖（比如 grid sync）。

那么 GigaThread Engine 怎么决定把 block 发给哪个 SM？

NVIDIA 没有公开 GigaThread Engine 的完整分发算法，从公开资料和逆向工程分析来看，策略基本是“轮询 + 资源感知”的混合。它的核心逻辑是：只要某个 SM 还有空闲的 block 槽位，且该 SM 的剩余资源（寄存器、共享内存）能满足新 block 的需求，就把 block 分发过去。如果所有 SM 都满了，block 就留在 Work Distributor 的队列里等待。

这个设计意图非常清晰：均衡。轮询保证块在 GPC/SM 间大致均匀分布，避免某些 SM 忙死、某些闲死。资源感知则避免分发一个 block 过去结果因寄存器不够又挤掉别的 block 的尴尬。

不过实践中有个值得注意的现象：当 grid 规模较小、block 数量和 SM 数量不完全匹配时，block 分布可能不均匀。比如一个 GPU 有 18 个 SM，而你的 grid 只有 16 个 block，那必然有 2 个 SM 闲置，整体利用率只有 88.9%。这就是所谓的tail effect（尾效应）。如果每个 SM 只能同时跑 1 个 block（block 很大、资源占用高），尾效应会更明显。

所以，工程上的第一个实在建议是：block 数量最好明显多于 SM 数量，既保证均衡负载，也缓冲 block 完成时间的差异。通常我建议 block 总数至少是 SM 数量的 2～3 倍，理想情况是 4～8 倍。

// 一个简单的 tail effect 演示：grid 大小固定时性能可能波动 int smCount; cudaDeviceGetAttribute(&smCount, cudaDevAttrMultiProcessorCount, deviceId); // 假设我们的 block 数等于 SM 数，理论上每个 SM 分到 1 个 block // 但实际上 block 执行时间不可能完全相同，总有快的、有慢的 // 慢 block 会让所在 SM 继续工作，快 block 所在 SM 已经空闲 // 于是整体执行时间被最慢的那个 block 决定 // 解决思路：让 grid = smCount * blocksPerSM * k（k 为整数倍率）

Grid 规模选择属于“看得见摸得着”的调度调优，几乎每次性能分析都会遇到。

2.3 streaming multiprocessor 上的 resident block 管理

Block 到达 SM 之后，并不是立即就能执行，它必须先“入住”SM。SM 会为 block 分配：

• block 槽位（每个 SM 最多可常驻的 block 数，不同架构不同，一般是 16～32 个）
• warp 槽位（每个 block 包含若干个 warp，SM 需要给这些 warp 分配执行上下文）
• 寄存器空间（warp 的寄存器文件在 SM 物理上是独享的，block 的每个 warp 都要占一组寄存器）
• 共享内存（block 内的线程通过__shared__访问的存储空间）
• 同步与屏障资源（如 barrier、named barrier 等）

只有这四类资源全部满足，block 才能进入“已调度”状态，也就是说它的 warp 才有可能被 warp scheduler 选中发射指令。

这里的核心矛盾是：这些资源的总量是有限的，block 之间是竞争关系。比如一个 SM 有 65536 个 32 位寄存器，如果每个线程用 64 个寄存器、一个 block 有 256 个线程，那么这个 block 要占 16384 个寄存器，SM 最多同时放 4 个这样的 block。如果你把每个线程的寄存器用量提高到 128，那一个 block 就要 32768 个寄存器，SM 最多放 2 个 block。寄存器用得越多，能同时常驻的 block 越少，能调度隐藏延迟的 warp 就越少。

共享内存同理。Ampere A100 每个 SM 有 163KB 共享内存（可配置），如果你的 block 要用 80KB 共享内存，那一个 SM 只能放 2 个 block。把共享内存砍到 40KB，就能放 4 个 block。

很多性能问题追根溯源，都是资源占用把并发度锁死了。这也是为什么 Nsight Compute 的 Occupancy 分析页面会把“理论占用率”和“实际占用率”分开列出来——理论是纯硬件上限，实际是资源限制后的结果。两者之间的差距，就是你的资源使用策略造成的。

一个具体的例子：

某个 kernel 每个 block 有 512 线程，每个线程 40 个寄存器。

• 512 × 40 = 20480 个寄存器/block
• SM 寄存器文件 65536，65536 / 20480 = 3.2，向下取整 = 3 个 block
• 每个 SM 最多 32 个 block（架构限制），3 < 32，所以 block 数量不是瓶颈
• block 总线程数 = 3 × 512 = 1536 线程
• 如果 SM 最大线程数为 2048，则占用率 = 1536 / 2048 = 75%

同样的 kernel，如果把寄存器压到 32 个：

• 512 × 32 = 16384，65536 / 16384 = 4，可以放 4 个 block
• 总线程数 = 4 × 512 = 2048，占用率 100%

多 25% 的并发 warp，在访存密集场景下往往意味着 20%+ 的性能提升。这就是寄存器分配和调度的耦合关系——后面第 4 节我们详细算。

3. 微观调度：Warp 的诞生与发射

3.1 Thread 到 Warp 的固定映射

Block 成功入住 SM 后，硬件会把 block 内的线程按连续 32 个一组，组成warp。这个映射是固定的、线性的：thread 0～31 是 warp 0，32～63 是 warp 1，依此类推。你无法改变这个映射关系，但可以预判它：如果 block 大小是 128，那 warp 0～3 分别对应 thread 0～31、32～63、64～95、96～127。

这个固定映射有两层直接影响。

第一，分支发散（divergence）的经济代价是 warp 级别的。同一个 warp 里，如果线程走了不同分支（if/else），硬件会先执行一个分支路径，再执行另一个，两条路径都执行完才算这个 warp 完成这条指令。对应到我们编程：尽量让分歧发生在 warp 边界上。比如判断if (threadIdx.x >= 32)，如果 block 大小为 64，线程 0～31 走 if，32～63 走 else，两个 warp 完全独立，不会产生 divergence，尽管代码本身有分支。

第二，共享内存 bank conflict、全局内存访问的合并性都取决于 warp 内 32 个线程的访存地址分布。固定映射意味着“哪些线程会在同一个 warp 里”是确定的，你可以据此组织数据结构。

3.2 Warp Scheduler 与 Instruction Dispatch

Warp 进入 SM 后，就成了warp scheduler眼中的调度单位。以 Ampere GA100 的 SM 为例，一个 SM 分为 4 个处理分区（processing block），每个分区有自己的：

• 1 个 warp scheduler
• 一组执行单元（FP32、INT32、FP64、SFU、Tensor Core、LD/ST 单元的集合）
• 一组寄存器文件切片

每个 warp scheduler 管理大约 16 个 warp（具体数值与架构相关，比如 Volta/Turing/Ampere 是 16，Kepler/Maxwell 是 8）。

Warp scheduler 的职责是：每个时钟周期（cycle）挑选一个或多个可以发射的 warp，把当前指令送到执行单元。

什么叫做“可以发射”？就是这条指令的操作数已经准备好、执行单元有空位、没有结构性阻塞（比如前一条指令还在用同一个单元写回、内存指令的队列满了、barrier 尚未到达等）。

不同代际的调度器发射能力不同：

• Fermi/Kepler/Maxwell/Pascal（Volta 之前）：每个 scheduler 每周期最多发射 1 条指令，称为 single-issue。
• Volta/Turing/Ampere/Hopper：每周期可以发射 2 条不同的指令（dual-issue），前提是两条指令属于同一个 warp 且目标执行单元不同（比如一条 INT、一条 FP32），或者属于不同 warp。

双发射策略给优化带来新的思考：如果 kernel 由纯 FP32 指令构成，调度器就算有双发射能力，执行单元也只有一个 FP32 流水线（每个分区），结果还是每周期 1 条 FP32 指令。真正的福利来自指令混合：比如 FP32 和 INT 混着写、特殊函数和普通算术混着写，让不同执行单元都在工作。

3.3 Warp 状态机：停滞、就绪与发射

Warp scheduler 内部维护每个 warp 的状态，本质上是一个有限状态机。工业界常说的 warp 状态包括：

• Running（执行中）：该 warp 正有一条指令在流水线中执行。这不是持续状态，只是某一条指令从发射到写回的窗口。
• Ready（就绪）：warp 的下一条指令已经就绪，操作数可用，等待发射。就绪 warp 越多，调度器的选择空间越大。
• Stalled（停滞）：warp 当前无法继续执行。原因可能是等待全局内存返回（最典型）、等待 shared memory 取数、等待 barrier、等待固定延迟的算术指令链、目标执行单元忙碌等。
• Dormant/Inactive（休眠/不可见）：block 尚未启动完成，或已被标记结束，warp 不参与调度。

调度器每周期从 Ready 集合中挑一个/两个 warp。挑选策略 NVIDIA 没有公布具体细节，一般认为是简单的优先级轮询（round-robin），保证公平性。但在实践层面，我们不需要精确模拟硬件调度，只需要理解核心结论：Ready warp 越多，隐藏延迟的能力越强。

有个形象的比方：warp scheduler 就像交通信号灯下的警察，面前有几十辆车。有车的车道（Ready warp）越多，他每分钟能放行的车辆总数越多。如果一半车道是空的（warp 被停滞），整体通行效率就下来了。

那么怎么量化“多少个 Ready warp 够用”？这就进入延迟隐藏的计算了。

3.4 延迟隐藏的数学直觉

设一条指令从发射到结果可用的延迟为 L 个周期。如果每个周期发射一条独立指令、且这些指令之间没有任何数据依赖，我们需要 L 条独立的指令同时“在飞”，才能在结果回来之前持续发射新指令。

但一个 warp 本质上是一条指令流，前后指令往往有依赖（比如a = b + c; d = a * e;，第二条要等第一条的 a）。所以单个 warp 很难独立提供 L 条并行指令。这时候就需要多个 warp 来凑。

一般经验值是：隐藏一个 HBM 全局内存访问的延迟，大约需要 12～20 个独立的 Ready warp（具体取决于访问是 L2 命中还是 DRAM 命中）。

举例计算：

• 假设内存延迟 400 周期（这是一个典型量级，A100/H100 的 DRAM 延迟大约在 300～500 周期，视访问模式而定）。
• 假设每个 warp 发射一个独立的 128 字节内存访问后，就停滞等待数据返回。
• 那么我们需要大约 400 条独立的 load 指令同时在飞，才能每周期发射一条新 load。
• 每个 SM 最多 64 个 warp（2048 线程 / 32），每个 warp 每周期最多发射 1 条 load；4 个 scheduler 并发，理论峰值每周期可以发射 4 条 load。
• 400 / 4 = 100，但每个 warp 只能发一条 load 然后停，所以至少需要 400 条 load 指令，而只有 64 个 warp，也就是平均每个 warp 要发射 6.25 条独立 load（不依赖前序 load 的）才能填满。

这还只是 load 一条指令，实际 kernel 中还有算术、store、分支，情况更复杂。结论是：哪怕 100% 占用率、64 个 warp 全在，如果每个 warp 只做一次访存就等结果，延迟依然无法完全隐藏。这就引出了我们在编程时要追求“每个 warp 内有尽可能多的、独立的访存指令”这个优化方向——比如循环展开、多个独立 load 变量并行。

4. 资源占用与 Occupancy 的关系

4.1 Occupancy 不是越高越好

说一个可能反直觉的结论：occupancy 越高不一定性能越好。

Occupancy 的定义是“SM 上活跃 warp 数与最大 warp 数的比值”。高 occupancy 通常意味着更强的延迟隐藏能力，但换来的是每个 warp 更少的寄存器、更多的指令缓存压力、更大的共享内存竞争。

我做过的不少 kernel 里，把 occupancy 从 100% 降到 50%（通过__launch_bounds__或cudaFuncSetAttribute限制最大线程数），性能反而提升了 30%。为什么？因为每个线程可用的寄存器多了，编译器能把更多临时变量留在寄存器里，减少局部内存溢出（local memory spill）和冗余的 load/store。访存次数少了，比多几个 warp 的延迟隐藏带来的收益更大。

所以 Occupancy 的本质是“延迟隐藏能力”和“每线程资源丰裕度”的权衡。真正的最优点需要针对具体 kernel 实验，没有一劳永逸的公式。

4.2 用 CUDA Occupancy API 精确计算

手工算资源比较麻烦，特别是寄存器分配还依赖编译器行为。CUDA 提供了 API 可以直接查询 kernel 在不同 block size 下的理论占用率。

#include <cstdio> #include <cuda_runtime.h> __global__ void myKernel(float* data, int n) { int idx = blockIdx.x * blockDim.x + threadIdx.x; if (idx < n) { data[idx] = data[idx] * 2.0f + 1.0f; } } int main() { int device = 0; cudaDeviceProp prop; cudaGetDeviceProperties(&prop, device); // 方式一：cudaOccupancyMaxActiveBlocksPerMultiprocessor int numBlocks = 0; cudaOccupancyMaxActiveBlocksPerMultiprocessor( &numBlocks, myKernel, // kernel 函数指针 256, // block 线程数 0 // 动态共享内存字节数 ); printf("blockSize=256, 每个 SM 最多活跃 block 数: %d\n", numBlocks); printf("理论占用率: %.1f%%\n", (float)numBlocks * 256 / prop.maxThreadsPerMultiProcessor * 100.0f); // 方式二：cudaOccupancyMaxPotentialBlockSize，自动找最佳 blockSize int bestBlockSize = 0; int bestNumBlocks = 0; cudaOccupancyMaxPotentialBlockSize( &bestNumBlocks, &bestBlockSize, myKernel, 0, 0 ); printf("推荐 blockSize=%d, 每个 SM 最多活跃 block 数=%d\n", bestBlockSize, bestNumBlocks); // 方式三：指定最小 blocks per SM，查看需要的资源（结合 launch_bounds） cudaFuncAttributes attr; cudaFuncGetAttributes(&attr, myKernel); printf("该 kernel 每线程寄存器数: %d, 每线程静态共享内存: %zu 字节\n", attr.numRegs, attr.sharedSizeBytes); return 0; }

这段代码编译后跑一下，能直接看到目标 GPU 上 kernel 的 occupancy 数据。我在实际项目中，性能优化第一件事就是跑这个，看理论占用率跟直觉差多少，然后决定要不要调 launch bounds。

4.3 用launch_bounds干预资源分配

__launch_bounds__可以告诉编译器：“这个 kernel 最多只会用这么多线程/每个 SM 最多要运行几个 block”，编译器会据此限制寄存器使用量，从而提高占用率，但代价是可能引入局部内存溢出。

__global__ __launch_bounds__(256, 8) void myKernel(float* data, int n) { // 编译时，编译器会尽量把每线程寄存器数限制在 // 65536 / (256 * 8) = 32 个以内（A100 类架构） ... }

这里256是每个 block 的线程数上限，8是每个 SM 期望的 block 数下限。两个参数配合，编译器算出每线程寄存器预算。如果预算太小，放不下所有局部变量，编译器就把一部分变量放到 local memory（本质是 spill，会拖慢速度）。

我遇到过一个典型的例子：一个深度学习算子，原本寄存器用了 64 个，occupancy 只有 50%。我加了__launch_bounds__(256, 6)，要求每 SM 至少 6 个 block，编译器把寄存器压到 42 个（65536 / (256*6) = 42.6，向下取整 42），occupancy 升到 75%，性能提升了 15%。虽然多了几次 spill，但延迟隐藏能力提升带来的收益更大。

反过来，如果你确定 kernel 的访存延迟低（比如纯算术密集），可以用__launch_bounds__(256, 2)这种参数给编译器“松绑”，让它放开手脚用寄存器，减少 spill。这全靠实测验证，不同 GPU、不同编译器版本结果都可能不同。

4.4 共享内存与 block 槽位的隐性约束

除了寄存器，共享内存是第二个常见瓶颈。这里要说一个新手容易踩的坑：动态共享内存大小是 runtime 才决定的，cudaOccupancyMaxActiveBlocksPerMultiprocessor计算时必须传入正确的动态共享内存字节数，否则算出来的 occupancy 虚高。

size_t dynSharedMem = 48 * 1024; // 每个 block 用 48KB 动态共享内存 int numBlocks = 0; cudaOccupancyMaxActiveBlocksPerMultiprocessor( &numBlocks, myKernel, 256, dynSharedMem );

还有一类容易忽略的资源是 block 槽位本身。现代 NVIDIA GPU 每个 SM 最多常驻的 block 数是 16～32，具体值取决于架构。假如你的 block 很小（比如 64 线程），就算资源完全够，SM 也放不下 32 个 block（架构硬上限），这时实际线程数就受限于maxBlocksPerSM * blockDim。这解释了为什么 block 大小不宣少于 128：太小会让 block 槽位变成瓶颈，白白损失占用率。

5. 典型调度场景：切换、同步与并发

5.1 线程块级切换：preemption 与 context switch

GPU 虽然号称“延迟隐藏机器”，但它做不到像 CPU 那样任意时刻抢占一个 warp 换另一个。它的技巧是通过大量常驻 warp 让慢操作“软切换”：内存访问没回来，warp 就停在 stalled 状态，调度器去执行别的 ready warp。这不需要保存/恢复上下文（因为每个 warp 的上下文一直在物理寄存器里），切换代价就是调度器的一次选择，接近零成本。

但有一种情况是有真实切换代价的：kernel 之间的切换。当两个 kernel 要分时共享 GPU 时，GPU 可能需要把前一个 kernel 的 block 执行完（drain），或者做一次完整的 context switch。Volta 引入了Independent Thread Scheduling和更细粒度的抢占支持，但开销仍然存在。

实践中，我尽量避免在同一 stream 中频繁启动小 kernel，或者用 CUDA Graphs 把一系列 kernel 捕获成一个图，一次性提交，减少启动和切换开销。对于需要“跑一点、停一下、再跑”的场景（比如解算器中的迭代），CUDA Graphs 的效果非常可观，有时候能省 30% 的启动时间。

5.2 Warp 级同步与 barrier 的调度影响

__syncthreads()是 block 内所有线程的屏障。它的实现依赖硬件 barrier 单元，调度器会保证 block 内所有 warp 都到达 barrier 后，才放行后续指令。问题在于：如果一个 block 的某个 warp 因为访存延迟迟迟不到达 barrier，其他已经到达的 warp 就全部卡住等待（停在 barrier）。

这就产生了一个 block 内 warp 的“木桶效应”。所以一个 block 内 warp 数不要太多（比如 block 1024 线程 = 32 warp），否则 barrier 等待成本更高。同时，尽量让每个 warp 在 barrier 前的工作量接近，避免个别 warp 拖慢整体。

Volta 引入的 Independent Thread Scheduling 改变了 warp 同步的底层语义：此前 warp 是额定的 SIMD 单位，线程间有隐式同步（一个 warp 的指令水平同步）；Volta 后每个线程有自己的程序计数器，可以独立推进。这解决了“一个线程进了死循环 or 一个线程等数据，整个 warp 都被卡死”的问题，但代价是如果代码里依赖了旧的隐式同步语义，可能引入 bug 或需要额外的显式同步。

5.3 多 Stream 并发与调度器交互

我见过最多的高性能场景翻车就在多 stream 上。NVIDIA 的任务调度模型支持多个 stream 并发执行 kernel，但“并发”不等于“同时开始、同时结束”。依赖关系、资源竞争都会影响实际调度。

如果两个 kernel 互相独立、资源要求低，它们可能在 GPU 上真并发（不同 SM 各跑各的）。但如果资源要求高（比如每个 kernel 都要 70% 的 SM），那么后启的 kernel 可能等前面的空出资源才能开始。GigaThread Engine 的动态 block 调度会让它们“见缝插针”——先进来的 kernel 的 block 先用资源，后进来的 kernel 排队。这本质上是greedy（贪心）的策略。

所以我建议多 stream 并发时：

• 每个 kernel 的资源占用（尤其共享内存和寄存器）要控制好，留出空间给其他 stream。
• 两个 kernel 的 grid 大小不必相等，也不要指望它们同时完成。
• 跨 stream 的事件同步（cudaEventRecord+cudaStreamWaitEvent）必须显式做好，否则可能因为调度顺序导致数据竞争。
• 如果想确定两个 kernel 是否真并发，用 Nsight Systems 看时间轴最直观，别猜。

6. 调度相关的排查与日志分析

6.1 从硬件计数器判断调度健康度

Nsight Compute 提供了很多与调度相关的 hardware counter。我在分析调度类问题时，重点看这几个：

指标	含义	排查线索
sm__warps_active.avg.pct_of_peak_sustained_active	平均活跃 warp 占比	如果长期低于 50%，说明延迟隐藏不足
sm__issue_active.avg.pct_of_peak_sustained_active	发射槽位使用率	通常比上面那个低很多，反映 IPC 上限
smsp__warp_issue_stalled_long_scoreboard	因 long scoreboard（全局访存）停滞的周期	占比高说明访存是瓶颈
smsp__warp_issue_stalled_barrier	因 barrier 停滞的周期	占比高说明 block 内负载不均衡
smsp__warp_issue_stalled_short_scoreboard	因共享内存/同步指令停滞	占比高说明共享内存访问冲突频繁
smsp__warp_issue_stalled_wait	因固定延迟（算术链）停滞	说明指令依赖太紧、ILP 不足

通常分析路径是：先看 active warp 占比，如果低，再看 stalled 分类，找到最大的停滞原因，然后针对性优化。比如：

• long scoreboard 高 → 优化数据局部性、增加并行访存、用__ldg或只读缓存。
• barrier 高 → 优化 block 内的负载均衡，或减小 block 大小。
• wait 高 → 循环展开、重排计算顺序，提高 ILP。
• short scoreboard 高 → 减少共享内存使用、消除 bank conflict。

6.2 Occupancy API 的实战排查步骤

遇到性能不符预期，我一般按这个顺序排查调度相关因素：

1. 用cudaOccupancyMaxActiveBlocksPerMultiprocessor算出理论占用率。如果理论就低（< 50%），先调 launch bounds 或改 block 大小。
2. 用 Nsight Compute 看实际占用率。如果实际比理论低，查是不是动态共享内存传参有误、是不是启动参数（比如cudaFuncSetAttribute的 max dynamic shared memory）没设对。
3. 看 stalled 分类。如果 long scoreboard 高，说明不是调度器不够用，而是延迟没藏住，优先加 ILP（独立指令并行）而不是加占用率。
4. 用 Nsight Systems 看 kernel 在时间轴上的位置。如果是多 stream，看是否有非预期的串行化。

这套流程能解决绝大多数“占用率看着很高但性能就是上不去”的怪问题。

6.3 调度对非 CUDA 场景的启示

最后说点延伸。NVIDIA 的任务调度模型影响的不只是 CUDA 开发者。TensorRT 做模型优化时会做 kernel fusion 和 graph optimization，本质就是减少 kernel 启动和切换开销、提高 SM 的执行密度。Jetson 嵌入式平台做多路视频解码，也依赖多个 engine（video decoder、GPU、DLA）之间的调度配合。明白 GPU 怎么分发和调度任务，很多上层框架的设计逻辑你一眼就能看穿。

就算你不写 CUDA，只用 PyTorch、TensorFlow，理解调度模型也能帮你解释很多怪现象：为什么同样的 batch size，A100 上比 3090 上快这么多；为什么一个小算子被频繁调用反而比一个大算子慢；为什么有时候torch.cuda.synchronize()的位置会影响性能。这些背后都是调度在起作用。

7. 调度模型演进与未来方向

从数据驱动的角度来看，NVIDIA 的调度策略一直在向“细粒度、低开销、高并发”演进。

Fermi 时代是单队列、block 级粗粒度分发。Kepler 引入 Hyper-Q 和动态并行（Dynamic Parallelism），让 GPU 内部也能启动 kernel。Volta 引入 Independent Thread Scheduling 和更细粒度的同步原语。Ampere 的 MIG（Multi-Instance GPU）在硬件层面把 GPU 切分成多个独立实例，每个实例有独立的 GigaThread Engine 和 L2 切片，调度隔离性大大增强。Hopper 的Thread Block Clusters则让多个 block 可以组成集群，在更小的范围内共享数据并协同调度，这是对“跨 SM 通信”需求的直接回应。

这些演进的核心都围绕一个矛盾：执行单元越来越多、越来越快，但片上存储和互连带宽的增长永远跟不上，所以调度器必须在更细的粒度上做决策，让每一笔数据搬运都在最合适的时间发生、被最合适的计算单元消费。

另一个值得关注的方向是GPU 上的硬件线程调度与 AI 工作负载的适配。Transformer 这类模型中有大量小矩阵乘法、layernorm、softmax 等 kernel，SM 的实际利用率往往不高。为此，NVIDIA 在 cuBLAS、TensorRT 里加入了不少“kernel fusion”和自动调度的策略，本质上是用软件调度弥补硬件调度的粒度不足，把多个小 kernel 合并成单个大 kernel，减少 block 切换和 barrier 开销。这对我们写 CUDA 的启示是：盯住调度器能看到什么样的“指令流”，比盯住单个 kernel 的计算量更重要。

我个人的体会是：调度模型的演进方向，始终是让“调度开销”在总执行时间里占据的比例越来越小。无论是 Hyper-Q、动态并行、Independent Thread Scheduling 还是 Thread Block Clusters，目标都很朴素：让计算单元尽量一直有活干，让排队的时间尽量短。理解这一点，很多架构层面的变化就不难理解了。