先别急着骂显卡是“虚标王”。GPU-Util 100%、SM满载、Warp调度打满，结果 nvidia-smi 里算力只有15%，这个场景在深度学习训练、高性能计算里太常见了。我最早遇到这问题是在调一个 transformer 推理服务，GPU 占用率显示接近 100%，但 QPS 上不去，卡片的热功耗和实际吞吐完全对不上。查了半天才发现，任务管理器里的 100% 和真正“把 SM 里每一份算力都榨干”是两码事。
这篇文章我打算从第一性原理把这笔账算清楚，涉及Warp 调度、Tensor Core执行路径和Roofline 模型三个核心概念。适合正在调 GPU 性能的算法工程师、高性能计算开发者，也适合被“利用率虚高”困扰的运维同学。看完你能建立一套自己的诊断框架：拿到一个 kernel，先判断它是计算受限还是访存受限，再用工具验证，最后对症下药。
1. 现象拆解：GPU-Util 100% 为什么是“伪满血”
1.1 任务管理器里的 GPU-Util 到底在测什么
Windows 任务管理器、nvidia-smi 里的 GPU-Util，甚至很多监控面板上的“GPU 使用率”，本质上都是一个非常粗糙的采样指标。以 nvidia-smi 为例，它轮询的是 GPU 上有没有 kernel 在执行：只要一个 kernel 占用了某个引擎的时间片，这一瞬间就被标记为“busy”。它不区分这个 kernel 是在跑 FP32 矩阵乘法，还是在线程块之间做一次空转同步。
这意味着什么？你一个 kernel 在 GPU 上跑了 10 毫秒，其中 9 毫秒在等显存数据返回，这 10 毫秒里 nvidia-smi 看到的就是“GPU 有活干”，utilization 100%。但对于 SM 里的计算单元来说，这 9 毫秒里 FP32 单元、Tensor Core 全都在“干瞪眼”。所以 GPU-Util 100% 只代表“时间片被占用”，不代表“计算单元被填满”。
这有点像餐厅翻台率的统计：只要桌子被占着，就算客人已经买单在闲聊一小时，系统也会认为翻台率是 100%。但真正创造价值的只有上菜到吃完那段时间。对 GPU 来说，有价值的“上菜到吃完”，就是计算指令真正在流水线里被发射、被执行的周期。
1.2 “满载”背后的真实空闲：从硬件计数器说起
要看清真实情况，必须绕开 GPU-Util 这种粗粒度采样，直接看硬件计数器。NVIDIA 的Nsight Compute（ncu）和Nsight Systems（nsys）能给你远比 nvidia-smi 细的视角。
我在实际排查中，第一眼会看这几个指标：
- SM Busy：SM 上有 warp 处于 active 状态的时间比例。注意，active 不代表在算，可能只是在等。
- Issue Slots Busy：每个周期 warp scheduler 真正向执行单元发射指令的比例，这个才是“算力被用上”的核心指标。
- Executed Ipc Active：每个 SM 每个周期实际执行的指令数。
- Warp Stall 原因分布：warp 因为什么原因卡住了，是等显存（Long Scoreboard）、等计算依赖（Short Scoreboard）、还是等同步（Barrier）。
我实测过很多次，GPU-Util 99% 的时候，Issue Slots Busy 可能只有 20%-30%。也就是说，3/4 的时间 SM 虽然有 warp 驻留在上面，但没有任何一条指令被发射到计算单元。这 75% 的“伪活跃”时间，就是算力损耗的黑洞。
2. 第一性原理：SM 执行单元的“分工”与“饥饿”
2.1 从线程到 Warp：GPU 调度的最小单位是 Warp，不是线程
GPU 执行模型里最反直觉的一点：GPU 不是一个个线程独立调度，而是以Warp（32 个线程为一组）为单位，执行“锁步”指令。一个 warp 里的 32 个线程，在任意时刻必须执行同一条指令。这个设计是为了省硬件：一套取指译码单元管 32 个线程的数据通路，代价是灵活性受限。
Warp 是 NVIDIA 的术语，AMD 里对应的叫 Wavefront（64 线程一组），虽然宽度不一样，但思想是一样的。在写 CUDA kernel 或者阅读 profiling 报告时，所有调度、发射、stall 的分析单位都是 warp，不是 thread。
举个例子，你启动一个 kernel，声明了 1024 个线程。GPU 会把它们切成 32 个 warp，然后分发到不同的 SM 上。每个 SM 上有若干 warp scheduler（比如 A100 是 4 个），每个 scheduler 每个周期最多发射一条指令。所以一个 SM 每个周期最多发射 4 条指令，这就是该 SM 的指令发射带宽上限。
2.2 分支发散与资源占用：为什么一个 Warp 可能在“空转”
既然一个 warp 里的线程必须执行同一条指令，那遇到 if/else 分支就麻烦了。如果 32 个线程里 16 个走 A 路径，16 个走 B 路径，硬件只能先把 A 路径的 16 个线程“激活”，另外 16 个线程被屏蔽（predicated off），执行完 A 再反过来执行 B。每轮都有 50% 的 lane 在空转，这也就是所谓的分支发散（branch divergence）性能损耗。
分支发散只是 warp 空转的一种原因。更常见的是延迟未隐藏。SM 调度器每个周期都会检查有哪些 warp 是“eligible”（也就是本条指令的操作数已经准备好、可以发射）。一个 warp 如果正在等显存返回数据，它会进入 stall 状态，直到数据到齐才能再次成为 eligible。这时候调度器会切换到其他 eligible warp 来执行，用并行性掩盖延迟。
问题来了：如果 SM 上驻留的 warp 数量太少，没有足够的“替补队员”来切换，那调度器就真的无事可做，compute pipeline 就空了。这就是为什么occupancy（占用率）如此重要。我见过很多新手把 block 设得特别大、每个线程又用了海量寄存器，结果一个 SM 上只能驻留 8 个 warp，根本盖不住访存延迟，SM 计算单元自然饿肚子。
2.3 SM 内部执行单元的类型与利用率真相
一个 SM 里不是只有一种计算单元。以 Ampere 架构 GA100 为例，一个 SM 内部有：
- 4 个 warp scheduler，每个 scheduler 配一组执行单元；
- 每组的 FP32 单元（单精度浮点）和 INT32 单元通常各 16 个，这意味着一个周期内如果发射的是整数指令，浮点单元就闲着；
- LD/ST（访存）单元，负责访问 global memory 和 shared memory；
- SFU（特殊函数单元），负责 sin、cos、倒数等；
- Tensor Core，专门做矩阵乘加（HMMA）。
关键认知：GPU-Util 显示 100%，但它不会告诉你这 100% 时间里是哪些单元在干活。你的 kernel 如果充斥着内存寻址的整数运算，那 FP32 单元可能从头到尾都在“看戏”。指令混合（instruction mix）不合理，是算力虚高的第二大来源。
我在查看 ncu 的指令统计时，会专门看 FP32 Pipe、INT32 Pipe、Tensor Pipe 各自的利用率。曾经优化过一个稀疏检索 kernel，发现 FP32 利用率 85%，Tensor Core 利用率 0。因为 kernel 里都是 gather 操作，根本没用到矩阵乘法。这种时候你不能怪 GPU 偷懒，是你的算法压根没喂给它合适的“食物”。
3. 深入 Warp 调度：从“能发射”到“非发射”
3.1 Warp 的状态机：驻留、就绪与发射
要理解 Warp 为什么实际只用了 15% 的算力，需要把 Warp 调度状态机完整过一遍。一个 warp 从被分配到 SM 到执行完毕，主要停留在以下几种状态：
- 驻留（Resident）：warp 被加载到 SM 上，占据寄存器、共享内存等资源。但驻留不等于在执行。
- 就绪（Eligible）：当前指令的操作数准备好，可以被调度器发射。
- 等待（Stalled）：由于数据依赖、访存延迟、屏障同步等原因，指令暂时无法发射。
- 休眠（Inactive）：warp 还没被创建，或者创建了但在等资源（比如等上一个 block 释放 shared memory）。
调度器每个周期要做的事，就是从所有 resident 的 warp 中挑一个 eligible warp 发射指令。重点来了：如果只有 16 个 warp 驻留在 SM 上，其中 12 个因为访存暂时停滞，剩下 4 个 eligible，那么调度器的选择余地就只有这 4 个。如果这 4 个 warp 的指令恰好都是 16 位宽的整数操作，而你的计算单元以 FP32 为主，那发射了也是低效发射。
3.2 延迟隐藏需要多少 Warp？一个算给你看
用延迟理论来算一笔账。假设访存延迟是 600 个周期，你想把这些延迟完全隐藏，需要多少个 warp？
公式是：所需并发指令数 = 延迟 × 每周期发射能力。
以 A100 一个 SM 为例，假设忽略 scheduler 数量，粗略算：
- 访存延迟：约 600 周期；
- 每周期发射 1 条指令（假设 4 个 scheduler 每个发射 1 条，那峰值是 4 条/周期）；
- 如果按每个 scheduler 来算，需要 600 个独立指令在那等着。
一条 warp 一条指令，所以单个 scheduler 需要 600 个 eligible warp 才能完全掩盖延迟。但一个 SM 最多也才 64 个 warp 槽位。全部塞满也只有 64 个 warp。按单 scheduler 算，64 / 4 = 16 个 warp 对应一个 scheduler，仍然远小于 600。所以访存延迟不可能被完全隐藏，只能靠提高每周期发射的独立指令数（ILP）来弥补。
这就是为什么纯粹的访存密集型 kernel（比如 memcpy、element-wise 加法）远达不到算力峰值：计算单元在等数据面前显得苍白无力。你的 GPU-Util 100%、算力只有 15%，很可能是这种情况下最典型的症状。
3.3 寄存器资源是 Warp 数量的硬约束
很多人忽略了寄存器资源才是 warp 驻留上限的最大瓶颈。一个 SM 的寄存器文件是固定的，比如 A100 是 65536 个 32 位寄存器。如果每个线程用 32 个寄存器，那么 64K / 32 = 2048 个线程可以驻留，也就是 64 个 warp 满编。如果每个线程用了 128 个寄存器（深度学习 kernel 里很常见），那最多只能驻留 512 个线程，也就是 16 个 warp。occupancy 直接掉到 25%。
这 25% 意味着什么？调度器的“替补池”缩水了。一旦这 16 个 warp 里有 10 个在等数据，你的 SM 里计算单元空闲时间就占了绝大多数。这时候你可以做的优化之一是用__launch_bounds__限制线程的寄存器使用量，或者手动调maxrregcount。但注意，强行压寄存器会导致局部变量溢出到 local memory（实际上是 global memory 的一部分），产生额外访存流量，反而可能更慢。我在实战中会先在 ncu 里看 Local Memory 相关指标，确认 spill 是否严重再动手。
4. Tensor Core 与 Roofline 模型：算力分水岭的两种视角
4.1 Tensor Core 为什么能“1 打 N”
Tensor Core 是 NVIDIA 从 Volta 架构开始引入的专用矩阵乘加硬件。它的核心思路：与其让 32 个 FP32 单元每个周期算一堆独立的乘加，不如做成一个专门算 4×4×4 矩阵乘加的硬件块。以 A100 为例，FP32 CUDA Core 的峰值是 19.5 TFLOPS，而 TF32 Tensor Core 的峰值是 156 TFLOPS（约 8 倍），FP16 Tensor Core 更是到了 312 TFLOPS（16 倍）。到了 H100，这个差距被拉得更大，FP16 Tensor Core 峰值接近 990 TFLOPS（含稀疏约 1979 TFLOPS）。
为什么差距这么大？因为 Tensor Core 把矩阵乘加这个特定运算的数据通路做到了极致：单个指令（HMMA.16816.F32）就能完成一个 16×16×16 的矩阵乘加。这需要程序员或库以特定方式组织数据，把计算以“块”为单位喂给 Tensor Core，而不是一个标量一个标量地算。
这就是为什么实际工程里，能用 cuBLAS、cuDNN 调用的矩阵运算，尽量别自己写手写 kernel。手写很容易退化成标量循环，编译器无法自动向量化成 Tensor Core 指令。除非你直接用 WMMA API 或者 CUTLASS，否则大概率只能吃到 FP32 CUDA Core 的吞吐。
4.2 为什么用了 Tensor Core，总体利用率还是上不去
我在项目里遇到过 Tensor Core 利用率 90%，但整个 GPU 算力还是只有 40% 的情况。这就要说到 Roofline 模型了。
Roofline 模型的核心：一个计算平台的性能上限是“峰值算力”和“峰值带宽 × 算术强度”的最小值。其中算术强度 = 总浮点运算量 / 总字节访问量，单位是 FLOP/Byte。把它画成图，x 轴是算术强度，y 轴是可达到的算力。
- 低算术强度区域（数据密集型）：性能被内存带宽卡住，算力上不去，就算 Tensor Core 再快也没用；
- 高算术强度区域（计算密集型）：性能被峰值算力卡住。
以 A100 80GB 为例，HBM2e 带宽约 2 TB/s，FP32 峰值 19.5 TFLOPS。交叉点（ridge point）在 19.5e12 / 2e12 ≈ 9.75 FLOP/Byte。如果某个 kernel 的算术强度只有 2 FLOP/Byte，那理论可达性能上限就是 2 × 2e12 = 4 TFLOPS，约为 FP32 峰值的 20.5%。换句话说，哪怕你用 Tensor Core 在算，整体流水线也被“喂数据”这个动作卡死了。
这个判断在优化前一定要做。我见过不少团队花了大量精力把 FP32 算子改成混合精度，结果发现收益微乎其微，原因就是它们的 kernel 是带宽受限的，瓶颈在 2 TB/s 那条斜线上，而不是峰值算力那条水平线上。改精度只提高了水平线，斜线不动，整体上限还是被带宽锁死。用 Roofline 模型先定位自己站在哪个区域，比盲目优化高效得多。
4.3 实测算术强度：用 Nsight 验证
Roofline 分析最难的环节是测出真实的算术强度。理论 FLOPs 可以从算法推导，但实际内存流量比理论复杂得多：有 shared memory 和 global memory 的差异、有 cache 命中率的影响、有寄存器重用的因素。最靠谱的方式还是直接读硬件计数器。
在 ncu 里直接看 Memory Workload Analysis 的 DRAM Throughput 和 Compute Workload Analysis 的 SM 吞吐。把两者相除，得到的就是实测的算术强度。然后你把它放到 Roofline 图里看位置：
- 如果点在斜线区域下方，说明 kernel 连带宽上限都没摸到，还有巨大的优化空间；
- 如果点在斜线上，说明已经顶满带宽了，想继续提速只能减少内存访问量（更好的数据复用、更低精度的存储）；
- 如果点在水平线区域下方，说明计算单元没被喂饱，该考虑提高占用率、提升 ILP、或者启用 Tensor Core。
5. 实战定位：从 15% 算力反推可优化空间
5.1 用 Nsight 工具确认瓶颈的前三步
拿到一个“GPU-Util 100%、算力只有 15%”的案例，我不会立刻去改代码。先按下面三步定位：
第一步，用 Nsight Systems（nsys）看整体时间线。重点看哪些 kernel 占比最高、CPU 和 GPU 之间有没有空隙、H2D/D2H 拷贝是否在关键路径上。很多“GPU 利用率虚高”的场景其实是 CPU 侧数据预处理跟不上，GPU 在等 feed 数据。
第二步，对热点 kernel 跑 Nsight Compute（ncu），使用--set full或者按 Section 分析。先看 GPU Speed Of Light 的四个关键指标：Compute（SM）吞吐、Memory 吞吐、DRAM 吞吐、L1/L2 cache 吞吐。哪个接近 100%，哪个就是瓶颈所在。
第三步，看 Warp State Statistics 里的 Stall 原因分布。如果Long Scoreboard（等待 global memory 返回）占比特别高，就是访存延迟主导；如果Short Scoreboard（等待计算依赖）高，是指令级并行不足；如果Barrier高，可能是同步开销过大。
这三个步骤能帮你把问题从“算力虚高”收敛成“访存瓶颈”或“延迟瓶颈”或“指令混合问题”，然后优化方向就清晰了。
5.2 四类瓶颈的针对性优化策略
基于上面的定位结果，我按实际效果排序给出建议：
访存带宽受限：典型特征是 DRAM Throughput 接近 100%，SM 计算单元相对空闲。优化手段是提高数据复用：tiling 分块、用 shared memory 缓存反复使用的数据、减少重复读取。另外把数据读取改成向量化加载（float4 一次读 16 字节）能显著减少指令数。实测在一些 element-wise 的 kernel 里，单纯把 float 改成 float4 读写，性能提升 2-3 倍，因为访存指令减到了原来的 1/4。
访存延迟受限：典型特征是 DRAM 吞吐不高，但 warp 长期 stall 在 Long Scoreboard。优化手段是增加 resident warp 数（调低寄存器占用、增大 block 数），以及提高 ILP：让一个线程同时算多个独立的数据，这样即使单个 load 有延迟，同一 warp 里也有其他独立的乘加指令可以发射。
计算吞吐受限：如果 compute throughput 已经接近 100%，那就是真·算力瓶颈。这时看你的 kernel 有没有可能切换到 Tensor Core。对矩阵乘法，直接用 cuBLAS/cuDNN 的 API；如果自定义算子，可以尝试 WMMA API 或者让编译器自动生成mma指令。精度允许的话把 FP32 换成 TF32/FP16，单算子可能直接翻几倍。
指令混合不合理：如果整数运算占据了过多的发射槽，想办法把地址计算从主循环里提出来，或者用更大的 tile 减少重复寻址。也可以把整数运算迁移到 CPU 侧预处理，GPU 只负责“重活”。这类问题往往被忽视，但收益也不小。
5.3 常见问题速查表
| 症状 | 可能原因 | 确认工具 | 优先处理方案 |
|---|---|---|---|
| GPU-Util 100% 但算力低 | kernel 只是占用了时间片，SM 实际在等数据 | ncu 看 Issue Slots Busy、Warp Stall | 提高占用率、增加 ILP |
| SM 计算吞吐高但整体算力低 | kernel 未使用 Tensor Core | ncu 的 Tensor Pipe 利用率 | 改用到 Tensor Core 的库/API |
| 算术强度低于 ridge point | 访存带宽成为瓶颈 | ncu 的 DRAM Throughput | 减少内存访问、数据复用/向量化 |
| DRAM 吞吐不高但 warp 大量 stall | 访存延迟未隐藏 | ncu Warp State Stats 看 Long Scoreboard | 提高 occupancy、增加并行度 |
| 同步开销大 | block 间同步频繁 | ncu 看 Barrier stall | 减少协同同步点、增大任务粒度 |
| kernel 过多且单个都很小 | launch 开销被放大 | nsys 时间线 | kernel fusion、减少 launch 次数 |
你会发现，大部分“GPU-Util 100% 算力 15%”的案例，本质上都不是 GPU 坏了，而是喂给它的数据流、并行度或者指令模板出了问题。
5.4 一个真实优化案例：矩阵乘法的算力从 15% 到 80%
最后分享一个我实际处理过的简化案例，展示完整的优化链路。一个 M=N=K=2048 的 FP32 GEMM，初始用 naive 写法，每线程算一个输出元素，GPU-Util 显示 99%，实际 FLOPs 测出来约为 A100 FP32 峰值的 15%。
ncu 分析结果：DRAM 吞吐只有 20%，Long Scoreboard stall 占比 65%，说明是访存延迟主导。优化分几步走：
第一步，用 tiling 思路重写：每线程块负责 16×16 的输出块，通过 shared memory 缓存 A 和 B 的 tile。这一步把反复从 global memory 读 A/B 的操作减小了一个量级，DRAM 吞吐上来了，但 stall 依然高。
第二步，提高 ILP：让每线程一次算 2 个输出元素，这样同一时刻有更多独立的乘加指令可以隐藏 load 延迟。stall 降到了 40%，性能提升到 30% 左右。
第三步，调整寄存器分配和 block 配置，让 occupancy 从 25% 提到 50%。性能到 45% 左右。
第四步，换用 CUTLASS 或者让 cuBLAS 接管：cuBLAS 在 A100 上直接使用 Tensor Core，FP32 输入会走 TF32 路径，实测 FLOPs 能到峰值的 78% 左右。如果接受混合精度，FP16 路径能到 90%+。
这就是第一性原理的完整应用：先找到瓶颈在哪个维度（延迟），再用 hardware counter 验证，然后针对性地调整并行策略、资源和指令模板。整个过程我没有依赖任何“猜”，每一步都有数据支撑。
写到最后，说点我自己的体会。GPU 性能优化最大的坑，就是过早优化和凭感觉优化。看到一个 kernel 慢，很多人第一反应是“换更大显存”“把 batch 调大”或者“上多卡”，但往往忽略了最基础的 Roofline 定位。我现在的习惯是：任何 kernel 优化前，先在 ncu 里跑一遍 Speed Of Light，看一眼瓶颈是 Memory 还是 Compute，再看一眼算术强度在哪条线附近。这两步通常花不了 5 分钟，但能帮你避免浪费一整天的无效尝试。对于新手，我建议先别急着背各种 trick，先把 Roofline 模型、Warp 调度这两块地基打牢，你会发现自己看一眼代码就能大概判断出瓶颈在哪。这个能力，靠调参经验堆不出来，只能从第一性原理里长出来。
