企业官网建设流程全解析

刚接触CUDA那阵子，我一直把Thread Block当成一个“程序员用来切数据”的抽象概念，启动参数填好，Grid和Block按公式分开就完事。直到有一次调一个粒子类模拟kernel，Nsight Compute里显示SM并不缺活跃warp，但整体耗时比理论估算高了快一倍。查到最后，问题不是指令效率，而是Thread Block在GPU上的调度方式出了岔子——一部分SM忙到排队，另一部分SM却空着看热闹。这篇文章把当时摸清楚的东西整理一遍：CUDA究竟怎么把Thread Block从Grid派发到SM，Block进入SM之后Warp又怎么被调度，这些机制又会怎样直接反映在占用率和运行时间上。适合已经会写基本kernel、但经常被Sm占用率偏低和负载不均困扰的人。

1. 先纠正一个常见误解：Block不是铺在GPU上均匀执行的

1.1 Block是一个整体，只能完整落在一个SM上

很多初学CUDA的人会下意识认为：既然GPU有那么多核心（或者说执行单元），那我启动1024个线程，硬件就会把这1024个线程平均摊给所有核心去跑。

实际完全不是这样。

Thread Block是一个“不可拆分”的执行单位。启动kernel之后，硬件的工作分发器会把整个Block派发给某一个SM，这个Block里的所有线程只在这个SM上执行，共享内存、寄存器分配、barrier同步这些资源也全部落在同一个SM上。一个Block永远不会被拆成两半扔到不同的SM上。

原因很直接：共享内存和线程间同步都要求线程在物理上靠在一起。Block内部之所以能用__syncthreads()做同步，能用__shared__共享数据，就是因为硬件保证这些线程在同一个SM内部，访问共享内存就是访问SM本地的物理存储，延迟只有几个周期。如果Block被切散到两个SM，同步和共享就变成跨SM操作，代价高得不可接受。

1.2 blockIdx的顺序，不是硬件执行顺序

编程模型里，Block的编号从0到gridDim-1，看起来很有序。但硬件在派发Block时并不会严格按blockIdx从小到大地执行。

我自己做过一个测试：在kernel入口处用clock64()记录每个Block开始执行的时钟周期，然后按blockIdx打印。结果多次运行时，先被调度的Block ID是乱序的，不同Block的启动时间差别很大。换句话说，代码里绝不能依赖“blockIdx小的Block一定先跑完”这种假设，Grid内部的执行顺序对程序员是不可观察、也不应依赖的。

这一点在写跨Block协作算法时尤其重要。比如某个Block要等另一个Block算完才能继续，你不能在kernel里直接假设顺序，必须用原子操作、全局同步或独立Grid之类的机制去显式同步。

1.3 硬件为什么坚持按Block切而不是按线程切

按Block分配会带来一个副作用：资源碎片化。比如一个SM只能驻留8个Block，哪怕每个Block只用了少量线程，SM上剩余的线程名额也未必能被别的Block利用。

但硬件依然选择Block作为分配粒度，这和GPU的物理架构有关。GPU执行模型是SIMT，核心尽量共享指令，但不同线程的同步、数据交换成本极高。把若干线程打包成Block整体分配，相当于把“谁来共享这间办公室”的问题提前定好，硬件不用每时每刻去调度单个线程。类比一下就是公司派项目：与其让每个员工在全大楼随机找座位，不如把一个项目组安排在同一层同一间会议室，协作成本大幅降低。

理解这一点是理解后面所有资源限制的基础：既然Block是整体，那么SM接收一个Block时，必须一次性给这个Block预留所有资源。

2. Block从Grid到SM的分配链路：分发器、Block槽位与并发调度

2.1 硬件层级里的工作队列与分发器

GPU内部并不是一个SM一个SM独立抢任务，有一个相对统一的分发机制把Grid里的Block派发给各个SM。

整体结构可以简化理解为：GPU分成若干GPC，每个GPC内部又有若干SM；启动kernel时，硬件维护一个全局的Block工作队列，各SM有空位就从这个队列里取下一个Block。

具体到SM侧，每个SM都有一定数量的Block槽位（slot）。以现代计算卡的常见参数为例，一个SM最多同时驻留16到32个Block。Block槽位和线程/Warp配额、寄存器文件、共享内存共同决定SM能同时“接多少活”。分发器做的事情，就是不断检查SM当前的资源余量，一旦发现某个SM还有空位，就把队列里的下一个Block塞进去。

这个过程完全由硬件完成，CUDA运行时会提供一个初始化的Block分布，但后续派发不经过CPU。

2.2 Block完成一个，立刻补一个：动态派发模型

