CUDA Thread Block 调度机制详解：从硬件视角优化 Kernel 性能

写 CUDA 程序写久了你一定会遇到一个问题我明明启动了 1000 个 thread blockGPU 到底按什么顺序把它们分配到 SM 上不同 block 是同时跑还是排队跑为什么有时候 block 数量翻倍性能反而更差这些问题都指向同一个核心CUDA 如何调度 Thread Block。这个知识点不搞透你写出来的 kernel 永远只能“能跑”很难“跑得快”。这篇文章我会从硬件调度器的视角拆开讲包括 block 是怎么被分到 SM 上的、资源上限如何卡住并发、调度顺序到底能不能依赖、以及怎么利用调度规律做优化。适合已经写过 CUDA kernel、但还没系统理解 GPU 内部行为的人。看完之后你会对 kernel 的执行过程有一个完全不同的认识。1. 先从两个层级看调度软件视角与硬件视角1.1 为什么 Thread Block 是调度的最小单位CUDA 的软件编程模型里一个 kernel 启动后对应一个 gridgrid 由一堆 thread block 组成每个 block 里又有一堆 thread。如果你只写过几个小 kernel可能觉得这只是“为了方便组织并行任务”而设计的层级。但从硬件的角度看Thread Block 不仅仅是一个代码组织单位它更是 GPU 调度的最小分配单位。CPU 上操作系统调度的是线程GPU 上硬件调度器把一个 block 当作一个整体要么整个 block 被分配到某个 SMStreaming Multiprocessor上要么就等着。block 内部的线程可以共享 shared memory、可以用__syncthreads()做块内同步这些都是 block 能作为一个调度单位的前提。GPU 不需要知道你的 block 里每个线程具体在算什么它只需要知道这个 block 占多少线程、多少寄存器、多少共享内存然后决定能不能把它塞进某个 SM。这里有个很容易误解的点GPU 真正执行指令的最小单位是 warp也就是 32 个线程不是 block。那为什么说“block 是调度的最小单位”因为 warp 不能跨 block 组合调度器不会把两个 block 里的线程凑成一个 warp。所以 block 虽然内部会被拆成多个 warp 来执行但 block 的“分配”和“驻留”却是整体性的。可以这样理解block 是一个集装箱warp 是集装箱里的货架GPU 的吊机一次只能吊一个集装箱但货架上的货可以在 SM 上单独搬运。1.2 SM 才是真正的执行现场任何一个 thread block最终都要落到某一个 SM 上执行。SM 是 GPU 内部最基本的计算核心单元它里面有一组 CUDA core、一块寄存器堆、一块 shared memory还有若干 warp scheduler。一个 GPU 上有几十甚至上百个 SM具体数量取决于型号。你在编程时看到的“GPU 核心数”其实是所有 SM 上的 CUDA core 总数但调度的视野是“SM 数量 × 每个 SM 能驻留多少 block”。每个 SM 能同时驻留的 block 数量不是无限大的。硬件会给每个 SM 设置几个硬性上限最大线程数、最大 block 数、最大共享内存、最大寄存器数。调度器在分配 block 时会检查这个 SM 当前的资源余量只要任何一个资源不够这个 block 就不能进入该 SM。多个 block 可以被分配到同一个 SM只要它们加起来不超出这些上限。所以你从软件层看到的 grid 是一个扁平结构所有 block 都是逻辑上并行的但从硬件层看每个 block 是在某一个 SM 上等着被调度它能不能立刻执行取决于那个 SM 当时的资源状态。理解了这一点后面所有关于调度的讨论就有了基础。2. 调度器的工作方式从 GigaThread Engine 到 SM 的空槽2.1 硬件工作分配器如何“发牌”当你在 host 端调用 kernel launch 时驱动会把 kernel 的启动信息提交给 GPUGPU 内部有一个硬件分发单元很多资料里称它为 GigaThread Engine。它的作用就是把 grid 里的 block 逐个分发到各个 SM 上。从宏观上看它像一个发牌员桌上有很多个 SM 座位每个座位有一定资源容量发牌员会按顺序往空位上放 block。这类硬件设计最早可以追溯到 Fermi 架构后来在 Volta、Ampere、Hopper 等架构上都有演进。不同架构的细分实现有差异但核心逻辑一直没变调度器会尝试让所有 SM 都“吃饱”而不是等某个 SM 上的 block 全部执行完再发下一批。也就是说一个 SM 只要还有足够的空闲资源调度器就可能往上面再塞一个新 block即使这个 SM 上已经有多个 block 在运行。这就是“流式调度”的核心。GPU 不会等 grid 里的一部分 block 跑完再把剩余 block 全部发下去。而是哪个 SM 有位置新的 block 就过去。这样能最大程度保证计算资源不被浪费。实际表现就是在 kernel 启动后的极短时间内所有 SM 都会被尽可能填满后续如果有 block 因为资源不足而排队它们会等待某个 SM 上的 block 完成再被分配进去。2.2 资源约束是硬天花板现在关键问题来了一个 SM 到底能同时驻留多少个 block这完全由四个因素决定。第一个是线程数。每个 SM 有最大线程数限制。比如很多现代 GPU 的 SM 最多支持 2048 个线程如果你的 block 大小是 512那么最多驻留 4 个 block如果 block 大小是 1024最多驻留 2 个。第二个是 block 数量上限。硬件还会限制每个 SM 最多驻留多少个 block通常是 16 或 32。假如你的 block 大小只有 32 个线程理论上 2048 / 32 64 个 block但如果硬件上限是 32那你也只能驻留 32 个。第三个是寄存器。每个线程用多少寄存器直接影响 block 能否被接纳。一个 SM 通常有 64K 个 32 位寄存器如果你的 kernel 每个线程用 40 个寄存器block 大小 256那一个 block 就要 256 × 40 10240 个寄存器64K 最多能放 6 个 block。这个限制有时比线程数更早卡住你。第四个是 shared memory。每个 SM 的 shared memory 容量有限比如 96KB 或 164KB如果每个 block 申请 40KB那最多驻留 2 个 block。哪怕线程数和寄存器都很宽裕shared memory 也会成为瓶颈。这四个限制必须同时满足。调度器在决定“这个 block 能不能进这个 SM”时取的是这四个限制的交集。这个逻辑很像停车场的闸机车位数量、入口宽度、限高、承重哪个不满足都进不去。2.3 占用率Occupancy到底是什么把这些资源限制综合起来就引出了 CUDA 优化里最常听到的指标Occupancy占用率。它一般定义为“每个 SM 上实际活跃的 warp 数”除以“每个 SM 最大能容纳的 warp 数”。为什么要关心占用率因为 GPU 靠大量线程并行来隐藏延迟。当一个 warp 在等待内存访问结果时调度器会切换到其他 warp 执行。如果 SM 上驻留的 warp 太少内存等待期间计算单元就会空转性能自然上不去。高占用率不一定等于高性能但低占用率几乎一定意味着性能受限除非你的 kernel 已经完全是计算密集且没有太多访存停顿。用前面寄存器限制的例子来算一下某 SM 最大 2048 线程最大 64 个 warp。如果 block 大小 256寄存器每线程 40 个寄存器上限 65536 个每个 block 需要 10240 个寄存器那么驻留 block 数 65536 / 10240 6 个对应线程 1536warp 数 48。如果该 SM 最大 warp 数是 64那么占用率就是 48 / 64 75%。而如果只看线程数上限2048 / 256 8也就是 100% 占用率但因为寄存器不够实际只能到 75%。这就是为什么不能只凭“线程数上限”估算并行度。NVIDIA 提供运行时 API 可以直接算出这个数字后面我会专门讲。现在先记住调度器做决策的依据就是这些资源限制而占用率只是这些限制综合之后的一个“结果”。3. 看似“随机”的调度顺序其实隐藏着规则3.1 Block ID 与执行顺序在大多数情况下GigaThread Engine 会按照 blockIdx 从小到大的顺序把 block 依次分发到 SM 上。你可以用clock64()在每个 block 的入口记录时间戳会发现前几十个 block 几乎同时启动后面 block 的启动时间会错开。这印证了一个事实block 不是一起开始的而是按资源空位分批进入的。第一个值得注意的点是一个 block 一旦被分配到某个 SM它就会一直待在那里直到执行完毕。它不会像 CPU 线程那样被换出或迁移到另一个 SM。所以你无法通过代码把某个 block“挪”到其他地方只能等它结束然后腾出资源给新 block。这也是为什么 grid 内 block 数量远远大于 SM 可驻留 block 数量时末尾会出现“长尾效应”——最后一批 block 要等前面的 block 陆续结束才能开始而这个“陆续结束”的时间差就是尾部延迟。第二个值得注意的点是不要依赖 block 的执行顺序。即使 block 通常按 ID 顺序分配但每个 block 里的线程执行速度和完成时间完全不可预测。某次运行里 block 5 可能先完成下次运行可能 block 7 先完成。如果 kernel 内部有多个 block 之间的隐式依赖你只能在 host 端启动多个 kernel或使用原子操作和全局内存同步。任何“假设 block 0 先跑完再轮到 block 1”的写法都是 bug 隐患。3.2 块间同步与协作调度既然 block 之间没有天然同步机制那 CUDA 生态里有没有办法让多个 block 互相等待有但你必须调用 cooperative groups 里的 grid sync 能力而不是普通 kernel launch。普通的 kernel launch 并不保证 grid 内所有 block 同时驻留。假设你的 grid 有 1000 个 block但某个 SM 只容得下 4 个 block那么同一时刻最多只有几十个 block 在跑。如果你在代码里让一个 block 等另一个 block 的数据而这个“另一个 block”还没被调度进来就会死锁。因此普通 launch 不支持 grid 级别同步。cooperative launch 则不同。它会提前计算资源保证整个 grid 的 block 数量不超过所有 SM 的驻留上限也就是说所有 block 都能同时驻留在 GPU 上。只有在这种前提下grid.sync()这样的块间同步才是安全的。这个机制对调度的影响非常直接它把“网格被分批调度”变成了“网格一次性全部驻留”代价是你的 grid 大小必须被限制在硬件能容纳的范围不能想开多少就开多少。如果你没有用 cooperative launch就不要在 block 之间做任何形式的忙等待同步。就算你用原子变量和全局内存写了一个类似 barrier 的逻辑也会因为调度顺序不确定而面临死锁风险。这是我见过新手最容易踩的坑代码逻辑上看起来是对的但跑起来时不时卡死。4. 调度行为怎么影响性能调优实战4.1 占用率不是越高越好前面讲占用率很多人的第一反应就是“把占用率拉满”。实际调优时你会发现100% 占用率并不总是最优解。一个常见场景是 shared memory 使用量大的 kernel。如果你把 block 大小从 256 降到 128每个 block 占用的 shared memory 也相应减少SM 上就能驻留更多 block占用率确实会提高。但 block 变小意味着每个 block 能处理的数据量变少块边界的同步和加载开销占比上升最后总时间反而变长。反过来有些 kernel 会在每个线程里使用大量寄存器。你把 block 大小设得很小虽然 block 数量上限允许更多 block 驻留但寄存器不够占用率被压低。此时你可以通过__launch_bounds__限制每个线程的最大寄存器数牺牲少量寄存器溢出换取更高占用率。这种操作能不能带来收益只能实测不能拍脑袋。我自己的经验是先把占用率从 50% 提到 75%一般会看到明显收益但从 75% 再往上收益会变得很微弱甚至出现反效果。因为更高的占用率意味着每个 thread block 分到的资源更少可能触发更多的 local memory 访问和缓存抖动。所以调优时不要只盯着 occupancy 这个单一数字要配合 Nsight Compute 看 stall 原因。如果 warps 是因为“等待数据”而停顿那提高占用率往往有效如果是因为“计算指令依赖”而停顿提高占用率帮助不大。4.2 Thread Block 大小的选择策略没有万能 block size但有无脑安全的选择。大多数 kernel 用 128 或 256 个线程是比较稳定的起点。它们都是 32 的整数倍不会产生“半个 warp”的浪费资源利用率又不会因为 block 太大而变得僵化。选 block size 的时候建议从两个方向考虑。第一个方向是“你想让 GPU 同时处理多少数据”。如果每个线程处理一个元素那 block 大小就是“每个调度单位负责的元素数”。block 太大比如 1024单个 block 占用的资源多SM 驻留 block 数量少调度灵活性低block 太小比如 32虽然调度粒度细但 block 内部共享数据的优势就没了而且会占用更多 block slot。第二个方向是“你的算法在 block 内是否需要大量协作”。比如矩阵分块乘法和卷积这类算法通常希望 block 能覆盖一个较大的 tile这样便于用 shared memory 复用数据。这时 block 大一点更好。如果你只是做一个逐元素的向量加或者归约block 大小在 128 到 512 之间都合理重点是让总 block 数足够多方便调度器在 SM 之间均衡分配。还有一个容易忽略的点block 总数不能太少。假设你的 GPU 有 80 个 SM每个 SM 最多驻留 8 个 block你只启动了 80 个 block那平均每个 SM 刚好分到 1 个 block。如果这 80 个 block 工作量不均匀负载均衡会很差某些 SM 忙死某些 SM 闲死。这时候把 block 切小一点比如启动 640 个 block每个 SM 可以有 8 个 block调度器就能在完成后动态补充尾部效应会明显缓解。4.3 用 Persistent Kernel 主动接管调度前面说的都是“把 block 交给硬件调度器由它决定谁先跑”。但在某些场景下你会希望自己控制 block 和数据的对应关系。最典型的就是 persistent kernel也就是常驻 kernel。做法是先通过资源计算确定“每个 SM 最多能驻留几个 block”然后让 grid 的大小刚好等于“SM 数量 × 每 SM 驻留 block 数”不多不少。kernel 启动后这些 block 一直驻留在 SM 上不再有新的 block 被调度进来也不会有旧的 block 退出。然后每个 block 自己进入一个循环从全局任务列表里取数据块来处理处理完了再取下一个。这种模式最适合任务量很大、且每个小任务执行时间不固定的场景比如稀疏矩阵计算、图数据处理。它把“硬件调度 block”变成了“软件调度任务”让开发者可以按需实现优先级、动态负载均衡。代码上通常用一个 atomic counter 当作全局任务指针每个 block 循环用atomicAdd获取下一个任务的起始位置。优点很明显block 反复启动和退出带来的硬件调度开销被降到最低负载也能自动均衡。缺点是程序复杂度变高而且你必须小心处理 shared memory 的生命周期。因为 persistent kernel 里的 block 不会退出shared memory 在循环过程中由你手动重置不能指望每次 .block 重新开始就自动清空。还有如果每个任务之间有隐式依赖需要额外加同步否则容易出竞态问题。4.4 原子操作与动态调度的微妙关系在 persistent kernel 里我们用一个原子计数器来分配任务。这里顺带说一个调度相关的问题如果你在普通 kernel 里给每个 block 分配固定区间比如 block i 处理i * chunk_size到(i 1) * chunk_size那 workload 不均匀只能靠硬件调度器来缓解。但如果你用原子操作从全局队列里取任务那么每个 block 实际分配到的工作是动态决定的调度器的行为就变得不那么重要了。这个思路同样可以用在非 persistent kernel 上。比方说你的数据被切成 10000 个小任务但你只启动了 256 个 block每个 block 循环取任务。这样 block 数量少但每个 block 里逻辑复杂shared memory 复用更好。硬件调度器面对这种 kernel 时不需要频繁分配新 block任务分配完全由原子计数器控制所以并发更可控负载更均衡。但要注意原子操作在全局内存上是有竞争开销的。如果任务粒度太小比如每个任务只算几个浮点运算那么线程大部分时间都在抢计数器性能会非常难看。我的经验是任务粒度至少要大到“原子分配开销 同步开销”远小于任务本身计算时间。这个比例很难凭空估计通常是先写一版固定分配再写一版动态分配对比之后再做选择。5. 常见问题与排查技巧实录5.1 如何查看真实的 Occupancy很多人在 kernel 启动前就宣称“我的占用率是 100%”实际上这个数字往往是“理论占用率”而不是实际达到的占用率。理论值只是根据资源限制算出来的上限实际运行中可能因为分支发散、共享内存 bank 冲突、线程提前退出等原因导致部分 warp 处于 inactive 状态。最直接的方式是用 Nsight Compute 这类性能分析工具。在 Occupancy 页面里你会看到两套数字Block Limit 和 Achieved Occupancy。前者是硬件和 kernel 资源共同决定的理论上限后者是实测值。如果两者差距很大说明你的 kernel 没有把驻留的 warp 用满可能是线程调度不均匀也可能是有大量线程在做等待。没有可视化工具时也可以用 CUDA runtime API 在 host 端预先计算。最常用的是cudaOccupancyMaxActiveBlocksPerMultiprocessor它接受 kernel 函数指针、block 大小和共享内存字节数返回每个 SM 能驻留的最大 block 数。然后你再结合设备属性里的 SM 数量就能算出 grid 内 block 总数是否合理。举个例子int numBlocks; int blockSize 256; cudaOccupancyMaxActiveBlocksPerMultiprocessor( numBlocks, myKernel, blockSize, 0); int deviceId; cudaGetDevice(deviceId); cudaDeviceProp prop; cudaGetDeviceProperties(prop, deviceId); int totalBlocks numBlocks * prop.multiProcessorCount;这样得到的totalBlocks就是“你的 GPU 上同时最多只能跑这么多 block”。如果 grid 明显大于它就要清楚剩余 block 会排队等待启动时间会存在长尾。5.2 问题速查表下面这张表是我在实际项目里总结出来的遇到调度相关性能问题可以先对照看看。现象可能原因排查思路kernel 启动时间很长但计算时间不长block 数量过多或资源不足导致排队检查 block 大小和占用率适当增大 block 或减小共享内存多个 SM 占用率显示差异大数据划分不均匀或者 block 数量太少改用更多小 block或用原子计数器做动态负载均衡理论上占用率 100%实际性能上不去访存停顿或分支发散严重用 Nsight Compute 看 stall 原因别急着调 block size程序偶尔死锁普通 kernel 内做了块间同步改用 cooperative launch或者去掉块间同步同一 kernel 多次运行耗时起伏大块间调度顺序导致的缓存局部性差异尽量让每个 block 的数据访问更独立减少跨 block 共享数据shared memory 大时占用率骤降每 block 共享内存申请过多考虑改动态共享内存或减少每个 block 的 tile 尺寸5.3 调试时不要依赖 Block 执行顺序有一个很坑的现象当你用printf在 device 端打印每个 block 的 blockIdx.x然后在 host 端观察输出顺序时看到的经常不是你以为的“block 0、1、2、3”顺序。因为打印输出本身有缓冲不同 block 的执行完成顺序不同输出顺序自然也不可靠。更靠谱的方法是让每个 block 把开始时间记录到全局数组里放在 kernel 结束前写回再用clock64()获取 SM 时钟周期。这样能看到每个 block 实际开始执行的时间点。但注意同 SM 上多个 block 是交错执行的clock64()的启动时间只能说明“这个 block 的代码从何时开始跑”不能说明“这个 block 是否一直占用 SM”。还有一种调试陷阱为了“看清调度顺序”在 kernel 里让每个 block 忙等不同的时间然后观察结果。这种做法会彻底改变调度行为因为你人为把每个 block 的执行时间拉长了资源释放顺序和正常情况完全不同。得出的结论没有参考价值。调试调度行为时最好使用不加额外干扰的 profile 工具或者只做轻量的时间戳记录。5.4 几个经验参考值从我做过的各种 kernel 来看以下经验值可以当作起点。通用计算、访存密集block size 128 或 256共享内存尽量控制在“能让 occupancy 超过 50%”的水平。计算密集、复用率高block size 256 到 512共享内存可以多用一些不必强求高 occupancy。任务粒度很小、需要动态负载均衡block 数量可以设为 SM 数量的整数倍然后持久化运行内部循环取任务。归约类 kernel先用每个 block 处理一部分数据最后用原子操作合并block size 用 128 或 256 时warp 内归约效率比较高。有__syncthreads()的 kernelblock 越大同步等待成本越高如果每次同步之间计算量不多建议 block 取小值。这些参考值不一定适合所有场景但可以帮你快速定位“是 block 大小问题还是调度模式问题”。6. 关于调度最后分享一点实际经验我对 CUDA 调度最深的体会是硬件调度器比大多数开发者想象的更“机械”它不关心你的 block 之间有没有逻辑依赖只看资源空位够不够。所以与其花时间猜调度顺序不如把资源占用算清楚把任务划分做得足够灵活。如果你刚开始接触这块我建议第一件事不是改代码而是先用工具把一个你现有的 kernel 跑一遍记录它的理论占用率和实际占用率再算一下理论最大并发 block 数与你实际 block 数的比例。这一步就能解释很多性能问题的根源。之后再去调整 block size、共享内存或者切换 persistent kernel 模式你会发现自己不再“盲调”而是真正知道硬件在干什么。最后再给一个建议每次调整调度相关参数后至少跑三次取中位数不要用单次耗时就下结论。GPU 调度行为受驱动版本、数据布局、甚至时钟频率的影响单次测量很容易误导人。把测试流程稳定下来你会少走很多弯路。