GPU计算架构与Warp调度：从Block投递到指令发射的底层机制

这一期我们聊计算架构与调度。排序放在第四期，是因为很多读者已经写过基础核函数，开始撞上性能瓶颈了：同样的逻辑，别人跑 0.5 毫秒，你跑 1.2 毫秒；多写了一个 if，整段程序像被踩了刹车。这些问题不解决，写再多核函数都只是在“碰运气”。我自己当年读 P M P P 相关章节时，最难建立的就是对“调度”的画面感——几千个线程到底在哪跑、怎么被喂给计算单元、为什么有的代码就是喂不饱硬件。这一期我会把这些底层的运行机制拆开讲清楚，适合已经能写出基础 CUDA 核函数、但还想把性能真正提上来的读者，也适合刚接触这类并行编程概念、想建立一个准确心智模型的初学者。

1. 从“一个 if 让性能掉了一半”说起：先建好计算架构的心智模型

1.1 CPU 在“等”，GPU 在“换”

传统 CPU 的设计哲学是降低单条指令的延迟。为了把一个循环跑得足够快，它把大量晶体管花在乱序执行、分支预测、多层缓存上，目的只有一个：让流水线尽量不要空转。可 GPU 不一样，它最高效的场景是同时跑成千上万个任务，单个任务慢一点无所谓，只要总量大、吞吐高就行。

这里有个很核心的概念叫延迟隐藏。你可以想象一个食堂窗口：CPU 的做法是只有一个厨师，但厨师动作极快，还提前把食材备好，避免你等太久。GPU 的做法是开一排窗口，每个窗口炒菜速度也就那样，但排队的人足够多，窗口永远不闲。显存访问要等几百个周期？没关系，这时候调度器切到另一个已经就绪的线程束继续发指令，计算单元照样满负荷。

所以同样一段代码，在 CPU 上优化方向是减少等待、提高单线程效率；在 GPU 上优化方向是增加并行度、让可切换的“等待者”足够多。这也是“大规模并发处理器”这个名字的由来——它赌的就是并发规模。

1.2 控制逻辑做减法，计算单元做加法

如果你看过 GPU 的核心面积分布，会发现它和 CPU 完全是两个路子。GPU 把绝大多数晶体管投给了算术单元、寄存器堆和共享内存，而把分支预测、乱序执行这类复杂控制逻辑砍到很薄。原因很简单：复杂的控制逻辑是为了让单个线程跑得更快，但 GPU 本来就打算用几千个线程来摊薄延迟，没必要在单个线程上花那么多资源。

但资源结构决定行为方式。正因为控制逻辑被简化，同一个调度器要“照顾”非常多线程的指令发射，它就必须采用一种相对整齐、规则的方式去组织线程。这就引出了 GPU 执行模型里最不能跳过的一层：Warp（线程束）。程序员惯常写代码时只关心 thread、block、grid，可在硬件眼里，一切最终都要落到 Warp 上。

2. 拆开 GPU 看内部结构：从网格到 SM 再到 CUDA Core

2.1 网格、块、线程对应到硬件里的哪几层

我们在 CUDA 里写的 grid 和 block 是软件视角的组织方式。硬件视角大致是这样：一颗 GPU 里有许多个 SM（流式多处理器），SM 是真正执行计算和调度的基本单位；一个 SM 内部又有若干个计算核心、寄存器堆、共享内存、缓存和调度器。

启动一个 kernel 时，grid 里的 block 会被分发到各个 SM 上，SM 负责执行这些 block 里的线程。一个 block 不会拆散分给多个 SM，它要么整块落在一个 SM 上，要么还在队列里等资源。这里你可以把 SM 想象成一间间“诊室”，block 是“病人”，病人不能被劈成两半让两个医生同时接诊。

不同厂商、不同代际的 GPU，SM 这个名字大体不变，但内部的执行核心数量、调度器数量、寄存器数量、共享内存容量都不同。所以一个很重要的习惯是：拿到一台设备，先用 API 查清楚它的属性，再决定 launch 配置。千万别把某个架构上“最优”的 block 大小直接搬到另一台机器上。

2.2 一条指令在 SM 内部的完整旅行