Grid里的Block数量经常远大于SM能同时驻留的Block总量，于是Block会分批次执行，一批在SM上跑，其余排队等。

很多人误以为必须等一整批Block全部结束后，下一批才启动。实际上只要SM上有一个Block完成、退出，空出来的Block槽位就会被下一个等待的Block立刻占据，这就是动态派发。硬件不需要等一个“wave”完全收尾。

动态派发带来的影响很实际：Block执行时间越不均匀，后续Block越能“插空”进入，整体流水越平滑。反过来，如果每个Block执行时间几乎一模一样，调度就会呈现出明显的成批推进节奏，也就有了后文要说的Wave和尾效应。

2.3 并发块调度：从“等一个完成”到“并行嵌套”

早期CUDA架构里，一个SM哪怕还有空余资源，只要当前驻留的Block还没全部结束，新的Block通常不会启动，必须等某个Block完全做完。这在Block执行时间普遍较短的kernel上问题不大，但一旦有Block长时间等待，SM上的其他Block槽位就可能白白空着。

后来架构引入了并发块调度能力：同一个SM可以同时驻留多个尚未完成的Block，并且这些Block可以处于生命周期的不同阶段。比如A Block刚启动，B Block已经执行到一半，C Block正在等待内存访问，三个Block在SM内交错推进。

这个能力大幅改善了一个场景：Block内部做了同步等待（比如自旋等待某个全局标志），如果硬件一次只能跑一个Block，等待期间SM算力闲置；有了并发块调度，其他Block可以顶上，SM的算力利用率高很多。

3. 资源门槛才是真正的“招聘条件”：寄存器、共享内存与Block上限

3.1 Block要过三关才能被分到SM

硬件分发Block时，SM不是来者不拒，而是先算一笔账。一个Block想落到某个SM上，必须同时满足三件事：

• SM当前的Block槽位还没满，驻留Block数量低于上限；
• 这个Block需要的线程数，加上已有Block的线程数，不超过SM最大并发线程数（典型值是2048，即64个Warp）；
• 这个Block需要的寄存器总量和共享内存总量，加上已有Block的用量，不能超过SM的寄存器文件容量和共享内存容量。

只有三个条件同时满足，分发器才会把Block派过去。由于是按Block整组分发，所以哪怕SM只剩下一点点碎片资源，只要不够一个新的完整Block使用，这些资源就只能浪费掉。

3.2 算一笔账：寄存器数量多8个，占用率可能从满配掉到75%

寄存器是影响Block驻留数量最隐蔽的因素。很多新手不知道，kernel每线程用多少寄存器，是由编译器根据局部变量和临时计算自动决定的，但这直接决定一个SM能塞下几个Block。

假设某卡SM寄存器文件是65536个32位寄存器，SM最大2048线程。一个Block设为256线程，每线程用32个寄存器，那么每个Block要占256×32=8192个寄存器，65536÷8192=8，SM正好能驻留8个Block，对应线程数2048，占用率拉满。

如果因为多加了一个数组，每线程寄存器从32涨到40，每个Block就要占256×40=10240个寄存器，65536÷10240=6，SM只能驻留6个Block，对应线程数1536，占用率直接掉到75%。

所以你看，多用的那8个寄存器，不只意味着多占一点寄存器文件，而是让Block驻留数从8掉到6，少了整整两个Block的并行度。这就是我经常跟人说“寄存器是一级优化指标”的原因。

3.3 共享内存一上去，Block数立刻缩水

共享内存同理。SM上的共享内存容量是有限的，典型计算卡用户可用的共享内存一般在一百多KB左右。

假设一个Block申请32KB共享内存，SM可分配共享内存约100KB，100÷32=3，SM最多驻留3个Block。哪怕寄存器和线程配额都宽松，共享内存这一条就把它卡住了。

如果Block只申请10KB共享内存，100÷10=10，但Block上限还要看线程配额和Block槽位，综合下来可能驻留8个或更多。所以共享内存设计要非常克制，很多kernel做局部数据复用一上来就是大数组，结果Block数被压得很低，性能反而不如每次重复读全局内存。

下面的表总结了我常用的判断逻辑：

最终能驻留的Block数是这四个条件同时满足时的最小值。用代码算的话，可以调用cudaOccupancyMaxActiveBlocksPerMultiprocessor。

4. Block进入SM后会被切成Warp，决定执行节奏的是Warp调度器

4.1 Block是一等公民，Warp是真正的执行单元

Block被分配到SM之后，并没有直接在SM上整体“跑起来”。SM会先把Block里的线程按每32个一组拆成Warp，Warp才是硬件实际调度的单位。

