CUDA计算内核调度模型：从线程块分配到warp发射

一条kernel启动语句，在CPU侧可能只是几行代码，真正交给GPU执行时，内部要经过“提交—分发—驻留—发射—执行”一整条链路。很多做CUDA性能优化的朋友都有过这种经历：明明把占用率调到了理论最大值，程序反而变慢了。这种反直觉现象，根源就在计算内核的调度模型——它不是简单地把线程丢给众多核心去跑，而是有一套分层的裁决规则。

这篇文章是NVIDIA调度分析系列的第4篇，主题就是计算内核调度模型。我会把一个kernel从CPU提交到GPU执行，再到多个内核并发、抢占的完整过程拆开讲，顺带解释GigaThread Engine、线程块分配、warp调度器、流并发这些机制之间是怎么配合的。看完以后，你再回看那些“看似合理但跑不快”的kernel，大概率能直接定位到问题出在哪一层。

如果你正在做CUDA性能优化，或者准备系统级的面试，这篇文章值得认真读一遍。不涉及太深的硬件微架构细节，但我会把关键参数和推算思路给足，方便你迁移到自己的场景里。

1. 从一个kernel启动说起：调度的参与者都有谁

1.1 CPU侧的提交路径不只是一次函数调用

很多人把cudaLaunchKernel当成普通函数调用，觉得执行到这一行，GPU就开始算了。实际上，CPU侧的CUDA Runtime只是把这次启动“记录下来”，塞进一个与GPU共享的环形命令缓冲区（pushbuffer）。驱动再把这个缓冲区里的命令结构体通过系统调用提交给GPU前端。

这个过程中的关键点在于：cudaLaunchKernel本身是异步的，它不等kernel执行完就返回了。 你后面必须用cudaStreamSynchronize或cudaDeviceSynchronize才能真正等到结果。很多刚入门的人在这里被坑过：测单次kernel耗时，把同步写在时间统计之外，结果测出来永远是个位数微秒——实际上那只是提交命令的耗时。

从命令提交到GPU实际开始执行，中间还隔着一个“前端解析”阶段。GPU前端引擎把命令结构体解析成内部描述符，其中就包含了grid维度、block维度、kernel参数指针、使用的流ID等等。随后，这些描述符会被转交给一个专门的硬件模块来决定下一步怎么分配——这就是GPU侧的全局调度器。

1.2 GigaThread Engine是GPU侧的大管家

老一点的NVIDIA架构资料里，会把GigaThread Engine叫作“工作分配器”（Work Distributor）。它在调度模型里的角色，类似操作系统的进程调度器，但处理粒度完全不同。GigaThread Engine负责把一个grid里的众多线程块（Thread Block）分配到各个SM上。

关键点在于：调度粒度是线程块，不是线程，更不是warp。 为什么不用warp作为全局调度单位？因为线程块是一个拥有共享内存、block级同步原语（__syncthreads）的独立执行单元，它必须完整地驻留在一个SM上才能工作。如果把warp打散了，共享内存和同步语义就全乱了。

GigaThread Engine在分配block时，不是简单轮询每个SM，它会读取每个SM当前已经驻留的block数量、已用资源，然后挑选资源充足的目标SM。实际实现细节没有完全公开，但从多流并发和block分配行为可以反推，它一定维护了一张“SM资源表”，并且会做一定程度的负载均衡。否则你很难解释为什么在多个流并发提交时，GPU能把空闲资源分配给后到的kernel的block。

2. 线程块如何被“安放”到一个SM上

2.1 SM的资源池决定你能放多少block

一个SM上有几个关键资源：warp槽位（比如Ampere A100是64个warp槽位，对应2048个线程）、寄存器文件（A100 SM是65536个32位寄存器）、共享内存（A100 SM有164KB，可配置）、最大驻留block数（通常32个）。这些数字不一定代代相同，H100的SM依然保留64个warp槽位，但寄存器和共享内存容量上有变化。

block要被“安放”到某颗SM，必须同时满足四类条件：block总数不超过SM的block槽位限制，总线程数不超过最大线程数限制，寄存器用量总和不超过寄存器文件大小，共享内存用量总和不超过共享内存大小。只要有一条不满足，这个block就得等下一波。

你可以把SM想象成一间旅店：warp槽位是床位数，寄存器文件是储物柜数量，共享内存是公共区域面积。线程块是一批旅客。旅客再多，床不够就是住不下；储物柜不够，旅客的东西没地方放，也不行。

