先别急着骂显卡是“虚标王”。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 Computencu和Nsight Systemsnsys能给你远比 nvidia-smi 细的视角。我在实际排查中第一眼会看这几个指标SM BusySM 上有 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 从线程到 WarpGPU 调度的最小单位是 Warp不是线程GPU 执行模型里最反直觉的一点GPU 不是一个个线程独立调度而是以Warp32 个线程为一组为单位执行“锁步”指令。一个 warp 里的 32 个线程在任意时刻必须执行同一条指令。这个设计是为了省硬件一套取指译码单元管 32 个线程的数据通路代价是灵活性受限。Warp 是 NVIDIA 的术语AMD 里对应的叫 Wavefront64 线程一组虽然宽度不一样但思想是一样的。在写 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 memorySFU特殊函数单元负责 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 到执行完毕主要停留在以下几种状态驻留Residentwarp 被加载到 SM 上占据寄存器、共享内存等资源。但驻留不等于在执行。就绪Eligible当前指令的操作数准备好可以被调度器发射。等待Stalled由于数据依赖、访存延迟、屏障同步等原因指令暂时无法发射。休眠Inactivewarp 还没被创建或者创建了但在等资源比如等上一个 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 TFLOPS16 倍。到了 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/sFP32 峰值 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 Systemsnsys看整体时间线。重点看哪些 kernel 占比最高、CPU 和 GPU 之间有没有空隙、H2D/D2H 拷贝是否在关键路径上。很多“GPU 利用率虚高”的场景其实是 CPU 侧数据预处理跟不上GPU 在等 feed 数据。第二步对热点 kernel 跑 Nsight Computencu使用--set full或者按 Section 分析。先看 GPU Speed Of Light 的四个关键指标ComputeSM吞吐、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提高占用率、增加 ILPSM 计算吞吐高但整体算力低kernel 未使用 Tensor Corencu 的 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%最后分享一个我实际处理过的简化案例展示完整的优化链路。一个 MNK2048 的 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 CoreFP32 输入会走 TF32 路径实测 FLOPs 能到峰值的 78% 左右。如果接受混合精度FP16 路径能到 90%。这就是第一性原理的完整应用先找到瓶颈在哪个维度延迟再用 hardware counter 验证然后针对性地调整并行策略、资源和指令模板。整个过程我没有依赖任何“猜”每一步都有数据支撑。写到最后说点我自己的体会。GPU 性能优化最大的坑就是过早优化和凭感觉优化。看到一个 kernel 慢很多人第一反应是“换更大显存”“把 batch 调大”或者“上多卡”但往往忽略了最基础的 Roofline 定位。我现在的习惯是任何 kernel 优化前先在 ncu 里跑一遍 Speed Of Light看一眼瓶颈是 Memory 还是 Compute再看一眼算术强度在哪条线附近。这两步通常花不了 5 分钟但能帮你避免浪费一整天的无效尝试。对于新手我建议先别急着背各种 trick先把 Roofline 模型、Warp 调度这两块地基打牢你会发现自己看一眼代码就能大概判断出瓶颈在哪。这个能力靠调参经验堆不出来只能从第一性原理里长出来。