这是一篇经验总结性质的技术分析。平时我们写 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分发到各个 GPCGraphics 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 程序时调用kernelgrid, 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 发给哪个 SMNVIDIA 没有公开 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 个 blockblock 很大、资源占用高尾效应会更明显。所以工程上的第一个实在建议是block 数量最好明显多于 SM 数量既保证均衡负载也缓冲 block 完成时间的差异。通常我建议 block 总数至少是 SM 数量的 23 倍理想情况是 48 倍。// 一个简单的 tail effect 演示grid 大小固定时性能可能波动 int smCount; cudaDeviceGetAttribute(smCount, cudaDevAttrMultiProcessorCount, deviceId); // 假设我们的 block 数等于 SM 数理论上每个 SM 分到 1 个 block // 但实际上 block 执行时间不可能完全相同总有快的、有慢的 // 慢 block 会让所在 SM 继续工作快 block 所在 SM 已经空闲 // 于是整体执行时间被最慢的那个 block 决定 // 解决思路让 grid smCount * blocksPerSM * kk 为整数倍率Grid 规模选择属于“看得见摸得着”的调度调优几乎每次性能分析都会遇到。2.3 streaming multiprocessor 上的 resident block 管理Block 到达 SM 之后并不是立即就能执行它必须先“入住”SM。SM 会为 block 分配block 槽位每个 SM 最多可常驻的 block 数不同架构不同一般是 1632 个warp 槽位每个 block 包含若干个 warpSM 需要给这些 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 个寄存器/blockSM 寄存器文件 6553665536 / 20480 3.2向下取整 3 个 block每个 SM 最多 32 个 block架构限制3 32所以 block 数量不是瓶颈block 总线程数 3 × 512 1536 线程如果 SM 最大线程数为 2048则占用率 1536 / 2048 75%同样的 kernel如果把寄存器压到 32 个512 × 32 1638465536 / 16384 4可以放 4 个 block总线程数 4 × 512 2048占用率 100%多 25% 的并发 warp在访存密集场景下往往意味着 20% 的性能提升。这就是寄存器分配和调度的耦合关系——后面第 4 节我们详细算。3. 微观调度Warp 的诞生与发射3.1 Thread 到 Warp 的固定映射Block 成功入住 SM 后硬件会把 block 内的线程按连续 32 个一组组成warp。这个映射是固定的、线性的thread 031 是 warp 03263 是 warp 1依此类推。你无法改变这个映射关系但可以预判它如果 block 大小是 128那 warp 03 分别对应 thread 031、3263、6495、96127。这个固定映射有两层直接影响。第一分支发散divergence的经济代价是 warp 级别的。同一个 warp 里如果线程走了不同分支if/else硬件会先执行一个分支路径再执行另一个两条路径都执行完才算这个 warp 完成这条指令。对应到我们编程尽量让分歧发生在 warp 边界上。比如判断if (threadIdx.x 32)如果 block 大小为 64线程 031 走 if3263 走 else两个 warp 完全独立不会产生 divergence尽管代码本身有分支。第二共享内存 bank conflict、全局内存访问的合并性都取决于 warp 内 32 个线程的访存地址分布。固定映射意味着“哪些线程会在同一个 warp 里”是确定的你可以据此组织数据结构。3.2 Warp Scheduler 与 Instruction DispatchWarp 进入 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 是 16Kepler/Maxwell 是 8。Warp scheduler 的职责是每个时钟周期cycle挑选一个或多个可以发射的 warp把当前指令送到执行单元。什么叫做“可以发射”就是这条指令的操作数已经准备好、执行单元有空位、没有结构性阻塞比如前一条指令还在用同一个单元写回、内存指令的队列满了、barrier 尚未到达等。不同代际的调度器发射能力不同Fermi/Kepler/Maxwell/PascalVolta 之前每个 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 全局内存访问的延迟大约需要 1220 个独立的 Ready warp具体取决于访问是 L2 命中还是 DRAM 命中。举例计算假设内存延迟 400 周期这是一个典型量级A100/H100 的 DRAM 延迟大约在 300500 周期视访问模式而定。假设每个 warp 发射一个独立的 128 字节内存访问后就停滞等待数据返回。那么我们需要大约 400 条独立的 load 指令同时在飞才能每周期发射一条新 load。每个 SM 最多 64 个 warp2048 线程 / 32每个 warp 每周期最多发射 1 条 load4 个 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(blockSize256, 每个 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向下取整 42occupancy 升到 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 数是 1632具体值取决于架构。假如你的 block 很小比如 64 线程就算资源完全够SM 也放不下 32 个 block架构硬上限这时实际线程数就受限于maxBlocksPerSM * blockDim。这解释了为什么 block 大小不宣少于 128太小会让 block 槽位变成瓶颈白白损失占用率。5. 典型调度场景切换、同步与并发5.1 线程块级切换preemption 与 context switchGPU 虽然号称“延迟隐藏机器”但它做不到像 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 的事件同步cudaEventRecordcudaStreamWaitEvent必须显式做好否则可能因为调度顺序导致数据竞争。如果想确定两个 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 的实战排查步骤遇到性能不符预期我一般按这个顺序排查调度相关因素用cudaOccupancyMaxActiveBlocksPerMultiprocessor算出理论占用率。如果理论就低 50%先调 launch bounds 或改 block 大小。用 Nsight Compute 看实际占用率。如果实际比理论低查是不是动态共享内存传参有误、是不是启动参数比如cudaFuncSetAttribute的 max dynamic shared memory没设对。看 stalled 分类。如果 long scoreboard 高说明不是调度器不够用而是延迟没藏住优先加 ILP独立指令并行而不是加占用率。用 Nsight Systems 看 kernel 在时间轴上的位置。如果是多 stream看是否有非预期的串行化。这套流程能解决绝大多数“占用率看着很高但性能就是上不去”的怪问题。6.3 调度对非 CUDA 场景的启示最后说点延伸。NVIDIA 的任务调度模型影响的不只是 CUDA 开发者。TensorRT 做模型优化时会做 kernel fusion 和 graph optimization本质就是减少 kernel 启动和切换开销、提高 SM 的执行密度。Jetson 嵌入式平台做多路视频解码也依赖多个 enginevideo decoder、GPU、DLA之间的调度配合。明白 GPU 怎么分发和调度任务很多上层框架的设计逻辑你一眼就能看穿。就算你不写 CUDA只用 PyTorch、TensorFlow理解调度模型也能帮你解释很多怪现象为什么同样的 batch sizeA100 上比 3090 上快这么多为什么一个小算子被频繁调用反而比一个大算子慢为什么有时候torch.cuda.synchronize()的位置会影响性能。这些背后都是调度在起作用。7. 调度模型演进与未来方向从数据驱动的角度来看NVIDIA 的调度策略一直在向“细粒度、低开销、高并发”演进。Fermi 时代是单队列、block 级粗粒度分发。Kepler 引入 Hyper-Q 和动态并行Dynamic Parallelism让 GPU 内部也能启动 kernel。Volta 引入 Independent Thread Scheduling 和更细粒度的同步原语。Ampere 的 MIGMulti-Instance GPU在硬件层面把 GPU 切分成多个独立实例每个实例有独立的 GigaThread Engine 和 L2 切片调度隔离性大大增强。Hopper 的Thread Block Clusters则让多个 block 可以组成集群在更小的范围内共享数据并协同调度这是对“跨 SM 通信”需求的直接回应。这些演进的核心都围绕一个矛盾执行单元越来越多、越来越快但片上存储和互连带宽的增长永远跟不上所以调度器必须在更细的粒度上做决策让每一笔数据搬运都在最合适的时间发生、被最合适的计算单元消费。另一个值得关注的方向是GPU 上的硬件线程调度与 AI 工作负载的适配。Transformer 这类模型中有大量小矩阵乘法、layernorm、softmax 等 kernelSM 的实际利用率往往不高。为此NVIDIA 在 cuBLAS、TensorRT 里加入了不少“kernel fusion”和自动调度的策略本质上是用软件调度弥补硬件调度的粒度不足把多个小 kernel 合并成单个大 kernel减少 block 切换和 barrier 开销。这对我们写 CUDA 的启示是盯住调度器能看到什么样的“指令流”比盯住单个 kernel 的计算量更重要。我个人的体会是调度模型的演进方向始终是让“调度开销”在总执行时间里占据的比例越来越小。无论是 Hyper-Q、动态并行、Independent Thread Scheduling 还是 Thread Block Clusters目标都很朴素让计算单元尽量一直有活干让排队的时间尽量短。理解这一点很多架构层面的变化就不难理解了。