我们放大看一次执行过程。SM 拿到一个 block 后，会把它切成一个个 Warp，每个 Warp 固定包含 32 个线程。调度器每次从其中一个 Warp 里取一条指令，发射给一组执行单元。所谓“执行单元”，就是你在各种资料里看到的 CUDA Core，本质上是一个个能完成整数、单精度浮点等基本运算的 ALU 单元。

一条指令的路径大约是：从指令缓存取指，经过调度器判断哪个 Warp 的指令已就绪，把这条指令广播给某个 Warp 的 32 个线程对应的执行单元；每个线程从自己的寄存器里取操作数，执行运算，结果写回自己的寄存器。也就是说，32 个线程共享同一个程序计数器，但在寄存器层面，每个线程都有自己的独立副本。

这也是为什么 GPU 里寄存器比 CPU 多得多。CPU 一颗核心的寄存器文件很小，因为只需要服务一两个线程；而 GPU 一个 SM 要服务几十上百个 Warp，每个线程都要有自己的寄存器上下文，寄存器文件自然得做成一个巨大的阵列。你可以把它理解成一个大型存储柜，每个线程的局部变量住在自己的抽屉里。

2.3 存储层次和隐藏延迟的关系

一个 Warp 里的线程执行一次访存指令时，数据可能来自寄存器、共享内存、高速缓存或者显存。访问延迟从几个周期到几百个周期不等。下表是我常用的一组数量级参考值，具体数值会随架构变化，但量级关系基本稳定：

为什么多 Warp 轮转很关键？因为单个 Warp 遇到高延迟访存时，只能在那里等。好在 SM 上不止一个 Warp，调度器会趁这个空档去执行其他 Warp 的就绪指令。也就是说，你不光要写并行代码，还要给硬件留出足够的 Warp 数量，它才有资格“换着做”。

3. Warp 不是抽象概念，而是物理事实：32 条线程为什么被绑在一起

3.1 threadIdx.x 为什么是最“忙”的那个维度

一个 block 被分配给 SM 后，硬件会按固定的顺序把线程切成 Warp。切分时，threadIdx.x变化最快。假设一个 block 是二维的dim3(16, 8)，线程按照线性顺序编号，编号规则就是常见的公式：

线性编号 = x + y * blockDim.x

连续 32 个线性编号会被分到同一个 Warp。也就是说，一个(16, 8)的 block 前两行正好组成 Warp 0，第三四行组成 Warp 1，依此类推。如果你用一维 block 且大小为 64，那么前 32 个线程是 Warp 0，后 32 个是 Warp 1。这个映射规则非常重要，因为它决定了同一 Warp 内的线程访问共享内存时，会不会撞出 bank conflict。

另外要留意：一个 block 的线程数不要求是 32 的整数倍。比如 block 只申请 60 个线程，硬件还是会把它占成两个 Warp 的槽位，只是第二个 Warp 里有 4 条 lane 是空闲的。空闲的 lane 不干活，但调度和发射开销仍然在。所以实际工程中，我基本都会把 block 大小设计成 32 的整数倍，省得最后一个 Warp 半空着浪费发射带宽。

3.2 为什么偏偏是 32

Warpsize 定为 32，并不是某个厂商拍脑袋定的，而是工程上权衡的结果。32 这个数字有两个好处：一来它是二进制体系里很规整的数，二来它对应一次指令发射的最合适宽度。

如果 Warp 太小，比如 8，那么每个 Warp 并行度太低，调度器管理大量小 Warp 的开销会明显变大，指令发射带宽也会被控制成本稀释。如果 Warp 太大，比如 128，那么一个 Warp 内出现分支分歧的概率指数上升，一旦某个分支条件在 128 个线程里不统一，硬件就得把多条路径串行执行，浪费成倍放大。32 是在“发射效率”和“分歧代价”之间站得比较稳的一个点。

还有一个容易忽略的细节：同一个指令发射给 32 个线程时，这些线程是在不同执行单元上并行算的，而不是在一条很宽的 SIMD 指令里算。所以 GPU 上的实际带宽比字面上的“每周期一条指令”要宽得多，因为这一个周期里有 32 条标量运算在同时发生。