2.2 block内部是怎么变成warp的

线程块被分配到一个SM之后，块内的线程会按x、y、z维度的线性顺序展开，每连续32个线程组成一个warp。比如一个256线程的block，展开后就是8个warp。warp是SIMT（单指令多线程）执行的基本调度单位：同一个warp里，所有线程某个时刻执行的是同一条指令，只是数据不同。

这里有个非常容易被忽略的点：block大小最好是32的倍数。 如果你设了100个线程，最后一个warp只有4个有效线程，剩下28个线程槽位被浪费。这些浪费的槽位依然占用SM上的warp slot，影响后续block的驻留，白白降低有效利用率。所以CUDA里最常见的block大小是128和256，不是因为好看，而是因为正好是32的整数倍，连锁色都不用管。

2.3 资源方程与占用率计算

实际决定一个block能不能上SM的，往往是寄存器和共享内存，而不是单纯线程数。我们可以用一个小例子感受一下。

假设你的kernel每个线程用64个寄存器，block大小256线程，那么一个block就需要256×64=16384个寄存器。如果SM寄存器文件是65536个寄存器，则最多驻留4个这样的block，也就是1024个线程，占用率只有50%。如果你把每线程寄存器压到32个，同样block大小下就能放8个block，达到满占用率。

共享内存同理：假设每个block静态分配64KB共享内存，SM一共164KB，那么最多驻留2个block，线程占用率可能很低，因为一块共享内存占掉了太多资源。

这就是为什么Nsight Compute的“Occupancy”页面会把限制因素拆成几项：Block Limit、Threads Limit、Registers Limit、Shared Memory Limit。你看一眼就知道当前kernel是被哪个资源卡住了。经验之谈，很多新手看到“寄存器数太高”就急着限制寄存器，结果引入local memory溢出，性能反而更差。寄存器个数不是越低越好，得看实际访存行为和指令依赖。

3. warp调度器内部的黑盒：一个周期里到底发生了什么

3.1 驻留warp的状态流转

block安放到SM上之后，SM的warp调度器成为主角。驻留在SM上的每个warp并不总是处于可执行状态，它们会在几种状态之间切换：空置（unallocated）、驻留（resident）、就绪（ready）、阻塞（stalled）。

阻塞的原因很多：可能是在等全局内存/局部内存的数据返回，可能是在等共享内存的访问完成，可能是被__syncthreads卡住等同一个block的其他warp，也可能是执行单元被占满。调度器每个周期会从所有“就绪”的warp里挑选一个，把它的下一条指令发射到执行单元。如果一个周期内没有任何warp就绪，SM就出现所谓的“执行气泡”，这是性能损耗的直接来源。

理解这个状态流转后，很多优化手段就好懂了。比如用__launch_bounds__限制每个block的线程数，本质上是减少驻留warp总数，让每个warp能分到更多资源，提高单个warp的执行效率——但代价是并发warp变少，延迟隐藏能力变弱。这从来不是免费午餐。

3.2 多调度器并行与发射宽度

一个SM通常有4个warp调度器（Ampere、Ada等主流架构如此）。每个调度器管理一组warp槽位。A100的SM里有64个warp槽位，4个调度器，每个调度器负责16个warp。每个调度器每个周期可以发射一条指令，理论上一个SM每周期能发射4条指令。

但这不意味着4条指令可以任意选择执行单元。执行单元（ALU、SFU、LD/ST单元等）在SM内部是共享的。如果4条指令全是访存指令，可能受限于LD/ST单元数量或访存带宽；如果全是分支指令，分支单元又成了瓶颈。所以你在Nsight里看到的“issue utilization”不高，有时候不是因为调度器偷懒，而是指令混合得不好，执行单元冲突严重。

做性能分析时，我推荐单独看每个调度器的活动周期，而不是只看整个SM的利用率。一些工具可以按SM分区看发射统计，如果某个分区明显比其他分区忙，说明block分配或warp迁移产生了不均衡。不过更多时候，同一SM内的block不会导致太大不均衡，因为warp是动态调度的。

3.3 延迟隐藏到底需要多少warp

warp调度的核心价值是隐藏延迟。当一个warp执行一条高延迟指令（比如全局内存load，可能几百个周期才能返回结果），如果SM只有这一个warp，它就只能干等。如果有其他就绪warp，调度器就可以在等待期间切换到它们，让执行单元继续运转。

