1. 显存省下来了GPU 为什么还是吃不满显存优化做到位了nvidia-smi一看显存占用从 22GB 降到 9GB心里美滋滋结果watch -n 0.5 nvidia-smi盯了十分钟GPU-Util 始终在 35% 到 55% 之间来回晃功耗连 TDP 的一半都摸不到。这个场景我在调 LLM 训练任务时遇到过太多次了尤其是刚做完梯度检查点、ZeRO 切分、LoRA 微调这一套组合拳之后显存确实降下来了但吞吐量并没有线性上涨甚至有时候比优化前还慢。这篇文章就围绕这个具体问题展开显存优化之后GPU 利用率为什么上不去以及怎么一步步定位和解决。涉及到的技术栈主要是 PyTorch、CUDA、LLM 训练流程适合已经跑过至少一次完整微调任务、对显存优化有基本认知的读者。如果你还在纠结怎么把模型塞进 6GB 显存那可以先看显存优化的部分再回来看这篇。核心结论先放在前面显存和算力是两条独立的流水线显存省下来不等于计算单元能吃饱。GPU 利用率上不去通常不是显存的问题而是数据供给、算子效率、通信开销、调度策略这四个环节中至少一个卡住了。下面逐个拆。2. 先搞清楚 GPU 利用率到底在衡量什么2.1 GPU-Util 的真实含义很多人把nvidia-smi里的 GPU-Util 当成“GPU 有多忙”的指标这个理解不够准确。GPU-Util 的定义是在采样周期内GPU 上至少有一个 kernel 在执行的 time slice 占比。注意它衡量的是“有没有 kernel 在跑”而不是“SM 的利用率有多高”。这意味着一个极端情况如果 GPU 上只有一个线程在跑一个极小的 kernelGPU-Util 照样显示 100%。反过来如果 kernel 在等数据、等通信、等 CPU 发指令GPU-Util 就会掉下来。所以当你看到 GPU-Util 只有 40% 的时候真正的问题是有 60% 的时间GPU 上没有任何 kernel 在执行。这段时间 GPU 在干什么在等。2.2 用 DCGM 看更细的指标nvidia-smi的 GPU-Util 粒度太粗建议用 DCGMData Center GPU Manager或者nvidia-smi dmon看更细的指标nvidia-smi dmon -s pucvmet -d 1这个命令会每秒输出一行包含smSM 占用率比 GPU-Util 更准确mem显存控制器占用率pwr功耗temp温度smclkSM 时钟频率mclk显存时钟频率如果sm只有 40% 但mem到了 80%说明瓶颈在显存带宽不在计算。如果sm和mem都低但pwr也低那大概率是在等数据或者等通信。2.3 一个容易被忽略的指标SM 活跃度 vs 占用率SM 占用率SM Occupancy和 SM 利用率SM Utilization是两回事。占用率指的是一个 SM 上同时驻留的 warp 数量占最大支持 warp 数的比例。利用率指的是 SM 的运算单元有多少时间在干活。占用率高不代表利用率高。比如一个 kernel 里全是__syncthreads()或者全局内存访问warp 驻留了很多但都在等内存运算单元闲着。LLM 训练里这种情况很常见尤其是 attention 层的 softmax 和 layer norm 部分。3. 显存优化和计算效率之间的错位3.1 梯度检查点省显存但增加计算梯度检查点Gradient Checkpointing的原理是前向传播时不保存中间激活值反向传播时重新计算。显存占用从 O(n) 降到 O(sqrt(n))但计算量增加了大约 30%。这个 30% 的计算量增加在 GPU 利用率本来就不高的情况下会让情况更糟。因为重计算的过程本身也是 kernel 执行如果这些 kernel 效率低GPU-Util 看起来可能上去了但实际吞吐量没变甚至下降。我实测过一个 7B 模型的 LoRA 微调开梯度检查点之前 GPU-Util 大约 65%开了之后 GPU-Util 到了 78%但每秒处理的 token 数从 3200 降到了 2800。GPU 看起来更忙了但干的活更少了。3.2 ZeRO 切分通信换显存ZeRO Stage 2 和 Stage 3 把优化器状态、梯度、参数切分到多张卡上单卡显存确实降了但每步训练多了 all-gather 和 reduce-scatter 通信。如果卡间通信带宽不够比如 PCIe 而不是 NVLinkGPU 就会在等通信的时候闲置。判断方法很简单在训练脚本里加torch.cuda.nvtx.range_push/pop然后用nsys profile抓一下 timeline。如果看到大量ncclAllGather或者ncclReduceScatter的 gap那就是通信瓶颈。3.3 LoRA 微调小算子多kernel launch 开销大LoRA 的思路是冻结原模型参数只训练低秩矩阵。显存占用大幅下降但计算图里多了一堆小矩阵乘法。这些小算子的 kernel launch 开销在 GPU 计算中占比很高尤其是当 batch size 不够大的时候。一个具体的数字在 A100 上一个 kernel launch 的开销大约是 5 到 10 微秒。如果模型里有 200 个小算子每步就是 1 到 2 毫秒的纯 launch 开销。如果每步计算时间只有 50 毫秒那 launch 开销就占了 2% 到 4%。看起来不多但叠加起来就很可观了。4. 定位瓶颈的实操流程4.1 第一步确认是不是数据加载拖后腿数据加载是 LLM 训练里最容易被忽视的瓶颈。尤其是当你在做指令微调数据需要动态 tokenize、拼接、padding 的时候CPU 可能根本喂不饱 GPU。检查方法import time from torch.utils.data import DataLoader # 在训练循环里加计时 data_start time.time() batch next(dataloader_iter) data_time time.time() - data_start compute_start time.time() # ... 前向反向 ... compute_time time.time() - compute_start if step % 100 0: print(fdata: {data_time:.4f}s, compute: {compute_time:.4f}s)如果data_time和compute_time在一个量级甚至更大那 GPU 就是在等数据。解决方案DataLoader的num_workers调到 CPU 核数的 2 到 4 倍pin_memoryTruepersistent_workersTruePyTorch 1.7如果 tokenize 是瓶颈提前把 tokenize 好的数据存成二进制文件训练时直接 mmap 读取4.2 第二步用 PyTorch Profiler 看算子耗时PyTorch 自带的 profiler 是最方便的工具from torch.profiler import profile, ProfilerActivity with profile( activities[ProfilerActivity.CPU, ProfilerActivity.CUDA], scheduletorch.profiler.schedule(wait1, warmup1, active3), on_trace_readytorch.profiler.tensorboard_trace_handler(./log), record_shapesTrue, profile_memoryTrue, ) as prof: for step, batch in enumerate(dataloader): if step 5: break train_step(batch) prof.step()跑完之后用 TensorBoard 打开./log看 CUDA kernel 的耗时分布。重点关注有没有耗时特别长的单个 kernel有没有大量耗时极短小于 10 微秒的 kernelkernel 之间的 gap 有多大如果 gap 很大说明是 CPU 侧的问题数据加载、Python 开销、kernel launch。如果单个 kernel 耗时很长说明是算子本身的问题。4.3 第三步检查 CPU 侧开销Python 的 GIL 和解释器开销在 LLM 训练里不可忽视。尤其是当模型 forward 里有大量 Python 控制流的时候。一个简单的判断方法把模型用torch.jit.script或者torch.compile包一下看吞吐量有没有提升。如果有明显提升说明之前 CPU 侧开销很大。# PyTorch 2.0 model torch.compile(model)torch.compile会把 Python 控制流编译成 TorchScript 或者 Triton kernel减少 CPU 和 GPU 之间的同步次数。4.4 第四步看通信占比多卡训练时用nsys抓 timelinensys profile -o profile_output -t cuda,nvtx python train.py然后用nsys-ui打开profile_output.nsys-rep看 NCCL 通信的时间占比。如果通信时间超过总时间的 20%就需要考虑换用 NVLink 而不是 PCIe增大 batch size减少通信频率用梯度累积把多个 micro-batch 的梯度攒起来再通信5. 常见瓶颈的排查速查表现象可能原因排查方法解决方案GPU-Util 低CPU 占用高数据加载瓶颈计时 data_time vs compute_time增加 num_workers预 tokenizeGPU-Util 低CPU 占用也低通信瓶颈或同步等待nsys 看 NCCL 占比增大 batch梯度累积换 NVLinkGPU-Util 高但吞吐低小算子过多profiler 看 kernel 数量torch.compile算子融合显存降了但速度没变计算量增加抵消了显存收益对比优化前后的 tokens/s权衡显存和计算选择性开启单卡正常多卡变慢通信开销对比单卡和多卡 timeline调整 ZeRO stage用 DDP 替代6. 几个我踩过的坑和对应的解法6.1 坑一num_workers 设太大反而慢DataLoader的num_workers不是越大越好。每个 worker 是一个独立进程进程间切换和内存拷贝有开销。我试过在 64 核机器上设num_workers32结果比num_workers8还慢。原因是 worker 之间的锁竞争和内存带宽争抢。经验值num_workers设为 CPU 物理核数的 1/4 到 1/2然后根据实测调整。6.2 坑二pin_memory 没开pin_memoryTrue会把数据放到页锁定内存里GPU 可以直接通过 DMA 读取不需要 CPU 中转。这个选项在数据加载快的时候效果不明显但在数据加载慢的时候能省不少时间。注意pin_memoryTrue会占用更多主机内存如果主机内存紧张要权衡。6.3 坑三梯度累积和 batch size 的关系没搞对梯度累积的常见错误是accumulation_steps4batch_size8以为等效于batch_size32。实际上如果模型里有 BatchNorm 或者依赖 batch 统计量的层梯度累积和真实大 batch 不等效。对于 LLM 来说LayerNorm 不依赖 batch 统计量所以梯度累积基本等效。但要注意学习率要相应调整通常梯度累积时学习率要按sqrt(accumulation_steps)缩放。6.4 坑四torch.compile 不是万能的torch.compile在 PyTorch 2.0 之后确实能提升不少模型的吞吐量但它对动态 shape 的支持有限。LLM 训练里如果序列长度变化频繁torch.compile会反复重新编译反而更慢。解决方案把序列长度 pad 到固定值或者用dynamicTrue让编译器处理动态 shape但性能会打折扣。6.5 坑五忽略了 GPU 频率和功耗墙有时候 GPU-Util 上不去是因为 GPU 被功耗墙或者温度墙限制了。用nvidia-smi -q -d POWER看当前功耗和功耗上限。如果功耗已经贴着上限SM 频率会自动降下来计算速度就慢了。在数据中心里这种情况通常不是问题。但在个人工作站或者租用的机器上散热不好或者电源功率不够就会触发降频。7. 一个完整的调优案例7.1 初始状态模型7B LLMLoRA 微调单卡 A100 40GB。 数据指令微调数据集约 50 万条平均长度 512 token。 配置batch_size4gradient_accumulation_steps8fp16开梯度检查点。现象显存占用 18GBGPU-Util 45%吞吐 1800 tokens/s。7.2 排查过程第一步计时数据加载data: 0.12s, compute: 0.35s数据加载占了 25% 的时间偏高。第二步profiler 看 kernel发现大量耗时小于 5 微秒的小 kernel主要是 LoRA 的 A/B 矩阵乘法。第三步看通信单卡训练没有 NCCL 通信。7.3 优化措施num_workers从 4 调到 12pin_memoryTruepersistent_workersTrue。数据加载时间降到 0.04s。把 LoRA 的 A/B 矩阵乘法合并成一个 kernel用torch.compile编译 LoRA 层。梯度检查点从每层都开改成每隔一层开减少重计算量。batch_size从 4 调到 8gradient_accumulation_steps从 8 降到 4。7.4 优化结果显存占用从 18GB 升到 22GB因为 batch size 大了GPU-Util 从 45% 升到 72%吞吐从 1800 tokens/s 升到 3400 tokens/s。这个案例说明显存省下来不是目的吞吐量才是。有时候适当多占一点显存换更高的 GPU 利用率整体收益更大。8. 一些底层原理的补充8.1 GPU 的流水线结构GPU 的计算流水线大致是指令发射 - warp 调度 - 运算单元执行 - 内存访问。任何一个环节卡住后面的环节就闲着。LLM 训练里最常见的卡点是内存访问。Attention 的 QK^T 和 softmax 需要频繁读写全局内存如果 batch size 不够大内存带宽利用率上不去运算单元就等数据。8.2 Occupancy 和 ILP 的权衡Occupancy 是 SM 上驻留的 warp 数量。ILPInstruction Level Parallelism是单个 warp 内指令的并行度。两者是互补的Occupancy 高可以掩盖内存延迟ILP 高可以减少对 Occupancy 的依赖。LLM 训练里矩阵乘法部分 ILP 很高Occupancy 低一点没关系。但 element-wise 操作如 LayerNorm、激活函数ILP 低需要高 Occupancy 来掩盖延迟。8.3 Kernel Launch 开销的量化在 A100 上一个 kernel launch 的典型开销是 5 到 10 微秒。如果模型有 500 个 kernel每步的 launch 开销就是 2.5 到 5 毫秒。如果每步计算时间是 100 毫秒launch 开销占 2.5% 到 5%。减少 kernel launch 的方法算子融合用 Triton 或者 torch.compile用 CUDA Graph 把整个计算图捕获下来一次性提交增大 batch size让每个 kernel 干更多活CUDA Graph 在 LLM 推理里用得比较多训练里因为动态 shape 的问题用得少。但如果序列长度固定CUDA Graph 也能用。9. 工具和命令速查# 实时监控 GPU nvidia-smi dmon -s pucvmet -d 1 # 抓取 timeline nsys profile -o output -t cuda,nvtx python train.py # 查看 GPU 功耗和频率 nvidia-smi -q -d POWER,CLOCK # 查看 NCCL 通信 export NCCL_DEBUGINFO python train.py 21 | grep NCCL# PyTorch profiler from torch.profiler import profile, ProfilerActivity with profile(activities[ProfilerActivity.CUDA]) as prof: train_step() print(prof.key_averages().table(sort_bycuda_time_total))# 检查数据加载时间 import time for step, batch in enumerate(dataloader): t0 time.time() batch {k: v.cuda() for k, v in batch.items()} t1 time.time() loss model(**batch).loss loss.backward() optimizer.step() t2 time.time() if step % 100 0: print(fdata transfer: {t1-t0:.4f}s, compute: {t2-t1:.4f}s)10. 最后几点个人体会显存优化和 GPU 利用率是两个独立的问题不要指望省了显存就能自动跑满 GPU。每次做完显存优化都要重新测一遍吞吐量确认没有引入新的瓶颈。数据加载是最容易被忽视的环节。很多人花大量时间调模型结构结果瓶颈在DataLoader上。先把数据管线调好再调模型。torch.compile和 CUDA Graph 是减少 kernel launch 开销的利器但都有适用条件。torch.compile适合静态 shapeCUDA Graph 适合固定计算图。用之前先确认自己的场景是否匹配。多卡训练时通信开销往往比计算开销更难优化。如果卡间带宽不够再好的单卡优化也白搭。选机器的时候NVLink 和 PCIe 的差距在 LLM 训练里非常明显。GPU 利用率不是越高越好。有些 kernel 天生就是内存密集型的SM 利用率上不去很正常。关键是看端到端的吞吐量而不是盯着nvidia-smi的数字。