一个Warp里的32个线程执行同一条指令，只是各自操作不同的数据，这就是SIMT模型的核心。假设一个Block有256个线程，硬件会把它拆成8个Warp，这8个Warp共享同一个Block的资源，受同一个Block生命周期管理，但执行时是独立被Warp调度器选择、发射指令的。

有一点值得注意：如果Block大小不是32的整数倍，最后一个Warp里会有部分线程处于未激活状态，这些线程不干活，但整个Warp依然占用一个槽位。所以Block大小选64、128、256、512这些32的倍数，不是为了好看，是避免浪费调度资源。

4.2 Warp Scheduler的工作节奏：每周期挑一个Warp发射指令

现代计算卡的SM里通常有4个Warp调度器，每个调度器每周期可以从自己负责的一组Warp里挑一个有资格执行的Warp，发射一条指令。

Warp的“有资格”意味着该Warp当前没有阻塞，下一条指令的操作数已经就绪，所属的执行单元空闲。如果某个Warp在等全局内存数据，或者被barrier.sync卡住，它就被标记为不可发射，调度器会跳过它去选别的Warp。

这就是GPU隐藏延迟的核心机制。全局内存延迟动辄几百个周期，如果把整个SM只放一个Block几个Warp，内存等待期间执行单元就空了。减少这种空白的主要手段就是提高Warp并发度，让一个Warp等待时，其他Warp能立刻顶上来继续发射指令。所以占用率低通常意味着延迟隐藏能力弱。

4.3 为什么Block大小对齐Warp之后，还要警惕Block内部的“重型同步”

Block内部使用__syncthreads()同步时，等效于把Block内所有Warp绑在一起：只有全部Warp都到达同步点，它们才能继续往下走。

如果Block很大，例如1024线程，一个同步点意味着32个Warp必须互相等待，快的Warp要等慢的Warp。如果Block较小，例如128线程，同步代价就小很多，调度器也能更灵活地穿插不同Block的执行。

代价是Block太小会导致Block数量变多，每个Block的启动、调度、结束都会带来额外开销。所以“Block大小到底选128还是256还是512”不能拍脑袋，要在同步开销和调度开销之间找平衡。实测里256是很多kernel的常用起点，但具体kernel一定要自己测。

5. 调度波次与尾效应：为什么你的kernel总差一口气

5.1 用数字看Wave与Tail：看似并行，其实在排队

设一台设备有16个SM，每个SM最多驻留8个Block，则同时最多有128个Block在跑。如果你启动256个Block，刚好两个满波次，每个SM从头到尾都有活干。

但如果你启动140个Block，第一批128个占满所有Slot，等有Block结束腾出位置后，剩下的12个Block才能进入，而这12个Block分布到少量SM上，其余SM可能已经空了。

按理想状态算，140个Block / 128个并发容量，消耗约1.09个波次的时间；但实际因为最后12个Block只有不到一个满波的工作量，最后一批跑起来时SM利用不均衡，整体耗时会明显大于1.09个波次，严重时接近2个波次。多出来的这部分就是尾效应。

尾效应在Block执行时间接近时最明显。如果Block执行时间本身差异大，动态派发反而能抹平一部分尾部浪费，这也是为什么负载不均衡的kernel有时“看起来”调度并不差。

5.2 动态补位救不了跨Block依赖

既然硬件支持“完成一个补一个”，尾效应是不是就没那么严重了？不一定。

动态补位解决的是“有独立新Block可以启动”的情况。如果一个Block执行到一半在等待另一个Block的结果，比如用全局内存标志位自旋等待，那么这个Block虽然占着资源，却干不了活；如果其他Block都还没有结束，SM也可能无新Block可补。

我踩过的一个典型坑：Kernel里让Block 0算一份结果，写到全局标志位，其余Block自旋等这个标志位。结果网格规模一上去，性能快速恶化。原因是等待的Block把SM的Warp槽位全部占了，后到的Block挤不进来，而Block 0又因为资源被占迟迟跑不完，形成互相等待。这种跨Block同步问题，线程块的动态调度解决不了，得从算法层面重新设计。

5.3 两个缓解思路：Wave对齐与持久化线程

缓解尾效应最直接的办法，是让Grid大小尽量对齐“波次”。先用cudaOccupancyMaxActiveBlocksPerMultiprocessor算出一个SM最多驻留多少Block，再乘以SM数量，得到一次能同时跑满的Block数，然后让Grid等于这个数的整数倍。

另一个思路是持久化线程（Persistent Kernel）：启动固定数量的Block，数量大约等于SM并发容量，让每个Block在kernel内部用循环处理多份数据。这样Block在kernel生命周期内一直占着SM，不会被反复销毁、重建，也天然消除了尾波次。

