GPU 调度这块我几年前刚开始认真写 CUDA kernel 的时候最直观的感受是明明照着官方文档把线程块大小设成了 256 occupancy 算出来也还行但性能就是上不去。后来把 NVIDIA 的任务调度模型真正啃了一遍才明白瓶颈到底卡在哪。这篇是这个系列的第二章前面聊完整体调度分析的大框架这次把 NVIDIA 的任务调度模型单独拉出来拆开揉碎讲清楚。先说清楚这篇要解决什么问题。很多人把“调度”理解成操作系统里那种进程/线程切换觉得 GPU 调度就是一堆线程排队上核心。这个理解错得很离谱。GPU 的调度是硬件级的、以线程束warp为最小单位的、基于状态机轮转的调度方式它压根不走操作系统内核也不需要什么上下文切换的软件开销。不理解这一层你写 kernel 的时候连“为什么 block 大小设成 128 有时比 256 好”这种问题都答不上来。这篇文章的受众我默认是对 CUDA 编程有一定基础、但没系统研究过 GPU 内部机制的开发者。如果你手头有 NVIDIA 的显卡跟着思路把 profiling 数据打开对照着看理解速度会快很多。下面先把最关键的几张“硬件底图”铺开。1. 先从硬件底图说起GPC、TPC 与 SM 的关系1.1 一张图理解 GPU 的物理结构分层NVIDIA GPU 的物理结构从上到下大概可以切成这么几层整个芯片叫 GPU芯片内部被划分成多个GPCGraphics Processing Cluster图形处理集群每个 GPC 内部包含若干个TPCTexture Processing Cluster纹理处理集群TPC 再往下拆就是SMStreaming Multiprocessor流式多处理器。每一层各管各的事GPC可以理解成芯片里的“大园区”它内部有独立的几何引擎、光栅化单元、以及若干个 SM。图形渲染任务里的三角形装配、光栅化这类活就是在 GPC 这个级别分工的。TPC一个 TPC 通常包含两个 SM老架构里也有一个 TPC 只挂一个 SM 的情况。TPC 主要起的是“分组管理”作用让纹理单元和 SM 之间的数据通路更紧凑。SM这才是真正的“计算工人”所有 CUDA 核心、Tensor Core、共享内存、寄存器堆、调度器全都住在 SM 里面。SM 是任务调度真正发生的地方。拿 Ampere 架构的 GA102 核心来举例它一共有 7 个 GPC、42 个 TPC、84 个 SM每个 SM 里有 128 个 CUDA 核心整卡总共就有 10752 个 CUDA 核心。这个数字看着唬人但你要记住这些 CUDA 核心并不是像 CPU 那样被“单个指令流”驱动它们是靠 SM 内部调度器喂数据才能干活的。这里插一句我踩过的坑早期我误以为一个 CUDA 核心对应一个线程那我开一万个线程就能让一万个核心满负荷跑。实际上线程和核心的关系是通过线程束和调度器间接建立的线程数超过核心数不代表就能并行执行得看调度器能喂多快。1.2 SM 内部的关键部件和调度相关寄存器SM 内部藏着几样和调度直接相关的“家当”我按重要性排个序线程束调度器Warp Scheduler每个 SM 里有 4 个Volta 之后普遍是 4 个这是调度的“大脑”负责从就绪的线程束里挑一个发指令。指令分发单元Dispatch Unit每个调度器旁边配一个负责把取到的指令真正发给执行单元。寄存器堆Register File每个 SM 里有 65536 个 32 位寄存器64K 是 Ampere 常见的配置。这些寄存器是给线程存局部变量的总容量固定所以线程越多、每个线程能分到的寄存器就越少。共享内存Shared Memory常见 128KB 或 164KB 每 SM用于线程块内部通信。Warp 状态表SM 内部维护一张表记录每个 warp 当前处于什么状态就绪、等待、阻塞等这是调度器做决策的依据。这些部件之间的关系我习惯用“餐厅后厨”来类比。GPC 是餐饮集团的一个门店TPC 是店里的几个档口SM 就是一个灶台warp 就是灶台上的炒锅。炒锅warp里装着一批菜线程灶台师傅调度器按照火候情况决定先翻炒哪口锅。每个灶台不止一口锅师傅也不止一个关键是谁能让锅不闲着。2. 任务调度的最小单位warp 的概念与意义2.1 为什么偏偏是 32 个线程NVIDIA 的硬件线程调度单位不是单个线程而是warp线程束一个 warp 恰好包含 32 个线程。这 32 个线程在硬件层面是“锁步”lockstep执行的也就是说同一个 warp 里的所有线程在同一时刻执行同一条指令只是处理的数据不同。为什么是 32 不是 16 也不是 64这里没有特别神秘的答案主要是硬件设计上的折中。32 个线程一条指令意味着指令 fetch 和 decode 的开销被 32 份工作分摊了这个倍数刚好能填满现代 GPU 执行单元的多周期流水线。历史上 Fermi 时代就定了 32 这个数后续架构一直沿用说明这个设计经过了充分的实践验证。你可以这样想CPU 一个核一次就处理一个线程的指令GPU 一个调度器一次处理 32 个线程的指令这就是 GPU 能靠“少控制、多并行”换取吞吐量的根本原因。代价是warp 内如果出现分支发散性能会断崖式下跌。2.2 分支发散为什么是性能杀手想象一个 warp 里有 32 个线程代码里写了一个 if-else 分支一半线程走 if一半走 else。硬件没法让这 32 个线程同时执行不同指令于是调度器只好先让走 if 的线程执行再让走 else 的线程执行最终这个 warp 需要串行执行两个分支的指令。这意味着什么本来一个周期能完成的指令现在要两个周期如果一个分支里嵌套了更多分支耗时还要成倍上涨。写 kernel 的时候最怕的不是分支本身而是同一个 warp 里出现“有的线程走 A、有的线程走 B”的分叉。代码层面怎么避免核心思想是让分支判断基于 warp 内统一的值而不是基于线程独有的数据。比如可以根据threadIdx.x / 32来分块处理让同一 warp 内的线程走同一条路径或者把分支逻辑改成算术运算用掩码来“选择”结果而不是真写 if。实操心得我早期写归约reduction核函数的时候用了if (tid % 2 0)这种写法做折叠操作结果性能比预期慢了一倍还多。后来改成让每个线程固定处理连续两段数据再配合 warp shuffle性能直接翻倍。分支发散这个坑真的是你不亲自拿 profiling 数据对比光靠感觉完全发现不了。2.3 线程束与线程块的换算关系一个线程块block由若干 warp 组成block 内的线程数必须是 warp 大小的整数倍吗硬件不强制但实践里强烈建议这样做。因为如果你把 block 大小设成 100硬件会把它补成 4 个 warp128 个线程其中 28 个线程是“空转”的它们不干任何活但照样占用调度资源和寄存器。所以你会看到业界几乎所有 kernel 的 block 大小都习惯性设成 32 的倍数128、256、512 都有但很少看见 100、300 这种数字。把 block 设成 warp 大小的整数倍是最基本的入门礼仪。3. SM 调度器的工作机制从取指到发射的完整链路3.1 四个调度器如何分工协作前面提到 Ampere 架构的 SM 里有 4 个 warp scheduler每个 scheduler 配一条独立的指令发射通路。这意味着一个 SM 在同一个时钟周期内最多能从 4 个不同的 warp 里各取一条指令发射出去。但这不意味着一个周期只能发射 4 条指令。NVIDIA 的调度器支持dual-issue也就是一个调度器在一个周期里可以发射两条指令前提是这两条指令互不依赖、并且执行单元有空闲。所以理论上一个 SM 一个周期最多可以发射 8 条指令4 个调度器 × 2 条。调度器怎么决定发射谁的指令底层逻辑是每个 warp 在每个周期都有“状态位”要么是“就绪可发射”eligible要么是“等待某种资源”stalled。调度器从 eligible 的 warp 里挑一个具体策略各家 NVIDIA 架构里略有不同常见的是轮转和优先级结合的方案。3.2 指令发射的延迟隐藏原理GPU 的调度器做一次“挑 warp 并发射指令”的动作在硬件层面只要几个时钟周期。但它真正的精髓在于用大量 ready 的 warp 来掩盖长延迟操作。比如从全局内存取数延迟可能有几百个周期CPU 的做法是停下等GPU 的做法是这个 warp 等着取数调度器立刻切换到另一个就绪的 warp 继续发射指令。这个机制叫latency hiding延迟隐藏。它要求 SM 里同时有足够多“蓄势待发”的 warp否则一旦所有 warp 都在等内存返回SM 就“空转”了也就是所谓的occupancy 不足。来算一笔账。一个 SM 有 4 个调度器每个调度器每个周期最多发射 2 条指令也就是每个周期最多需要 8 个 warp 供它轮转。假设每个 warp 的平均停顿周期数是 200等内存那 SM 里至少要同时有 8 × 200 1600 个 warp 在轮转才能把每周期 8 条发射管线喂满。而实际上 SM 里能容纳的 warp 总数是有限的——Ampere 架构最多 64 个 warp2048 线程per SM。这个远小于 1600所以现实是 GPU 只能做到“尽量隐藏延迟”不可能完全隐藏。这就是为什么occupancy 高不代表性能一定好但 occupancy 太低性能一定好不了。如果你的 kernel 每个线程用了太多寄存器导致 SM 里能同时常驻的 warp 数减少延迟隐藏能力就下降内存延迟就会暴露出来。3.3 调度器怎么知道你“准备好了”这个细节容易被忽略。每个 warp 的每条指令在进入执行流水线之前要经过所谓的scoreboard记分板机制。记分板记录每个 warp 的每条指令的输入操作数是否已经准备好。比如全局内存 load 指令的结果还没返回时这条指令会卡在 scoreboard 阶段warp 被标记为 stalled。等到数据写回寄存器了scoreboard 更新状态warp 变成 eligible调度器才会把它纳入候选。这个机制和 CPU 里的动态调度有点类似但 GPU 的实现是分布在各个 warp 状态里的所有跟踪逻辑都是硬件电路没有任何软件参与。这也是为什么 GPU 调度能做到纳秒级别而操作系统进程调度要微秒甚至毫秒级别的原因。4. 线程块到 SM 的映射策略4.1 线程块如何分配到 SM前几节讲的是 warp 在 SM 内部的调度还有一个重要问题是线程块怎么分配到不同的 SM 上。这个分配是硬件完成的由GigaThread Engine全局调度器来负责。当你 launch 一个 kernel 时CUDA runtime 把整个 grid网格的任务描述交给 GigaThread Engine它会按顺序把一个一个线程块block分发到各个 SM 上。分配的依据是 SM 的资源容量寄存器和共享内存还够不够容纳一个新的 block如果够就分配如果不够就等当前在跑的 block 里有线程退出、资源释放了再分发新的 block 进来。这就是为什么你 launch 一个超大 grid比如一百万个 block时GPU 并不会一次全收下而是“边跑边补”。4.2 并发上限怎么算一个 SM 能装下多少 block这里给一个具体的计算例子。假设我在 Ampere 架构的 A100 上写 kernel每个 SM 最多 2048 个线程、32 个 block软件限制、65536 个寄存器、164KB 共享内存。我的 block 设成 256 线程每线程用 32 个寄存器按线程上限2048 ÷ 256 8 个 block按寄存器上限65536 ÷ (256 × 32) 8 个 block按 block 上限32 个 block三个限制取最小值所以这个配置下每个 SM 只能装 8 个 block换算成 warp 数就是 8 block × 8 warp/block 64 个 warp刚好顶到硬件上限occupancy 100%。如果把每线程寄存器数改成 64按寄存器上限65536 ÷ (256 × 64) 4 个 block这时候每个 SM 只能装 4 个 block总共 32 个 warpoccupancy 掉到 50%。如果内核本身没有太多寄存器压力只是编译器默认给了很宽的寄存器窗口那降低 occupancy 就得不偿失。你可以用__launch_bounds__(256, 8)提示编译器限制每线程寄存器数让它去 spill 一些非常用变量从而换取更高的 occupancy。4.3 block 大小和调度效率的微妙关系block 设得越大单个 block 包含的 warp 越多调度器在 block 内部可轮转的 warp 也越多。block 设得越小block 间切换更灵活更容易做到负载均衡。但是 block 太小也有问题假设一个 block 只有 32 线程1 个 warp那 SM 要管 64 个 block如果 occupancy 100%block 调度的簿记开销会变大而且每个 block 能用的共享内存量也少。我个人的经验是常规计算型 kernel 用 256 线程/block 是最不容易出错的起点因为它在延迟隐藏效率和块级负载均衡之间取得了一个比较均衡的位置。需要调优的时候再往 128 或 512 两个方向测试结合 profiling 数据做决定而不是拍脑袋硬改。5. 从 GPU 到 CUDA 软件栈的任务下发链路5.1 CPU 侧如何把任务交给 GPU你写 CUDA 代码的时候kernel launch 在 CPU 侧只做一件事把任务描述扔进命令缓冲区然后立即返回异步行为。GPU 驱动和硬件之间通过命令处理器Command Processor或叫 Host Interface沟通驱动把核函数入口地址、grid/block 尺寸、参数列表都写进一个叫control buffer的内存区域然后写一个门铃寄存器doorbell register告诉 GPU“有新任务了”。GPU 侧的工作调度器Work Distributor收到信号后开始读取控制缓冲区里的任务描述把 grid 拆成 block通过 GigaThread Engine 分发到各个 SM。整个过程都是硬件帮衬的CPU 只负责“下单”GPU 自己决定“怎么做、做多快”。实操心得如果你发现 kernel launch 的开销特别大用 Nsight Systems 能看到 launch 耗时几个微秒以上通常不是 GPU 慢而是 CPU 侧驱动栈和命令缓冲区的开销。这时候可以试试 CUDA Graph把多个 kernel 的依赖关系预先构建成一张图一次 launch 整张图能把启动开销摊薄到极致。这一点对短小 kernel 频繁调用的场景尤其有效。5.2 任务调度模型视角下的流Stream与事件EventCUDA 的Stream在任务调度模型里扮演什么角色你可以把它理解成一条“任务流水线”stream 之间的任务是可以并行的前提是硬件资源足够同一个 stream 内部的任务保持顺序执行。底层来看每个 stream 对应一条独立的命令队列GPU 的多个队列之间可以并发取任务。但要注意“可以并行”不等于“一定并行”如果两个 stream 的任务都要用到同一个 SM 上的全部资源那它们实际上还是串行的。事件Event则用于跨 stream 同步一个 stream 里可以cudaStreamWaitEvent(stream, event)强制这个 stream 等某个事件完成才开始后序任务。这本质上是把依赖关系告诉硬件调度器让它别做“优化排序”把有先后依赖的任务搞乱。5.3 MPS、多进程场景下的调度差异多进程共享 GPU 的情况和单进程多 stream 的调度模型有很大不同。MPSMulti-Process Service允许来自不同进程的 kernel 同时被调度到同一个 SM 上这相当于把 GPU 的调度单位从“进程”进一步细粒度化了。普通模式下多个进程要时间片轮转地独占整个 GPUMPS 模式下不同进程的 block 可以被交错调度到同一个 SM 上利用率显著提升但代价是单个 kernel 的延迟可能变高因为要和别人共享执行单元。如果你在做推理服务下游可能有多个模型实例MPS 几乎是把 GPU 利用率拉满的必备技能。6. 任务调度模型的性能关键参数与调优策略6.1 Occupancy 的计算与解读Occupancy 的定义一个 SM 上实际活跃的 warp 数占硬件最大容量的比例。它是衡量“资源利用充分度”最直观的指标。用 CUDA 自带的 occupancy calculator或cudaOccupancyMaxActiveBlocksPerMultiprocessorAPI可以直接查到给定 block 大小和寄存器/共享内存用量下的理论 occupancy。但我要给一句忠告不要把 occupancy 当唯一指标。我见过很多 kerneloccupancy 从 50% 提到 100%性能反而下降了。原因是高 occupancy 意味着每个线程能用的寄存器更少编译器不得不把局部变量吐到本地内存本地内存虽然走 L1/L2 缓存但比寄存器慢一两个数量级。这时候低 occupancy 但高寄存器命中率的配置反而更快。调优的正确路径应该是用 profiling 工具看 stall 原因memory dependency、execution dependency、barrier 等是“因为等内存而空转”就加 occupancy是“计算单元太挤”反而要降。6.2 常见 stall 原因与排查思路Nsight Compute 的 profiler 里你能看到每个 kernel 的 warp stall 周期分布。几个常见的 stall reason 对应的调优方向Long Scoreboard等待全局内存或共享内存数据返回。核心思路是提高内存访问局部性、用__ldg走只读缓存、加大 block 规模以提供更多独立 warp。Barrier在等待同一个 block 内其他 warp 到达 barrier如__syncthreads()。这说明 block 内同步太频繁或负载不均可以考虑减少同步次数、让每个线程做更多独立工作。MIO Throttle访问共享内存或特殊指令如 shuffle触发指令队列拥塞。可以降低共享内存访问密度把部分数据挪到寄存器。Not Selectedwarp 已就绪但调度器选了其他 warp。这个通常无害说明调度器有得选不算性能瓶颈。实操心得把Nsight Compute的 Sampling Data 打开按 stall reason 排序基本上一眼就能定位 kernel 的短板。这种“用数据说话”的方式比反复猜“是不是 block 大小不对”高效得多。我在调一个流体模拟 kernel 的时候一直以为瓶颈是全局内存带宽结果 profiling 一看是 barrier 同步太多改完同步策略性能提升 40%。6.3 调优实例从 60% occupancy 到 90% plus分享一个实际的调优案例。当时写一个金融领域的 Monte Carlo 模拟内核基础版本参数是 block256、每个线程 96 个寄存器occupancy 只有 37.5%8 个 block 上限被寄存器限制卡死实测带宽不到 40%。第一步用__launch_bounds__(256, 6)限制每线程寄存器数编译器把一些冷变量 spill 到本地内存寄存器降到 40 个。occupancy 提升到 75%但因为 spill 增加了内存流量性能只提升了 15%。第二步改代码结构把大多数 spill 的变量改成float精度原来是double并把部分数组改成手工分块放入共享内存。这一步将 spill 消除大半occupancy 达到 87.5%性能又提升了 32%。第三步调整 block 到 128每个 SM 的 block 数从 14 个升到 28 个受 block 上限限制实际活跃 warp 数不变但 block 调度更均衡最终性能比初版提升了 60% 以上。这个例子说明occupancy 只是表象真正的坑在资源使用的细粒度上。你要同时盯着寄存器、共享内存、spill 三条线才知道到底是谁限制了性能。7. 常见问题与排查技巧实录7.1 同一个 kernel 在不同 GPU 上表现天差地别换卡之后某些 kernel 效率骤降是特别常见的事。原因通常是不同架构的 SM 资源参数不同寄存器文件大小、共享内存大小、warp 调度器数量、甚至 warp 大小基本都是 32都可能变化。老卡上压出来的魔数参数比如 block512、每线程 64 寄存器直接搬到新卡上很可能把新卡的 occupancy 压得很低。排查思路跑一遍deviceQuery拿设备参数再重新走一遍 occupancy 计算。遇到这种情况代码里的硬编码参数应该抽成宏或参数对象按设备属性动态计算。7.2 Kernel launch 失败但没报错一种让人挠头的情况是kernel 调用之后后面的cudaMemcpy报“invalid argument”或直接卡死。常见原因是 block 数超过了硬件的 grid 上限x 方向一般是 2^31 - 1y/z 方向有更严格限制或者共享内存申请量超过了默认上限导致 launch 失败但错误码被忽略了。排查方法很简单每次 kernel launch 后检查返回值或者设置cudaDeviceSetLimit调大共享内存上限。这个坑我踩了不止一次奉劝大家 launch kernel 后立刻cudaGetLastError()成本几乎为零能救你几个小时。7.3 分支发散比想象中隐蔽有些发散不是 if-else 那种明显的写法而是循环的退出条件不同。同一个 warp 里 32 个线程执行一个 while 循环循环次数有的线程是 10 次、有的是 20 次那这个 warp 实际上要执行 20 次迭代前 10 次所有线程都活跃后 10 次只有一部分线程活跃。这种发散在 profiling 里不容易一眼看出来但计算量被白白拉高了。这种场景的优化方式是把循环拉平flatten或改成针对固定迭代次数的展开循环或者预计算每个线程的“终止位置”用掩码跳过多余迭代。实在无法解决时考虑用“分桶”的思路让相同循环次数的线程分到同一 warp。7.4 cudaDeviceSynchronize 卡住不动这个几乎都会碰到。它本身不是调度模型的锅但理解了调度就知道为什么它危险cudaDeviceSynchronize会让 CPU 阻塞等待 GPU 上所有已提交任务完成如果 GPU 上一个 kernel 是死循环或等待一个永远不满足的条件典型是 kernel 里访问未正确初始化的全局内存CPU 就会一直挂着。排查思路先用 cuda-gdb 或 nsight 把卡住的 kernel 找出来然后检查 kernel 内是否有循环条件永远为真、共享内存索引是否越界、以及是否误用了threadIdx和blockIdx的组合。不要小看这个我见过有人在if (blockIdx.x gridDim.x)这种永远是假的条件里加了死循环排查了两小时才找到。7.5 多 stream 并行时性能不升反降多 stream 并行没提速甚至更慢了大概率是 kernel 太小、CPU 侧 launch 开销占了大头或者是多个 stream 的 kernel 太大同时挤占同一个 SM 的资源。还有一种情况是 stream 之间没有足够独立的数据导致它们互相争抢 L2 缓存。排查思路用 Nsight Systems 看 stream 的时间轴确认各 stream 是不是真的在并行。用小 kernel 调大 stream 数测试找到“并行收益阈值”。检查 L2 缓存命中率命中率下降严重说明 stream 间的数据复用率太低这时候可以考虑单 stream 顺序执行。8. 需要纠正的几个错误认知8.1 线程数超过核心数就等于充分利用这是最常见的误解。核心数是固定的但 SM 里能驻留的线程数远大于物理核心数靠的就是调度器在 warps 之间来回切换。线程多只能说明“池子大”但池子里的线程如果都在等同一个资源比如同一块显存利用率照样上不去。真正应该关注的是“就绪的 warp 数”而不是“总线程数”。8.2 高 occupancy 一定等价于高性能前文已经反复提到了这里再强调一次。高 occupancy 可以掩盖延迟但代价是资源受限、寄存器溢出、甚至 cache 争抢加剧。低 occupancy 的内核如果每个线程的局部性特别好、指令级并行度高完全可能比高 occupancy 版本更快。指标服务于瓶颈而不是反过来。8.3 GPU 调度和 CPU 调度可以走同一套思路GPU 调度是硬件的、无操作系统介入的、极度高频的纳秒级CPU 调度是软件/内核的、有优先级策略的、相对低频的微秒以上。你把 CPU 多线程的调优手段比如锁、条件变量、线程池栈大小调优搬进 CUDA kernel基本全是反模式。GPU 上正确的“并发”方式是开足够多的独立线程、避免同步、减少资源占用。9. 后续还能深挖的方向这一篇说到底是从“模型”层面把 NVIDIA 任务调度的骨架梳理了一遍。再往下挖还有几块硬骨头值得单独写Volta 之后的独立线程调度从 Volta 开始NVIDIA 引入了独立线程调度Independent Thread Schedulingwarp 内不再是完全锁步这让以前的“隐式同步假设”失效也带来了一些新的同步语义比如__syncwarp的引入。这背后有一套更精细的程序计数器PC管理机制值得单独讲。持久线程Persistent Threads与手工负载均衡在一些大规模计算场景里让线程数量等于 SM 容量而非数据量然后靠循环自己拉取任务能做到更精细的负载均衡这算是对调度模型的“手动接管”。用户态任务图CUDA Graphs与流序调度面向低延迟推理场景把 kernel 启动开销压到接近零的做法调度模型会变成一张图而不是一条链。这些内容等后续有时间我再慢慢写。最后再分享一个实际的建议手头有 NVIDIA GPU 的朋友别光看这篇文章拿 Nsight Compute 跑一个自己以前的 kernel把 warp state 那一栏翻出来对照着这一篇的概念去看你会发现很多以前觉得“玄学”的性能问题其实都是调度模型里写得很明白的事情。
NVIDIA GPU任务调度模型：深入Warp与Occupancy的CUDA性能优化
GPU 调度这块我几年前刚开始认真写 CUDA kernel 的时候最直观的感受是明明照着官方文档把线程块大小设成了 256 occupancy 算出来也还行但性能就是上不去。后来把 NVIDIA 的任务调度模型真正啃了一遍才明白瓶颈到底卡在哪。这篇是这个系列的第二章前面聊完整体调度分析的大框架这次把 NVIDIA 的任务调度模型单独拉出来拆开揉碎讲清楚。先说清楚这篇要解决什么问题。很多人把“调度”理解成操作系统里那种进程/线程切换觉得 GPU 调度就是一堆线程排队上核心。这个理解错得很离谱。GPU 的调度是硬件级的、以线程束warp为最小单位的、基于状态机轮转的调度方式它压根不走操作系统内核也不需要什么上下文切换的软件开销。不理解这一层你写 kernel 的时候连“为什么 block 大小设成 128 有时比 256 好”这种问题都答不上来。这篇文章的受众我默认是对 CUDA 编程有一定基础、但没系统研究过 GPU 内部机制的开发者。如果你手头有 NVIDIA 的显卡跟着思路把 profiling 数据打开对照着看理解速度会快很多。下面先把最关键的几张“硬件底图”铺开。1. 先从硬件底图说起GPC、TPC 与 SM 的关系1.1 一张图理解 GPU 的物理结构分层NVIDIA GPU 的物理结构从上到下大概可以切成这么几层整个芯片叫 GPU芯片内部被划分成多个GPCGraphics Processing Cluster图形处理集群每个 GPC 内部包含若干个TPCTexture Processing Cluster纹理处理集群TPC 再往下拆就是SMStreaming Multiprocessor流式多处理器。每一层各管各的事GPC可以理解成芯片里的“大园区”它内部有独立的几何引擎、光栅化单元、以及若干个 SM。图形渲染任务里的三角形装配、光栅化这类活就是在 GPC 这个级别分工的。TPC一个 TPC 通常包含两个 SM老架构里也有一个 TPC 只挂一个 SM 的情况。TPC 主要起的是“分组管理”作用让纹理单元和 SM 之间的数据通路更紧凑。SM这才是真正的“计算工人”所有 CUDA 核心、Tensor Core、共享内存、寄存器堆、调度器全都住在 SM 里面。SM 是任务调度真正发生的地方。拿 Ampere 架构的 GA102 核心来举例它一共有 7 个 GPC、42 个 TPC、84 个 SM每个 SM 里有 128 个 CUDA 核心整卡总共就有 10752 个 CUDA 核心。这个数字看着唬人但你要记住这些 CUDA 核心并不是像 CPU 那样被“单个指令流”驱动它们是靠 SM 内部调度器喂数据才能干活的。这里插一句我踩过的坑早期我误以为一个 CUDA 核心对应一个线程那我开一万个线程就能让一万个核心满负荷跑。实际上线程和核心的关系是通过线程束和调度器间接建立的线程数超过核心数不代表就能并行执行得看调度器能喂多快。1.2 SM 内部的关键部件和调度相关寄存器SM 内部藏着几样和调度直接相关的“家当”我按重要性排个序线程束调度器Warp Scheduler每个 SM 里有 4 个Volta 之后普遍是 4 个这是调度的“大脑”负责从就绪的线程束里挑一个发指令。指令分发单元Dispatch Unit每个调度器旁边配一个负责把取到的指令真正发给执行单元。寄存器堆Register File每个 SM 里有 65536 个 32 位寄存器64K 是 Ampere 常见的配置。这些寄存器是给线程存局部变量的总容量固定所以线程越多、每个线程能分到的寄存器就越少。共享内存Shared Memory常见 128KB 或 164KB 每 SM用于线程块内部通信。Warp 状态表SM 内部维护一张表记录每个 warp 当前处于什么状态就绪、等待、阻塞等这是调度器做决策的依据。这些部件之间的关系我习惯用“餐厅后厨”来类比。GPC 是餐饮集团的一个门店TPC 是店里的几个档口SM 就是一个灶台warp 就是灶台上的炒锅。炒锅warp里装着一批菜线程灶台师傅调度器按照火候情况决定先翻炒哪口锅。每个灶台不止一口锅师傅也不止一个关键是谁能让锅不闲着。2. 任务调度的最小单位warp 的概念与意义2.1 为什么偏偏是 32 个线程NVIDIA 的硬件线程调度单位不是单个线程而是warp线程束一个 warp 恰好包含 32 个线程。这 32 个线程在硬件层面是“锁步”lockstep执行的也就是说同一个 warp 里的所有线程在同一时刻执行同一条指令只是处理的数据不同。为什么是 32 不是 16 也不是 64这里没有特别神秘的答案主要是硬件设计上的折中。32 个线程一条指令意味着指令 fetch 和 decode 的开销被 32 份工作分摊了这个倍数刚好能填满现代 GPU 执行单元的多周期流水线。历史上 Fermi 时代就定了 32 这个数后续架构一直沿用说明这个设计经过了充分的实践验证。你可以这样想CPU 一个核一次就处理一个线程的指令GPU 一个调度器一次处理 32 个线程的指令这就是 GPU 能靠“少控制、多并行”换取吞吐量的根本原因。代价是warp 内如果出现分支发散性能会断崖式下跌。2.2 分支发散为什么是性能杀手想象一个 warp 里有 32 个线程代码里写了一个 if-else 分支一半线程走 if一半走 else。硬件没法让这 32 个线程同时执行不同指令于是调度器只好先让走 if 的线程执行再让走 else 的线程执行最终这个 warp 需要串行执行两个分支的指令。这意味着什么本来一个周期能完成的指令现在要两个周期如果一个分支里嵌套了更多分支耗时还要成倍上涨。写 kernel 的时候最怕的不是分支本身而是同一个 warp 里出现“有的线程走 A、有的线程走 B”的分叉。代码层面怎么避免核心思想是让分支判断基于 warp 内统一的值而不是基于线程独有的数据。比如可以根据threadIdx.x / 32来分块处理让同一 warp 内的线程走同一条路径或者把分支逻辑改成算术运算用掩码来“选择”结果而不是真写 if。实操心得我早期写归约reduction核函数的时候用了if (tid % 2 0)这种写法做折叠操作结果性能比预期慢了一倍还多。后来改成让每个线程固定处理连续两段数据再配合 warp shuffle性能直接翻倍。分支发散这个坑真的是你不亲自拿 profiling 数据对比光靠感觉完全发现不了。2.3 线程束与线程块的换算关系一个线程块block由若干 warp 组成block 内的线程数必须是 warp 大小的整数倍吗硬件不强制但实践里强烈建议这样做。因为如果你把 block 大小设成 100硬件会把它补成 4 个 warp128 个线程其中 28 个线程是“空转”的它们不干任何活但照样占用调度资源和寄存器。所以你会看到业界几乎所有 kernel 的 block 大小都习惯性设成 32 的倍数128、256、512 都有但很少看见 100、300 这种数字。把 block 设成 warp 大小的整数倍是最基本的入门礼仪。3. SM 调度器的工作机制从取指到发射的完整链路3.1 四个调度器如何分工协作前面提到 Ampere 架构的 SM 里有 4 个 warp scheduler每个 scheduler 配一条独立的指令发射通路。这意味着一个 SM 在同一个时钟周期内最多能从 4 个不同的 warp 里各取一条指令发射出去。但这不意味着一个周期只能发射 4 条指令。NVIDIA 的调度器支持dual-issue也就是一个调度器在一个周期里可以发射两条指令前提是这两条指令互不依赖、并且执行单元有空闲。所以理论上一个 SM 一个周期最多可以发射 8 条指令4 个调度器 × 2 条。调度器怎么决定发射谁的指令底层逻辑是每个 warp 在每个周期都有“状态位”要么是“就绪可发射”eligible要么是“等待某种资源”stalled。调度器从 eligible 的 warp 里挑一个具体策略各家 NVIDIA 架构里略有不同常见的是轮转和优先级结合的方案。3.2 指令发射的延迟隐藏原理GPU 的调度器做一次“挑 warp 并发射指令”的动作在硬件层面只要几个时钟周期。但它真正的精髓在于用大量 ready 的 warp 来掩盖长延迟操作。比如从全局内存取数延迟可能有几百个周期CPU 的做法是停下等GPU 的做法是这个 warp 等着取数调度器立刻切换到另一个就绪的 warp 继续发射指令。这个机制叫latency hiding延迟隐藏。它要求 SM 里同时有足够多“蓄势待发”的 warp否则一旦所有 warp 都在等内存返回SM 就“空转”了也就是所谓的occupancy 不足。来算一笔账。一个 SM 有 4 个调度器每个调度器每个周期最多发射 2 条指令也就是每个周期最多需要 8 个 warp 供它轮转。假设每个 warp 的平均停顿周期数是 200等内存那 SM 里至少要同时有 8 × 200 1600 个 warp 在轮转才能把每周期 8 条发射管线喂满。而实际上 SM 里能容纳的 warp 总数是有限的——Ampere 架构最多 64 个 warp2048 线程per SM。这个远小于 1600所以现实是 GPU 只能做到“尽量隐藏延迟”不可能完全隐藏。这就是为什么occupancy 高不代表性能一定好但 occupancy 太低性能一定好不了。如果你的 kernel 每个线程用了太多寄存器导致 SM 里能同时常驻的 warp 数减少延迟隐藏能力就下降内存延迟就会暴露出来。3.3 调度器怎么知道你“准备好了”这个细节容易被忽略。每个 warp 的每条指令在进入执行流水线之前要经过所谓的scoreboard记分板机制。记分板记录每个 warp 的每条指令的输入操作数是否已经准备好。比如全局内存 load 指令的结果还没返回时这条指令会卡在 scoreboard 阶段warp 被标记为 stalled。等到数据写回寄存器了scoreboard 更新状态warp 变成 eligible调度器才会把它纳入候选。这个机制和 CPU 里的动态调度有点类似但 GPU 的实现是分布在各个 warp 状态里的所有跟踪逻辑都是硬件电路没有任何软件参与。这也是为什么 GPU 调度能做到纳秒级别而操作系统进程调度要微秒甚至毫秒级别的原因。4. 线程块到 SM 的映射策略4.1 线程块如何分配到 SM前几节讲的是 warp 在 SM 内部的调度还有一个重要问题是线程块怎么分配到不同的 SM 上。这个分配是硬件完成的由GigaThread Engine全局调度器来负责。当你 launch 一个 kernel 时CUDA runtime 把整个 grid网格的任务描述交给 GigaThread Engine它会按顺序把一个一个线程块block分发到各个 SM 上。分配的依据是 SM 的资源容量寄存器和共享内存还够不够容纳一个新的 block如果够就分配如果不够就等当前在跑的 block 里有线程退出、资源释放了再分发新的 block 进来。这就是为什么你 launch 一个超大 grid比如一百万个 block时GPU 并不会一次全收下而是“边跑边补”。4.2 并发上限怎么算一个 SM 能装下多少 block这里给一个具体的计算例子。假设我在 Ampere 架构的 A100 上写 kernel每个 SM 最多 2048 个线程、32 个 block软件限制、65536 个寄存器、164KB 共享内存。我的 block 设成 256 线程每线程用 32 个寄存器按线程上限2048 ÷ 256 8 个 block按寄存器上限65536 ÷ (256 × 32) 8 个 block按 block 上限32 个 block三个限制取最小值所以这个配置下每个 SM 只能装 8 个 block换算成 warp 数就是 8 block × 8 warp/block 64 个 warp刚好顶到硬件上限occupancy 100%。如果把每线程寄存器数改成 64按寄存器上限65536 ÷ (256 × 64) 4 个 block这时候每个 SM 只能装 4 个 block总共 32 个 warpoccupancy 掉到 50%。如果内核本身没有太多寄存器压力只是编译器默认给了很宽的寄存器窗口那降低 occupancy 就得不偿失。你可以用__launch_bounds__(256, 8)提示编译器限制每线程寄存器数让它去 spill 一些非常用变量从而换取更高的 occupancy。4.3 block 大小和调度效率的微妙关系block 设得越大单个 block 包含的 warp 越多调度器在 block 内部可轮转的 warp 也越多。block 设得越小block 间切换更灵活更容易做到负载均衡。但是 block 太小也有问题假设一个 block 只有 32 线程1 个 warp那 SM 要管 64 个 block如果 occupancy 100%block 调度的簿记开销会变大而且每个 block 能用的共享内存量也少。我个人的经验是常规计算型 kernel 用 256 线程/block 是最不容易出错的起点因为它在延迟隐藏效率和块级负载均衡之间取得了一个比较均衡的位置。需要调优的时候再往 128 或 512 两个方向测试结合 profiling 数据做决定而不是拍脑袋硬改。5. 从 GPU 到 CUDA 软件栈的任务下发链路5.1 CPU 侧如何把任务交给 GPU你写 CUDA 代码的时候kernel launch 在 CPU 侧只做一件事把任务描述扔进命令缓冲区然后立即返回异步行为。GPU 驱动和硬件之间通过命令处理器Command Processor或叫 Host Interface沟通驱动把核函数入口地址、grid/block 尺寸、参数列表都写进一个叫control buffer的内存区域然后写一个门铃寄存器doorbell register告诉 GPU“有新任务了”。GPU 侧的工作调度器Work Distributor收到信号后开始读取控制缓冲区里的任务描述把 grid 拆成 block通过 GigaThread Engine 分发到各个 SM。整个过程都是硬件帮衬的CPU 只负责“下单”GPU 自己决定“怎么做、做多快”。实操心得如果你发现 kernel launch 的开销特别大用 Nsight Systems 能看到 launch 耗时几个微秒以上通常不是 GPU 慢而是 CPU 侧驱动栈和命令缓冲区的开销。这时候可以试试 CUDA Graph把多个 kernel 的依赖关系预先构建成一张图一次 launch 整张图能把启动开销摊薄到极致。这一点对短小 kernel 频繁调用的场景尤其有效。5.2 任务调度模型视角下的流Stream与事件EventCUDA 的Stream在任务调度模型里扮演什么角色你可以把它理解成一条“任务流水线”stream 之间的任务是可以并行的前提是硬件资源足够同一个 stream 内部的任务保持顺序执行。底层来看每个 stream 对应一条独立的命令队列GPU 的多个队列之间可以并发取任务。但要注意“可以并行”不等于“一定并行”如果两个 stream 的任务都要用到同一个 SM 上的全部资源那它们实际上还是串行的。事件Event则用于跨 stream 同步一个 stream 里可以cudaStreamWaitEvent(stream, event)强制这个 stream 等某个事件完成才开始后序任务。这本质上是把依赖关系告诉硬件调度器让它别做“优化排序”把有先后依赖的任务搞乱。5.3 MPS、多进程场景下的调度差异多进程共享 GPU 的情况和单进程多 stream 的调度模型有很大不同。MPSMulti-Process Service允许来自不同进程的 kernel 同时被调度到同一个 SM 上这相当于把 GPU 的调度单位从“进程”进一步细粒度化了。普通模式下多个进程要时间片轮转地独占整个 GPUMPS 模式下不同进程的 block 可以被交错调度到同一个 SM 上利用率显著提升但代价是单个 kernel 的延迟可能变高因为要和别人共享执行单元。如果你在做推理服务下游可能有多个模型实例MPS 几乎是把 GPU 利用率拉满的必备技能。6. 任务调度模型的性能关键参数与调优策略6.1 Occupancy 的计算与解读Occupancy 的定义一个 SM 上实际活跃的 warp 数占硬件最大容量的比例。它是衡量“资源利用充分度”最直观的指标。用 CUDA 自带的 occupancy calculator或cudaOccupancyMaxActiveBlocksPerMultiprocessorAPI可以直接查到给定 block 大小和寄存器/共享内存用量下的理论 occupancy。但我要给一句忠告不要把 occupancy 当唯一指标。我见过很多 kerneloccupancy 从 50% 提到 100%性能反而下降了。原因是高 occupancy 意味着每个线程能用的寄存器更少编译器不得不把局部变量吐到本地内存本地内存虽然走 L1/L2 缓存但比寄存器慢一两个数量级。这时候低 occupancy 但高寄存器命中率的配置反而更快。调优的正确路径应该是用 profiling 工具看 stall 原因memory dependency、execution dependency、barrier 等是“因为等内存而空转”就加 occupancy是“计算单元太挤”反而要降。6.2 常见 stall 原因与排查思路Nsight Compute 的 profiler 里你能看到每个 kernel 的 warp stall 周期分布。几个常见的 stall reason 对应的调优方向Long Scoreboard等待全局内存或共享内存数据返回。核心思路是提高内存访问局部性、用__ldg走只读缓存、加大 block 规模以提供更多独立 warp。Barrier在等待同一个 block 内其他 warp 到达 barrier如__syncthreads()。这说明 block 内同步太频繁或负载不均可以考虑减少同步次数、让每个线程做更多独立工作。MIO Throttle访问共享内存或特殊指令如 shuffle触发指令队列拥塞。可以降低共享内存访问密度把部分数据挪到寄存器。Not Selectedwarp 已就绪但调度器选了其他 warp。这个通常无害说明调度器有得选不算性能瓶颈。实操心得把Nsight Compute的 Sampling Data 打开按 stall reason 排序基本上一眼就能定位 kernel 的短板。这种“用数据说话”的方式比反复猜“是不是 block 大小不对”高效得多。我在调一个流体模拟 kernel 的时候一直以为瓶颈是全局内存带宽结果 profiling 一看是 barrier 同步太多改完同步策略性能提升 40%。6.3 调优实例从 60% occupancy 到 90% plus分享一个实际的调优案例。当时写一个金融领域的 Monte Carlo 模拟内核基础版本参数是 block256、每个线程 96 个寄存器occupancy 只有 37.5%8 个 block 上限被寄存器限制卡死实测带宽不到 40%。第一步用__launch_bounds__(256, 6)限制每线程寄存器数编译器把一些冷变量 spill 到本地内存寄存器降到 40 个。occupancy 提升到 75%但因为 spill 增加了内存流量性能只提升了 15%。第二步改代码结构把大多数 spill 的变量改成float精度原来是double并把部分数组改成手工分块放入共享内存。这一步将 spill 消除大半occupancy 达到 87.5%性能又提升了 32%。第三步调整 block 到 128每个 SM 的 block 数从 14 个升到 28 个受 block 上限限制实际活跃 warp 数不变但 block 调度更均衡最终性能比初版提升了 60% 以上。这个例子说明occupancy 只是表象真正的坑在资源使用的细粒度上。你要同时盯着寄存器、共享内存、spill 三条线才知道到底是谁限制了性能。7. 常见问题与排查技巧实录7.1 同一个 kernel 在不同 GPU 上表现天差地别换卡之后某些 kernel 效率骤降是特别常见的事。原因通常是不同架构的 SM 资源参数不同寄存器文件大小、共享内存大小、warp 调度器数量、甚至 warp 大小基本都是 32都可能变化。老卡上压出来的魔数参数比如 block512、每线程 64 寄存器直接搬到新卡上很可能把新卡的 occupancy 压得很低。排查思路跑一遍deviceQuery拿设备参数再重新走一遍 occupancy 计算。遇到这种情况代码里的硬编码参数应该抽成宏或参数对象按设备属性动态计算。7.2 Kernel launch 失败但没报错一种让人挠头的情况是kernel 调用之后后面的cudaMemcpy报“invalid argument”或直接卡死。常见原因是 block 数超过了硬件的 grid 上限x 方向一般是 2^31 - 1y/z 方向有更严格限制或者共享内存申请量超过了默认上限导致 launch 失败但错误码被忽略了。排查方法很简单每次 kernel launch 后检查返回值或者设置cudaDeviceSetLimit调大共享内存上限。这个坑我踩了不止一次奉劝大家 launch kernel 后立刻cudaGetLastError()成本几乎为零能救你几个小时。7.3 分支发散比想象中隐蔽有些发散不是 if-else 那种明显的写法而是循环的退出条件不同。同一个 warp 里 32 个线程执行一个 while 循环循环次数有的线程是 10 次、有的是 20 次那这个 warp 实际上要执行 20 次迭代前 10 次所有线程都活跃后 10 次只有一部分线程活跃。这种发散在 profiling 里不容易一眼看出来但计算量被白白拉高了。这种场景的优化方式是把循环拉平flatten或改成针对固定迭代次数的展开循环或者预计算每个线程的“终止位置”用掩码跳过多余迭代。实在无法解决时考虑用“分桶”的思路让相同循环次数的线程分到同一 warp。7.4 cudaDeviceSynchronize 卡住不动这个几乎都会碰到。它本身不是调度模型的锅但理解了调度就知道为什么它危险cudaDeviceSynchronize会让 CPU 阻塞等待 GPU 上所有已提交任务完成如果 GPU 上一个 kernel 是死循环或等待一个永远不满足的条件典型是 kernel 里访问未正确初始化的全局内存CPU 就会一直挂着。排查思路先用 cuda-gdb 或 nsight 把卡住的 kernel 找出来然后检查 kernel 内是否有循环条件永远为真、共享内存索引是否越界、以及是否误用了threadIdx和blockIdx的组合。不要小看这个我见过有人在if (blockIdx.x gridDim.x)这种永远是假的条件里加了死循环排查了两小时才找到。7.5 多 stream 并行时性能不升反降多 stream 并行没提速甚至更慢了大概率是 kernel 太小、CPU 侧 launch 开销占了大头或者是多个 stream 的 kernel 太大同时挤占同一个 SM 的资源。还有一种情况是 stream 之间没有足够独立的数据导致它们互相争抢 L2 缓存。排查思路用 Nsight Systems 看 stream 的时间轴确认各 stream 是不是真的在并行。用小 kernel 调大 stream 数测试找到“并行收益阈值”。检查 L2 缓存命中率命中率下降严重说明 stream 间的数据复用率太低这时候可以考虑单 stream 顺序执行。8. 需要纠正的几个错误认知8.1 线程数超过核心数就等于充分利用这是最常见的误解。核心数是固定的但 SM 里能驻留的线程数远大于物理核心数靠的就是调度器在 warps 之间来回切换。线程多只能说明“池子大”但池子里的线程如果都在等同一个资源比如同一块显存利用率照样上不去。真正应该关注的是“就绪的 warp 数”而不是“总线程数”。8.2 高 occupancy 一定等价于高性能前文已经反复提到了这里再强调一次。高 occupancy 可以掩盖延迟但代价是资源受限、寄存器溢出、甚至 cache 争抢加剧。低 occupancy 的内核如果每个线程的局部性特别好、指令级并行度高完全可能比高 occupancy 版本更快。指标服务于瓶颈而不是反过来。8.3 GPU 调度和 CPU 调度可以走同一套思路GPU 调度是硬件的、无操作系统介入的、极度高频的纳秒级CPU 调度是软件/内核的、有优先级策略的、相对低频的微秒以上。你把 CPU 多线程的调优手段比如锁、条件变量、线程池栈大小调优搬进 CUDA kernel基本全是反模式。GPU 上正确的“并发”方式是开足够多的独立线程、避免同步、减少资源占用。9. 后续还能深挖的方向这一篇说到底是从“模型”层面把 NVIDIA 任务调度的骨架梳理了一遍。再往下挖还有几块硬骨头值得单独写Volta 之后的独立线程调度从 Volta 开始NVIDIA 引入了独立线程调度Independent Thread Schedulingwarp 内不再是完全锁步这让以前的“隐式同步假设”失效也带来了一些新的同步语义比如__syncwarp的引入。这背后有一套更精细的程序计数器PC管理机制值得单独讲。持久线程Persistent Threads与手工负载均衡在一些大规模计算场景里让线程数量等于 SM 容量而非数据量然后靠循环自己拉取任务能做到更精细的负载均衡这算是对调度模型的“手动接管”。用户态任务图CUDA Graphs与流序调度面向低延迟推理场景把 kernel 启动开销压到接近零的做法调度模型会变成一张图而不是一条链。这些内容等后续有时间我再慢慢写。最后再分享一个实际的建议手头有 NVIDIA GPU 的朋友别光看这篇文章拿 Nsight Compute 跑一个自己以前的 kernel把 warp state 那一栏翻出来对照着这一篇的概念去看你会发现很多以前觉得“玄学”的性能问题其实都是调度模型里写得很明白的事情。