3.3 SIMT 和 SIMD：长得像，骨子里不同

很多新手会把 SIMT 和 SIMD 混为一谈。SIMD 是一条指令操作一个很宽的向量寄存器，所有元素必须步调一致地执行同一次运算，比如向量加法。SIMT 是单指令多线程，32 个线程执行同一条指令，但每个线程有自己的寄存器、自己的地址、甚至自己的控制流判断能力。

差别在哪里？SIMD 是你把四个数装进一个 128 位寄存器里，一次性算完；SIMT 则是 32 个标量线程各自从自己的寄存器里取数，虽然指令相同，但数据完全不同。这带来的编程体验就是：写 CUDA 内核的时候，你几乎感觉自己在写单线程的标量代码，完全不用把一个循环改写成向量化形式。这种“伪装”很舒服，可代价藏在分支里：如果同一 Warp 内线程走向不同分支，硬件只能把各条路径都执行一遍，这就是下一部分要重点聊的分支发散。

4. 两级调度接龙：从 Block 投递到 Warp 指令发射

4.1 第一跳：全局调度器把 Block 分发到 SM

kernel 启动时，grid 中所有 block 会进入一个由硬件管理的待分配队列。每个 SM 能同时驻留多少个 block，取决于它的资源上限：线程数上限、寄存器数量上限、共享内存上限。全局调度器会逐个把 block 投递到有资源余量的 SM，资源不够就先等着。

这个过程有几个特征值得记住。第一，block 一旦被投递到某个 SM，就不能再迁移，直到结束。如果某个 block 执行得很慢，同一个 SM 上的其他 block 也许要等它让出资源。第二，不同的 SM 可能拿到不同数量的 block，这不取决于程序员，而取决于资源占用情况。第三，调度器是公平的，它按顺序分配，和 block 的复杂度无关。

理解了这一跳，你就能明白为什么“某个 block 里的线程能不能和另一个 block 里的线程同步”是个伪命题。block 分布在不同的 SM、不同的调度器下面，它们之间连全局同步的原语都不存在，只能靠 kernel 边界来同步。

4.2 第二跳：Warp Scheduler 每个周期挑一个“就绪者”

Block 进入 SM 后，被拆成 Warp，接下来就轮到了真正的指令调度。我们可以这样理解 Warp Scheduler 的工作：每个周期，它从自己负责的那组 Warp 里，挑一个“下一条指令已经就绪”的 Warp，把这周期的一条指令发射出去。如果这个 Warp 正卡在等待显存返回、等待共享内存数据、或者等待同步屏障，它就还没就绪，调度器不会白白等它。

现代 SM 一般会有多个 Warp Scheduler，每个负责一部分 Warp。以常见情况为例，如果一个 SM 能驻留 64 个 Warp，且内部有 4 个 Scheduler，那每个 Scheduler 大约管理 16 个 Warp，大家轮流发射，互不干扰。

这个机制解答了一个常见困惑：“为什么我的 Warp 明明很多，利用率却上不去？”答案多半是你的 Warp 大量处于未就绪状态，调度器想发射也没得发射。好比食堂窗口排了 20 个人，但其中 15 个人还没想好要点什么菜，窗口照样空闲。

这里我还是放一个简单的时刻表，演示一个周期里调度器的选择逻辑：

实际硬件里调度器会利用优先级、Warp 年龄、指令类别来做更细致的裁决，但宏观逻辑就是“就绪才发，谁合适发谁”。我们要写的代码，本质上就是让大多数 Warp 在大多数时刻处于就绪状态。

4.3 零开销切换的秘密：寄存器槽位是一早锁好的

我们常说 GPU 做 Warp 切换是零开销，这听起来很玄，其实原理很简单。CPU 切换线程时要把寄存器、程序计数器、栈指针等一大堆上下文保存到内存里，再恢复另一套，开销不小。而 GPU 的 Warp 切换，调度器只是换一个 Warp 发射指令罢了，被切换走的 Warp 的寄存器还在物理寄存器堆里原位放着，根本不需要搬来搬去。