要估算“够不够用”，有个粗略思路：假设一次全局load延迟是L个周期，而调度器每个周期最多发射1条指令。如果每个warp每轮只有一条load指令，那么在等待返回的L个周期里，需要大约L个其他warp的指令来填补发射槽。当然实际计算没这么简单，因为一个warp通常包含多级指令流水，且load不一定每轮只有一条，还可能穿插计算指令。但趋势是对的：访存越密集、单次访存延迟越高，需要的驻留warp就越多。 这也是为什么访存密集kernel把占用率拉到80%以上通常有正收益，而计算密集kernel占用率50%也可能已经跑满。

我自己的经验，对于访存密集的kernel，每个调度器至少要有8~12个就绪warp，延迟隐藏才会比较充分；计算密集kernel，4~8个就够。这些都是经验值，具体还是得用Nsight Compute看“Warp Stall”统计。

4. 内核之间如何交错：流、并发与抢占

4.1 Stream和Event是内核级调度的用户入口

前面聊的是单个kernel内部怎么调度，现在上升到内核与内核之间。CUDA流本质上是GPU侧的命令队列。提交到同一个流里的kernel严格按顺序执行，因为它们在用户语义上是串行的；不同流里的kernel没有显式依赖，GPU有机会交错执行。

Event的作用是在流里插一个“标记”，用来跨流同步。但要注意，调用cudaEventRecord本身并不阻塞CPU，只有cudaEventSynchronize或cudaStreamWaitEvent才会让某个流等待另一个流的Event。

很多人在多流编程里发现“并发不生效”，回放时间线发现kernel根本没有重叠。最常见的原因不是硬件不支持，而是你的代码里插了隐式同步点。比如在流之间调用了cudaMalloc、cudaMemcpy这类函数，它们可能会在设备端同步整个上下文；或者你的流事件等待关系写得太宽，把本可并行的流串起来了。用Nsight Systems看一眼时间线，比瞎猜快很多。

4.2 并发内核执行与block填充

当两个kernel在不同流里提交，且第一个kernel的block数量不足以占满整个GPU时，GigaThread Engine可以让第二个kernel“插空”执行自己的block。这也是为什么GigaThread Engine需要全局资源表：它得清楚每个SM还剩多少block槽位、多少寄存器、多少共享内存。

CUDA还提供了流优先级：创建流时用cudaStreamCreateWithPriority指定更高的数值，高优先级流里的kernel会更早获得资源分配。这个机制对延迟敏感的应用很有用，但别指望它是硬实时抢占。优先级影响的主要是block分发顺序和资源竞争时的倾向，并不是高优先级kernel一到，低优先级kernel立刻让位。

Volta架构以后，NVIDIA引入了细粒度计算抢占，内核可以在指令级被打断并切换上下文，需要时可让高优先级工作快速上位。但这类机制对绝大多数应用是透明的，我们做优化时不需要过度依赖，关键还是经营好block数量和流依赖结构。

4.3 Persistent Kernel是把调度权抢回自己手里

有一种工程做法叫persistent kernel：启动固定数量的block，让它们常驻SM，用循环在内部处理大量任务。这样做的动机是减少GigaThread Engine反复分配block的开销，顺便让程序员自己控制任务划分。

但这招有代价。硬件调度器的负载均衡能力很强，persistent kernel如果拆分任务不均衡，就会出现某些block早早干完活空转，其他block还在处理长尾任务。真要用persistent kernel，最好搭配动态任务队列（比如global memory上的atomic计数器或者工作队列），让每个常驻block空闲时主动去取下一个任务。这等于把软件调度的活又接回来了。

老实说，除非你的kernel启动次数极多、每次任务量极小，否则我不建议一上来就搞persistent kernel。先让硬件调，不行再抢调度权。

5. 实测中容易误判调度行为的几个场景

5.1 占用率越高不代表越快

占用率只是“有多少warp在SM上驻留”的指标，不等于“这些warp真的在执行”。一个典型的反例：为了让占用率从50%到100%，把每线程寄存器数限制到很低，结果部分寄存器变量被编译器溢出到local memory。local memory虽然物理上在显存，但走的是L1/L2缓存路径，延迟比寄存器高得多。这时驻留warp变多了，每个warp的执行效率却大幅下降，总性能反而倒退。

