GPU利用率100%就是满负荷？剖析内存墙与Roofline模型下的性能瓶颈

先问你一个问题nvidia-smi里显示的 GPU-Util 已经 100%你的训练任务真的跑满了吗我上个月在一台 8 卡 H100 的机器上排查一个微调任务GPU Util 稳定在 99%Loss 却比预期慢了大半拍。用 Nsight 抓了一圈才发现SM 里的计算单元其实大部分时间在“等货”——数据从显存搬到寄存器的路上堵得死死的内核根本没在算只是在等。这其实就是 AI Infra 性能工程里最常遇到的内存墙问题。到了 Blackwell B200 这一代算力翻了 2.25 倍可“运数据的路”并没有等比例变宽结果就是如果你还在按老思路优化内核B200 上的有效算力可能连 H100 都跑不过。这篇文章我想从硬件参数、Roofline 模型、Profiler 实操和优化手段四个角度把“内核在等数据”这件事彻底讲透。不管你是做训练框架、推理服务还是单纯想把自己写的 CUDA kernel 跑得更快这篇都值得看完。1. 算力涨 2.25 倍“运数据的路”到底卡在哪里1.1 GPU 算力不等于有效算力先理解“内存墙”很多做 AI 的同学对算力有个误解以为 GPU 的 FLOPS 越高程序就一定越快。这就像厨房里厨师颠勺再快如果传菜通道只有一条窄走廊客人吃到的速度还是被走廊卡死。GPU 里的“厨师”是 SM 里的计算单元而“传菜通道”就是各级存储和总线——显存带宽、L2 带宽、片内互联、PCIe、网络全都算。所谓“内存墙”指的就是处理器算力增长速度远快于内存带宽增长速度导致程序性能越来越依赖数据搬运效率。这个矛盾在 CPU 时代就存在到了 GPU 时代只是被放大了。更扎心的是GPU 是个高度并行的机器它解决算力不足的办法是堆 SM、堆核心但数据要从显存搬进 SM物理上还是得走那几条有限的通道。通道宽度不涨就算 SM 再多也只能干等着。做 AI Infra 的人必须建立第一个直觉先看数据能从多快搬进来再谈计算有多快。如果数据搬不进来FLOPS 再高也只是纸面性能。1.2 B200 一代硬件参数变化到底意味着什么H100 到 B200单卡算力的提升幅度确实惊人。以比较典型的 2.25 倍算力增速为例纸面 FLOPS 上去了但显存带宽、Host 侧 IO、多卡互联带宽的增速完全没跟上。我根据自己的测试和公开资料整理了一张示意性对比表方便你看清差距维度H100 时代B200 时代变化趋势单卡算力FLOPS基准约 2.25 倍大幅提升显存带宽基准约 1.x 倍提升有限单卡可获得的“每算力带宽”基准相对下降更缺数据Host 到 GPU 的 PCIe/C2C 带宽基准基本持平或小幅提升几乎没宽多卡互联带宽基准稳步提升但落后于算力相对吃紧集群网络带宽基准受制于网卡与交换机最容易成瓶颈这里我不是要背参数而是想强调一个关键结论单位算力对应的带宽在下降。换句话说同样的代码在 H100 上是计算密集型的搬到 B200 上可能就变成访存密集型的了。这可不是代码写坏了而是硬件比例变了。再往深一层说B200 的显存带宽虽然也涨了但涨的幅度远不如算力涨的幅度。于是每个 FLOP 能分到的字节数更少了这就要求内核必须在“数据复用”上做得更极致。你如果只是把 H100 上能跑的 kernel 原封不动搬过来大概率会发现 SM 占用率没变但执行时间并没有按算力提升比例缩短因为瓶颈早就从“算不过来”变成了“喂不过来”。1.3 数据通路有四条别只盯着显存提到“运数据的路”大多数人第一个想到的是显存带宽。但我在实际调优中发现瓶颈常常不止这一条。按数据走向可以拆成四条通路每一段都可能成为“等数据”的源头第一条是 SM 到显存HBM的通路也就是 DRAM 带宽。这是最常见的瓶颈做访存密集型内核时基本都会踩到。第二条是 Host 到 GPU 的通路走 PCIe 或者 NVLink-C2C。很多人做推理服务时数据要从 CPU 内存拷进 GPU这条路的带宽和延迟直接影响端到端性能。第三条是 GPU 到 GPU 的通路走 NVLink 或 PCIe Switch。多卡训练时梯度同步、张量并行里的通信全部依赖这条通路。第四条是 GPU 到集群网络RDMA/网卡的通路。多机训练时数据要从远端存储或者另一台机器的显存里搬过来网络带宽往往是全局最窄的一段。我见过太多人优化了半天显存访问最后发现瓶颈居然在 PCIe 拷贝也见过有人在多卡训练里拼命调 kernel结果被 AllReduce 通信卡死。排瓶颈之前先搞清楚你的数据到底堵在哪条路上。这四条路的优化手段完全不同千万别拿同一个方子去治所有的病。2. 内核在等数据从 Roofline 模型看 Memory-Bound2.1 算术强度不达标内核就是在空转要判断一个内核是不是在等数据光靠感觉不行得有一个量化工具。业界最常用的是 Roofline屋顶模型——把每个内核的“算术强度”和硬件能力放到一张图里对比。算术强度的定义很简单算术强度 总计算量FLOPs / 总数据搬运量Bytes单位是 FLOPs/Byte。它描述的是“每搬进来一个字节能换来多少次计算”。如果这个值很高说明数据复用好内核偏计算密集型如果很低说明每个数据只用了一两次就算完了内核偏访存密集型。硬件本身也有一个临界值叫“拐点算术强度”。以 B200 为例算力越高、带宽相对越窄拐点就越靠右。也就是说同样的算术强度在 H100 上可能刚过拐点在 B200 上就掉到内存墙区域了。拿最简单的 Vector Add向量相加算一笔账读两个数组、写一个数组每个元素产生 3 次内存访问假设都是 4 字节 float那就是 12 字节计算只有 1 次加法和 1 次写回算 2 FLOP。算术强度大约是 2/12 ≈ 0.17 FLOPs/Byte。这个数值离 B200 的拐点差了两个数量级妥妥的访存密集——你哪怕把计算单元超频到冒烟性能也上不去因为时间全花在搬数据上了。2.2 为什么 GPU 利用率高但吞吐上不去我经常遇到有人拿 GPU-Util 100% 来证明“GPU 已经满负荷了”这是性能工程里最经典的误解。GPU-Util 这个指标NVIDIA 官方采样的是“在一个采样周期内是否有 kernel 在 SM 上运行”。只要有 kernel 在跑哪怕 SM 里所有计算单元都在等待显存数据返回也算作“忙”。所以你看到的 100%只能说明“GPU 没有闲着”不能说明“GPU 在拼命算”。真实情况往往是SM 里有大量 warp 在等待数据加载ALU算术逻辑单元的空闲率很高只是调度器还在不停地切换 warp让你从外面看起来像是满的。换句话说GPU-Util 衡量的是“占用”不是“吞吐”。要判断是不是在等数据必须去读 Profiler 里的 Memory Throughput 和 Compute Throughput。我自己的习惯是只要看到 Memory Throughput 接近峰值、Compute Throughput 却很低就直接判定为 Memory-Bound反过来才是 Compute-Bound。两个都高才是理想状态说明数据搬运和计算刚好重叠得很好。2.3 快速判断你的内核是不是在等数据没有 Profiler 在手边的时候也可以用四个问题快速自检第一个问题内核里每个数据元素被读进来几次如果只读一次就写完那基本可以断定是访存密集得想办法提高复用。第二个问题有没有可以让数据暂存在 shared memory 或寄存器里的机会比如相邻线程访问相邻数据理论上可以合并加载、多次复用如果没利用上就是在白白浪费带宽。第三个问题显存带宽占用率是多少如果 DRAM Throughput 常年在 90% 以上说明带宽已经打满优化方向应该是减少数据搬运量而不是优化计算指令。第四个问题kernel 运行时间是否远小于启动和调度开销如果一个 kernel 只跑几十微秒那 launch overhead 可能比计算本身还贵这时要优先考虑合并 kernel 或用 CUDA Graph。如果这四个问题里有三个以上“中招”你的内核大概率就是在等数据。3. 用 Profiler 给内核做一次“数据体检”3.1 Nsight Compute 的核心指标怎么读说起 GPU 性能分析首选工具就是 NVIDIA Nsight Computencu。它不是像nvidia-smi那样给你一个总览而是可以深入到每个 kernel、每个指令、每个 warp 的层面告诉你时间到底花在哪了。第一次用 ncu 的人往往会淹没在几百项指标里其实核心就几组。第一组是 Speed of LightSOL里的 ComputeSMThroughput 和 Memory Throughput它们直接告诉你这个 kernel 离计算峰值和带宽峰值各有多远。第二组是 Memory Workload Analysis里面会分 DRAM、L2、L1/shared 三层统计吞吐。DRAM 高说明数据确实在显存和 SM 之间反复搬运L2 高说明有一定局部性但没落到 L1L1/shared 高反而是好事说明数据已经在芯片内部复用了。第三组是 Occupancy 和 Warp State。Occupancy 代表 SM 上能同时驻留多少 warp太低会导致无法用并发掩盖访存延迟Warp State 里的 Stall Long Scoreboard 和 Stall Wait 是“等数据”的直接证据——scoreboard 就是等待显存数据返回的硬件机制它占比高就说明 warp 真的在等数据。3.2 一条命令跑出瓶颈分类在实际工作中我不太喜欢用 ncu 的 GUI 去逐个点指标更习惯在命令行里直接拉几个关键 metric一秒钟就能出结论。ncu --metrics gpu__time_duration.avg,sm__throughput.avg.pct_of_peak_sustained_elapsed,memory__throughput.avg.pct_of_peak_sustained_elapsed,dram__throughput.avg.pct_of_peak_sustained_elapsed,l1tex__throughput.avg.pct_of_peak_sustained_elapsed ./my_kernel跑完你会得到一张表格每一行对应一个 kernel。我一般先看两个数字sm__throughput和dram__throughput。如果sm__throughput很高比如 80% 以上而dram__throughput只有 30%说明计算单元很忙但显存很闲这是典型的 Compute-Bound优化方向是减少冗余计算或者用更低精度。如果反过来sm__throughput只有 30%dram__throughput却到了 90% 以上那就不用怀疑了内核就是在等数据。这时候再去调指令、调线程块大小都没用老老实实想办法减少访存量或者提高数据复用。如果两个都很低那就要注意了——可能是 launch overhead 太大或者线程块太小、SM 没喂饱这时候优先查 occupancy 和 wave 数。3.3 一个 Reduce Kernel 的剖析案例我带过一个实习生他写了个最朴素的 Reduce归约求和内核逻辑很简单每个线程读一个 float加到一个全局变量上。第一版代码大致长这样__global__ void reduce_v1(const float* in, float* out, int n) { int idx blockIdx.x * blockDim.x threadIdx.x; if (idx n) { atomicAdd(out, in[idx]); } }单看逻辑没毛病但 ncu 一跑DRAM Throughput 连 20% 都不到SM Throughput 更低。问题出在哪第一每个线程只读一个 float4 字节的访问没有利用 128-bit 向量化加载等于每次访存指令只搬了 1/4 的数据宽度。第二所有线程都往同一个全局地址做 atomicAdd会串行化冲突变成隐形的锁等待。第三根本没用 shared memory数据没有在块内先归并。后来改成两阶段每个线程连续读多个 float累加在寄存器里再用 shared memory 做块内归并最后每个 block 只产生一次 atomicAdd。改完以后DRAM Throughput 从 20% 拉到了 80% 以上kernel 时间直接降了一个数量级。这个案例特别典型——不是算法变复杂了而是让数据搬运的方式更符合硬件的脾气。4. B200 时代给内核减负的四个实操方向4.1 Kernel Fusion少搬一次数据就赢一次B200 算力上去了带宽相对更紧所以要尽量减少“数据进出显存”的次数。Kernel Fusion内核融合就是最直接的思路把多个连续 kernel 的中间结果留在 SM 内部不要写回显存再读出来。举一个常见的 LayerNorm 例子。很多初版实现会分三个 kernel先算均值、再算方差、最后做归一化。每个 kernel 都要把整个 tensor 从显存读一遍把中间结果写一遍。三个 kernel 加起来数据要在显存和 SM 之间来回跑六趟。融合版本只开一个 kernel每个线程块把一段数据加载到 shared memory然后在块内算均值、算方差、做归一化最后一次性写回。数据进出显存的次数从六趟降到了两趟读一次、写一次。在 B200 这种“算力充裕、带宽稀缺”的硬件上这种优化往往比从 FMA 指令里抠性能划算得多。4.2 Tiling 与 Shared Memory把数据切片放在手边对于矩阵乘、Attention 这类有大量数据重用的内核光靠 fusion 还不够必须做 Tiling分块。简单说就是把一个大矩阵切成很多个小块让每个小块被加载到 shared memory 后能被反复利用多次再换下一块。FlashAttention 就是把这个思路做到极致的代表。它不把完整的注意力矩阵写回显存而是按块分别计算 softmax 的局部统计量在片上完成整个 attention 输出。这么做最大的收益就是避免了大矩阵的中间读写把显存带宽的压力降了一个量级。我见过不少人想自己写 Attention kernel第一版就是教科书式的三步QK^T 写显存、Softmax 读写、再乘 V。在 H100 上还能忍换到 B200 上就彻底露馅了——因为中间的每一步都在搬运 O(N²) 的数据而 B200 的带宽又没跟上算力。所以我的建议很直接除非你只是想学习原理否则别手写 naive attention直接基于 FlashAttention 的思路改写。4.3 向量化加载与内存对齐让带宽跑满的细节很多 kernel 性能上不去不是算法不行而是连最基础的访存姿势都不对。GPU 显存带宽是按“事务”来算的一个事务通常是 32 字节甚至更大。如果每个线程只读 4 字节的 float整个 warp 读 32 个 float虽然地址连续但硬件要发多个事务才能搞定如果每个线程一次读 16 字节float4一个 warp 就能覆盖更大范围的连续数据事务数量大幅减少带宽利用率自然上去了。float4 v reinterpret_castconst float4*(in)[idx];这里要注意内存对齐。float4 要求 16 字节对齐如果数组头不对齐或者每个元素大小不是 4 的倍数强制转 float4 会直接崩。所以实际工程里我会先把数据补齐或者单独处理头部尾部不能对齐的部分。还有一个容易被忽略的点是“合并访问”。同一 warp 内相邻线程应该访问相邻地址这样硬件才能把多次访问合并成一个大的事务。如果线程和地址之间的映射是乱的比如按列访问行主序数组哪怕是带宽再大也白搭因为事务全被拆碎了。4.4 通信与计算重叠把等待时间“藏掉”多卡或多机训练里数据搬运的最高优先级往往不是显存而是通信。梯度同步、张量并行里的 AllReduce每一步都在等数据从别的地方送过来。通信不像计算它没有“算得快一点就能解决”的解药唯一的出路是让通信时间和计算时间重叠。具体做法有很多反向传播是按层从后往前算的每一层算完梯度后可以立刻把这个层的梯度发出去做 reduce而不是等所有层都算完再统一通信。这就是常见的梯度分桶bucket策略。再配合独立的通信 streamAllReduce就能和下一层的计算并行执行。到了多机场景网络带宽更是硬瓶颈。能走 RDMA 就别走 TCP能用 NVLink 域内通信就别跨机走网络。很多人以为多卡训练慢是 kernel 太差结果一测发现计算只占了 30% 时间剩下全在等梯度同步。遇到这种情况先别调 kernel去把通信矩阵画出来看看。在 B200 时代算力越强通信和计算重叠的收益就越大因为计算变快了通信反而成了显性的时间黑洞。5. 常见问题与排查技巧实录5.1 一张速查表应对“等数据”问题我把这些年排查“内核在等数据”问题的高频场景整理成了一张速查表碰到类似现象可以直接照着定位现象指标特征大概率原因优化方向内核很慢SM 和 DRAM 吞吐都很高Compute 和 Memory 都在 80% 以上数据量大且算法无法减少访存数据复用、低精度、分块SM 吞吐高DRAM 吞吐低Compute 80%Memory 30%计算密集但计算本身有冗余精简指令、查分支发散SM 吞吐低DRAM 吞吐高Compute 30%Memory 90%典型访存密集带宽打满向量化、shared memory、fusionSM 和 DRAM 吞吐都低两者都低于 30%线程块太少或 launch 开销大调 occupancy、合并 kernel、CUDA Graph多卡训练整体慢通信占比超过计算梯度同步没和计算重叠桶通信、独立 stream 异步数据拷贝到 GPU 很慢传输吞吐远低于峰值用了 pageable memory 或小包拷贝换 pinned memory、加大单次拷贝粒度这张表不一定覆盖所有情况但基本能帮你在一分钟内圈定问题方向。5.2 三个容易被忽略的实战坑第一个坑是在 kernel 内部做动态内存分配。malloc在 GPU kernel 里是允许的但代价很大——它可能触发隐式的全局锁导致所有线程排队而且分配出的内存全局可见访存局部性极差。B200 算力再强也填不了这个坑。正确做法是在 host 端预先分配好 bufferkernel 只负责读写。第二个坑是用cudaMemcpy传输小块数据。PCIe 传输有个特性单次传输的数据量越大有效带宽越接近峰值如果一传就是几百字节的小包绝大部分时间都耗在握手和延迟上。我见过有人循环里每次拷贝一个 batch 的标签数据结果 PCIe 利用率连峰值的一成都不到。正确做法是先把所有要拷的数据拼到一块连续内存一次拷完如果确实需要高频小数据考虑用 pinned memory 异步拷贝或者干脆放到统一内存里让硬件按需搬。第三个坑是只盯着 SM 吞吐忽略 L2 和 DRAM 的比例。有时候sm__throughput很低但l1tex__throughput很高说明数据其实已经缓存在片上只是线程间存在等待或分支发散。这时候你再调共享内存大小没什么用反而应该去看 warp 的调度和依赖关系。要综合读指标别靠单一数字下结论。5.3 我的实操体会宁可先算一道算术题再去调内核这几年做 AI Infra我最大的体会是拿到一个新内核别急着开 Profiler先拿出一张纸把算术强度算一遍。哪怕只是粗略估算也能立刻知道这个内核理论上该偏计算还是偏访存该往哪个方向优化。很多人调试了半天都是在错误的方向上用力。另一个体会是硬件的代际更替会悄悄改变优化优先级。H100 时代很多“不敏感”的访存模式到了 B200 时代就成了大瓶颈。所以每隔一段时间就要把线上正在跑的核心 kernel 重新 profiling 一遍看看瓶颈分类是不是已经变了。我吃过这个亏——有个老 kernel 一直没动过B200 上线后性能倒退一查才发现它早就 Memory-Bound 了只是以前算力不够显不出来。最后再分享一个小技巧排查“等数据”问题的时候可以顺手看一眼 memory clock 是否降频了以及 GPU 温度。很多时候性能上不去不是代码问题而是散热导致显存降频数据传输自然变慢。这种物理层面的坑Profiler 不会替你标注出来但实测影响能到 20% 以上。养成“先查物理再查逻辑”的习惯能省下大量无效调优的时间。