代价是什么？寄存器是静态分配的。一个线程从 kernel 开始到结束，它能用的寄存器数量是固定下来的，不能像 CPU 那样按函数调用动态伸缩。如果你想提高并发度，就得控制每个线程使用的寄存器数量；寄存器用太多，SM 能驻留的 Warp 就会变少，延迟隐藏能力随之下降。这是一种非常典型的空间换时间的权衡，理解这一点，是下一章所有性能分析的基础。

5. 调度器用这三堵墙“卡”性能：寄存器上限、尾效应与分支发散

5.1 三个资源墙：线程槽、寄存器堆、共享内存

一个 SM 能驻留多少 Warp，首先看三个硬指标：线程数上限、寄存器堆大小、共享内存大小。三者的任何一个到顶，SM 就不再接收新的 block 或 Warp。为了直观，我假设一个典型配置：SM 的理论最大线程数是 2048，寄存器堆有 65536 个 32 位寄存器，Warp 大小 32。

如果你每个线程只用 32 个寄存器，那么 2048 个线程正好把寄存器堆全部占满，SM 满线程运行，利用率 100%。如果编译器给每个线程分配了 64 个寄存器，那最多只能驻留 2048 个线程的一半——1024 个线程。如果编译器顽固地给你每个线程 128 个寄存器，那就只剩 512 个线程了，整个设备空一半。

注意，表格里的数字只是演示用的。其实，编译器默认会在寄存器使用和潜在溢出之间找平衡，但你写复杂内核时，寄存器数量很容易失控。解决方案不外乎两种：用编译选项限制最大寄存器数，或者用__launch_bounds__告诉编译器“这个 kernel 最多同时跑多少线程”，它就会按这个目标去压缩寄存器分配。代价是如果寄存器真不够，变量会被挤到局部内存里，访问延迟变大。所以这是一个需要实测权衡的调优点。

共享内存同理。如果你的 block 申请了很多动态共享内存，SM 上能同时放的 block 数就会减少。合理使用共享内存能带来巨大加速，但过量使用会反过来限制并发度，导致其他 Warp 没法驻留，访存延迟没人掩盖，性能照样掉。

5.2 尾效应：最后几个 Block 为什么拖后腿

调度是动态的，但硬件“同时能跑的数量”有上限。假设你的设备有 16 个 SM，每个 SM 最多能驻留 2 个 block，那么在理想情况下，它同时最多运行 32 个 block。如果你的 grid 正好是 33 个 block，前 32 个 block 把位置占满后，最后一个 block 必须等其中某个块结束才能被调度。

问题在于：最后这个阶段，只剩 1 个 block 在跑，其他 SM 的硬件资源全部闲置。如果每个 block 执行时间差不多，总耗时约等于两批 block 的执行时间；你要是 grid 只有 33 个 block，而设备能并行 32 个 block，那两批和一批只差一点，但如果每个 block 耗时不短，这种“尾巴”造成的空转就很可观。

这不是程序 bug，而是资源粒度造成的负载不均。缓解方式通常是：让 grid 里的 block 数量远大于“SM 可驻留 block 数”的整数倍，使大部分执行阶段都能吃满硬件。如果你只需要 100 个 block 的工作量，设备却只能同时跑 128 个，那也不要去硬凑成 256，因为多出来的 block 只会增加启动和调度开销，尾效应反而被放大。

5.3 分支发散：Warp 内的“众口难调”真会伤筋动骨

SIMT 模型给了每个线程独立的控制流能力，同样也埋下了效率陷阱：一个 Warp 内，如果 32 个线程的某个条件判断结果不一致，硬件必须把 Warp 拆成多个执行路径，先执行条件为真的线程，再执行条件为假的线程，执行时间大约是各路径的和，而不是路径的最大值。

例如一个 if-else，如果一半线程走向分支 A，一半走向分支 B，那么这个 Warp 的实际执行成本就是 A+B，理论上几乎翻倍。如果分支是三层 if-else，情况更糟。反过来，如果整个 Warp 的 32 个线程都判断为真，则没有分歧，硬件只执行 A 路径，非常快。