碰到占用率与预期不符的情况，我会先用Nsight Compute看当前kernel的Achieved Occupancy，再对比理论占用率。如果两个差距很大，说明调度过程中有大量warp阻塞或资源等待；如果理论占用率就低，再回看资源限制是哪一项。

__launch_bounds__是个好工具，但使用时要带着问题去用：我是想减少寄存器占用，还是想控制block数量？先分析瓶颈，再动手减寄存器，否则容易按下葫芦浮起瓢。

5.2 block数量过多或过少都会出问题

block总数如果太少，比如GPU有108个SM，你只启动108个block，每个SM一个block，那么只要有一个block因为访存或同步慢了一点，整个SM就可能闲下来，形成明显长尾。一般建议block总数至少是SM数的几倍，最好能形成多个“wave”，这样调度器有机会在一个block执行完后立刻补上新的block。

反过来，如果把block拆得特别碎，每个block只有32或64线程，虽然每个SM能驻留很多block，但每个block都有创建、同步、回收的开销，调度器要频繁切换管理对象。实测下来，block大小低于64通常是不划算的，除非算法本身有强约束。我自己的通用起点是128线程/block，遇到访存密集或资源紧张再试256。

5.3 从Nsight的Warp Stall入手验证调度模型

前面讲了一堆理论，实际验证靠Nsight Compute最直接。在Kernel profile的Warp State里，你会看到各种stall原因，比如：

• Long Scoreboard：warp在等全局/局部内存返回，说明访存延迟隐藏不够，或访存并行度低。
• Short Scoreboard：等共享内存等短延迟操作，相对问题不大，但如果占比过高，可能是共享内存bank conflict。
• Barrier：等同一个block内其他warp到达同步点，要么block内任务划分不均，要么block太小同步开销占比高。
• Wait：等固定延迟指令流水完成，属于指令依赖，说明ILP（指令级并行）不够。
• MIO Throttle：访存指令队列满了，说明同一时刻发射访存指令太多，冲击了存储器子系统。

我通常的调试顺序是：先看所有stall原因中的Top 1，如果是Long Scoreboard，优先增加并行warp数或并行访存数量（比如循环展开增加独立load）；如果是Wait，优先改代码减少指令依赖；如果是Barrier，思考是不是block太大导致同步代价高。

命令行跑一个kernel的分析也很简单：

它会输出大量指标，不用全看，先找Warp Stall Reasons和Achieved Occupancy两块就够了。很多网上说“这个kernel带宽没跑满”的判断，其实都能从stall原因里直接看出来，根本不用瞎猜。

5.4 多流并发不生效时先查隐式同步

最后分享一个排查多流并发的经验。Nsight Systems是这时候的主角，它能画出每个流上kernel的时间线。如果多个流的kernel在时间线上排成了一条直线，没有重叠，我会去检查调用链里有没有这些“同步刺客”：

• 在流之间使用了cudaDeviceSynchronize
• 流中执行了cudaMemcpy（默认是同步语义，虽然可以拷到cudaMemcpyAsync，但也需要页面锁定内存）
• 调用了cudaMalloc或cudaFree（可能触发全设备同步）
• 同一个流里前一个kernel运行时间太长，导致后一个kernel没机会“插空”

实在排不掉这些同步点，还可以试试CUDA Graph。CUDA Graph把多个kernel的依赖关系固化下来，提交一次就能连续执行多个kernel，省掉重复启动的开销，对于短kernel多流的场景效果尤其明显。它本质上是在利用调度模型对整张任务图做优化，而不是每次启动后让GigaThread Engine临时决策。

写在最后的调度排查习惯

这套调度模型单看每个概念都不难，难的是同时把它们放在一起思考。我自己的习惯是：拿到一个性能问题，先打开Nsight Compute看warp stall，确定是延迟隐藏问题还是执行依赖问题；再回头检查block大小和占用率；最后用Nsight Systems看多流时间线，排除并发结构上的“串行陷阱”。按照这个顺序走下来，大部分“说不清楚为什么慢”的kernel都能定位到具体原因。

NVIDIA的硬件每一代都在微调调度器的细节，但计算内核调度的总体分层框架——CPU提交、GigaThread分配block、SM驻留、warp调度器发射指令——是很稳定的。把这一层吃透，读后续的microbenchmark数据、分析别的GPU架构，你会发现思路基本都是通的。