代价是持久化线程需要自己维护工作队列、处理负载均衡，不能再用简单的blockIdx.x直接对应数据切片。它适合每个Block计算量小、启动次数多的场景，如果计算量本身很大，普通Grid/Block模型就够了。

6. 想验证调度行为？我常用的两个工具和三个自测实验

6.1 Nsight Compute里最值得看的三组计数

讲理论容易，眼见为实。排查Block调度问题时，我用Nsight Compute比较多，主要盯三组数据：

• Achieved Occupancy：kernel实际执行期间平均每个SM驻留的活跃Warp数占理论最大值的比例，能直接反映SM“有没有吃饱”；
• 每个SM的Active Blocks数：可以确认到底驻留了几个Block，结合Block语义判断理论值和实际值对不对得上；
• SM Busy / Issue Activity：调度器每个周期实际发射的指令比例，如果Occupancy不低但Issue很低，说明Warp大多处于等待状态而不是可发射状态。

还有一个很实用的API，可以直接在本机设备上查理论最大驻留Block数：

int dev = 0; cudaDeviceProp prop; cudaGetDeviceProperties(&prop, dev); int maxBlocksPerSM = 0; cudaOccupancyMaxActiveBlocksPerMultiprocessor( &maxBlocksPerSM, myKernel, 256, 0); printf("maxBlocksPerSM for blockSize=256: %d\n", maxBlocksPerSM);

cudaOccupancyMaxActiveBlocksPerMultiprocessor会综合考虑注册资源、共享内存和线程配额，返回当前kernel在当前设备上单个SM最多能驻留几个Block。我写性能测试脚本时，第一步永远是跑这个函数，先把理论值拿稳。

6.2 实验A：小Grid到底占用了几个SM

拿一个没有任何复杂逻辑的kernel，Grid设为8个Block，Block 256线程，跑在一个有16个SM的设备上。用Nsight Compute看每个SM的Active Blocks分布。

常见的结果是这8个Block被分散到8个不同的SM上，每个SM一个Block，剩下8个SM全程空闲。这验证了两个结论：硬件分发Block倾向于先“铺开”到不同SM，而不是把所有Block塞进同一个SM；同时Grid小于SM并发容量时，空闲SM不会帮你“匀”任何计算。

所以当你启动的Grid很小，就别指望GPU整个都在忙。想测试Tail效应的话，把Grid设成waveCount * maxConcurrentBlocks + 1，对比最后加不加那1个Block的耗时，往往能看出明显的耗时台阶。

6.3 实验B：同一kernel在不同BlockSize下的占用率曲线

选一个中规中矩的kernel，把BlockSize依次设成64、128、256、512、1024，分别跑一遍，记录Achieved Occupancy和总耗时。

很多人的第一反应是Block越大越好，因为每个Block处理的数据多、调度开销小。实测结果往往在256或512附近最好，原因就是BlockSize变大之后，Block内同步代价上升，寄存器压力也可能让驻留Block数下降，最终降低真实并行度。

这个实验我建议用脚本批量跑，把每次的寄存器数、共享内存量、驻留Block数、耗时打在同一个表里，你会很直观地看到“资源占用-驻留Block数-性能”之间的联动。

6.4 实验C：用共享内存把SM“挤成”单Block驻留

做一个故意申请大共享内存的kernel，比如在Block里定义一个64KB的__shared__数组（写成16×1024个float）。在共享内存总量约100KB的设备上，一个SM最多驻留1个Block，原本能跑满8个Block的kernel会因为共享内存这一条直接降到1个。

这个实验很有说服力：kernel里的并行度上限不是线程数决定的，而是资源账单决定的。每Block共享内存越大，Block驻留数就越少，一旦Block驻留数不足，Warp调度器手里没有足够多的Warp用来切换，内存延迟就藏不住了，性能立刻反应出来。

我记得第一次跑完这个实验，之后再写kernel都会习惯性先看一眼每Block共享内存量，能省就省。很多“莫名变慢”的问题，追到底都是这里出了问题。

一些我现在养成的习惯

现在每次调CUDA性能，我都会先做三件固定动作：用cudaOccupancyMaxActiveBlocksPerMultiprocessor拿到理论驻留Block数；查一眼每线程寄存器数和每Block共享内存量；再估算当前Grid是不是接近整波次。这套组合拳帮我排掉了大量“看起来算力够、但就是跑不满”的怪问题。Thread Block的调度机制没有多神秘，但它是GPU上一切并行度的地基，把地基上的资源账算清楚，很多性能问题其实在写kernel阶段就可以避免。