怎么优化？常见策略有三个：一是让分支尽量只在 block 粒度或 warp 粒度上出现，例如按warpId = threadIdx.x / 32做分流，保证整束一致；二是把简单条件表达式改写成算术方式，让编译器生成 predicated 指令而不是真正分支跳转；三是调整数据排布，把需要走同一分支的数据聚合到同一个 Warp 里，这在某些算法里能带来成倍提升。记住判断标准很简单：看这段分支在 Warp 内是否“步调一致”，一致就是赚，不一致就要想法子消除。

6. 从计算架构倒推出几条可落地的优化习惯

6.1 写代码前，先把设备底牌查清楚

很多 launch 配置问题都源于“不看设备参数瞎猜”。我现在的习惯是，写任何 kernel 之前，先跑一小段代码把关键参数打出来：

cudaDeviceProp prop; cudaGetDeviceProperties(&prop, 0); printf("SM count: %d\n", prop.multiProcessorCount); printf("max threads per SM: %d\n", prop.maxThreadsPerMultiProcessor); printf("regs per SM: %d\n", prop.regsPerMultiprocessor); printf("shared mem per SM: %zu bytes\n", prop.sharedMemPerMultiprocessor);

这几个字段基本决定了你可以怎么配 block 和 grid。比如知道每个 SM 最多 2048 线程后，block 选 256 意味着每个 SM 最多驻留 8 个 block；如果 SM 数恰好是 16，想填满整卡，grid 至少要有 128 个 block，还要考虑数据规模是否匹配。很多教程只告诉你 block 取 256 比较“安全”，但你真正迁移到不同算力的机器上时，还是要用这两行代码验证。

6.2 用占用率计算器和 Profiler 验证调度假设

遇到性能问题时，不要只靠直觉猜“是不是算法不够好”。我会先用官方提供的占用率计算表格把每线程寄存器数、共享内存量输进去，看理论占用率是多少。如果占用率已经被寄存器压到 50% 以下，那先不要动算法，先把寄存器限制调整好再说。

之后再用 Profiler 这类工具看实际执行情况。重点关注两类信息：一类是“stall reason”，也就是 Warp 卡住的原因，常见有等待显存返回、等待共享内存、等待同步屏障、分支解析等。如果大量的 stall 是等待访存，说明并发 Warp 可能还不够多，或者访问路径跑到了显存而不是缓存。另一类是占用率和执行带宽的差距，如果占用率很高但吞吐还是上不去，那可能要检查是否存在 bank conflict、尾效应或者负载不均。

6.3 我常年使用的三个编码习惯

第一个习惯：block 大小固定为 32 的整数倍，且尽量让一个 block 内的 Warp 数也固定。这样每个 Warp 都满载，不会出现半空 Warp 浪费发射槽的情况。

第二个习惯：把大循环尾数处理收敛到单独一个 block 里，而不是让 grid 最后一个 block 承担不规则的工作量。比如一个数组长度是 100000，block 是 256，最自然的启动方式是 391 个 block，最后一个 block 只处理 192 个元素。问题不大，但如果你后续还有 block 间同步或归并逻辑，最好在数据划分阶段就设计成“大部分 block 工作均匀，尾数集中处理”。

第三个习惯：能用数据选择替代分支的地方，尽量不用 if。例如判断一个值是否超过阈值并缩放：

// 不推荐：warp 内可能严重发散 if (val > threshold) { out[i] = val * 2.0f; } else { out[i] = val * 0.5f; } // 推荐：编译器更容易生成选择指令 float scale = (val > threshold) ? 2.0f : 0.5f; out[i] = val * scale;

注意，这里的核心不是“三元表达式一定快”，而是给编译器更多机会把分支转换成无发散的预测执行。复杂分支就别指望这么容易绕过去了。

最后说一个我自己的调试习惯。现在遇到核函数性能异常，我第一反应不是反复改代码盲试，而是先开 Profiler 看两件事：理论占用率被哪种资源卡住了，实际的 Warp stall 主要集中在哪里。把这两个问题定位清楚，再回头改 launch 配置或者改写分支逻辑，效果往往立竿见影。计算架构和调度这层窗口，一旦打开，很多性能问题就不再是玄学，而是可以按图索骥的工程问题了。