GPU SIMT执行模型：SM、Warp与算力碎片
GPU 算力利用率上不去，多数时候不是软件写得烂，而是撞上了硬件的调度粒度。本文从 SM 讲到 SIMT 与 Warp，再到 Grid/Block/Thread 的软硬件映射，说清算力碎片的来源、哪种空闲能复用、哪种不能——这是 VarLen：变长序列推理如何消灭Padding 这类布局优化的能力边界所在。
GPU 的并行骨架：SM
GPU 和 CPU 的本质区别在于优化目标。CPU 为串行延迟优化：几个大核心，配复杂的分支预测和多级缓存，目标是把单条指令流跑快。GPU 为吞吐设计：用大量简单的计算单元并行处理海量数据，靠并行度隐藏延迟。构成 GPU 并行算力的基本硬件单元叫 SM（Streaming Multiprocessor，流式多处理器），一块 GPU 有几十到上百个。
可以把整块 GPU 看成一个大型工厂，每个 SM 是一个独立车间。每个 SM 有自己的计算单元——CUDA Core 和 Tensor Core，有自己的寄存器堆，还有一块速度极快但容量很小的共享内存。车间之间完全独立：各自调度各自的任务，不能直接访问对方的资源，要通信只能走工厂的中央仓库，也就是全局显存，延迟高得多。
这种独立性是后面所有分析的关键边界——判断空闲算力能不能复用，切分线就画在 SM 之间：不同 SM 上的任务可以并发互不干扰，而 SM 内部的算力怎么分配，是硬件自己的事。
SIMT 与 Warp：硬件调度的最小粒度
车间里的任务怎么执行？这就要说到 CUDA 最核心的设计：SIMT（Single Instruction Multiple Threads）。它经常和 CPU 的 SIMD 搞混，但完全不是一回事：SIMD 是一条指令操作一个向量寄存器，本质仍是单条执行流，没有独立线程的概念；SIMT 的核心是硬件把线程打包成固定大小的组——Warp，NVIDIA 架构下一个 Warp 固定 32 个线程。同一个 Warp 里的 32 个线程共享同一个程序计数器，同一时刻执行完全相同的指令。
Warp 是 GPU 硬件调度的最小单位，没有之一。调度器从不单独调度某个线程，只调度完整的 Warp。Warp 内的线程虽然执行同一条指令，但各自持有独立的寄存器，可以加载不同数据、处理不同输入——指令流共享，数据独立，这就是「单指令多线程」的全部含义。
这套模型处理整齐数据时效率极高：一张图片每个线程算一个像素，步调一致，零浪费。一旦遇到分支就出问题：同一个 Warp 里一部分线程走 if、另一部分走 else，硬件的解法是两个分支都跑——跑 if 时走 else 的线程被掩码屏蔽，不写结果、原地空转；跑 else 时反过来。整个 Warp 的执行时间等于两个分支之和。这就是 Warp 分支发散（branch divergence），GPU 算力浪费最根本的来源。
为什么掩码线程不能偷跑别的任务
Warp 内 32 线程共享同一个程序计数器，同进退是硬件约束而非软件选择。被掩码的线程和活跃线程走的是同一条指令流，在「跑 else 分支」的那个时钟里，它没法去执行另一段毫不相干的代码。
短序列推理填不满 Warp 就是这个机制的直接后果：VarLen 铺平后 batch 里只有 17 个有效 token，硬件照样分配一个完整 Warp，剩下 15 个线程全程掩码，而这个 Warp 占据的调度槽位和硬件资源一点不少。硬件层面算力利用率的口径，见 显卡的利用率是如何计算的？。
从 Kernel 到硬件：Grid、Block、Warp 的映射
往上看软件层。跑在 GPU 上的函数叫 Kernel，在 CPU 端调用 Kernel 本质是给 GPU 提交一个计算任务。任务在软件上组织成三层：Grid、Block、Thread——整个任务是一个 Grid，Grid 划分成多个 Block，每个 Block 包含多个 Thread。
三层结构最终映射到硬件：一个 Block 被完整分配到一个 SM 上执行，绝不跨 SM 拆分；一个 SM 可以同时驻留多个 Block，靠快速切换隐藏显存访问延迟；Block 内部的线程被硬件自动按 32 个一组切成 Warp，交给 SM 的调度器。Grid 和 Block 是软件给硬件的任务打包方式，硬件真正认识和调度的只有 Warp。
这里有个高频坑：写算子时 Block 大小随手设成 17、25 这种非 32 倍数的值，硬件会自动补齐成完整 Warp，多出的线程从一开始就被掩码。所以 CUDA 算子的基本准则之一是 Block 线程数必须取 32 的倍数——本质就是向 Warp 粒度对齐，砍掉天然空转。
Block 维度设计
blockDim.x优先取 32 的倍数（128、256 是常见起点）；线程数不足一个 Warp 的小任务，考虑在 Block 层面合并，而不是让硬件替你补零。
把这些串起来，一个 Kernel 从提交到执行完的完整路径：
- CPU 端调用 Kernel，任务进入 GPU 的任务队列
- GPU 把 Grid 拆成 Block，分配到有空闲资源的 SM
- Block 进 SM 后内部线程切成 Warp，调度器维护所有 Warp 的状态——就绪、等显存、等同步
- 每个时钟周期，调度器挑就绪的 Warp 发射指令；遇到分支就分开跑，产生空转
- Block 的全部 Warp 执行完，释放寄存器和共享内存，SM 加载下一个 Block
- 所有 Block 跑完，Kernel 结束
矩阵乘这类规整计算怎么在这个模型上跑满，见 GEMM。
两种空闲：Warp 碎片是死的，SM 闲置是活的
「空闲算力能不能复用」要按空闲发生在哪一层分开回答，边界就画在 Warp 和 SM 之间。
第一种是 Warp 内部碎片：一个 Warp 里只有部分线程处理有效 token，其余被掩码。这部分线程绝无可能拿去干别的——SIMT 模型里 Warp 是最小调度单元，同进退是铁律，不能把一个 Warp 拆开、把闲置线程抽出去跑其他任务。这部分算力产生即报废，零复用可能。短序列推理的算力上限永远到不了 100%，根子就在这里。
第二种是 完整的 SM 空闲：任务特别小时整个 Kernel 只用到少数几个 SM，剩下的 SM 完全没有任务。这部分是活算力，工业界已有成熟方案——多 CUDA Stream 并发，一个 Stream 跑短序列 Prefill 的同时用另一个 Stream 提交 Decode 或其他小算子，让空闲 SM 转起来；推理框架里常用的重叠调度（overlapping schedule）把 Prefill 和 Decode 的执行时间重叠，用的也是短任务跑完后 SM 的空闲窗口。
别把两种浪费混为一谈
零散的 Warp 内碎片是死的，只有完整的 SM 级闲置才是活的。「为什么不能把空着的线程利用起来」——不是工程师不想做，是硬件根本不支持。
工程缓解：在硬件边界内挤性能
碎片消除不了，但可以稀释。最直接的思路是加大并行规模：把更多请求塞进同一个 batch，让每个 Warp 里的有效线程占比尽量高。
| 并行规模 | 占用 Warp 数 | 有效线程 | Warp 利用率 |
|---|---|---|---|
| 17 token | 1 | 17 / 32 | 53% |
| 60 token | 2 | 60 / 64 | 94% |
代价是排队延迟升高——吞吐和延迟的经典权衡，没有免费午餐。
再进一步是分群调度：把长度相近的请求放进同一个 Pack，避免长短混合。长序列和一堆短序列混在一起，长序列拖慢整个节奏、短序列又填不满 Warp，两头浪费；分群后长序列组算力利用率高，短序列组尽量打包，整体收益更好。至于 Prefill/Decode 重叠执行、多任务混部，本质上都是在 SM 层面做时间复用，把空闲的硬件时间片填满。
推理栈里的对应物
vLLM、SGLang 的 continuous batching 负责「加大并行规模」这一层；FlashAttention 的 varlen 布局负责消灭 padding 的逻辑浪费；PD 重叠、多 Stream 负责填充 SM 空闲窗口。三者叠加才逼近硬件允许的利用率上限，见 VarLen：变长序列推理如何消灭Padding。
说到底，性能优化到最后拼的是对硬件底层的理解。很多看起来是软件的问题，根子都在硬件模型里——分清什么能改、什么不能改，哪里能挤牙膏、哪里是物理天花板，才能在边界之内把性能做满。