把一台双卡 4090 的工作站放到你面前PyTorch 跑起来nvidia-smi里 GPU-Util 只有 30%温度一片冰凉——这种场面我见过太多次了。第一反应通常是显存不够驱动有问题但真正的答案往往在标题这行字里算力、带宽与搬运是一张三角账。AI基础设施这个词这两年被讲得越来越玄广告里动不动就是多少 P FLOPS。但落到实际项目最值钱的认知恰恰是这张三角账怎么算算力是加工能力带宽是运输通道的宽度搬运是数据真实发生的移动量和移动方式。现代计算尤其是大模型训练和推理绝大多数性能瓶颈早就不在能不能算而在数据能不能按时到位。这篇文章就是要把这三者的关系彻底讲透。适合两类人一类是自己部署 GPU 服务器、做大模型微调和推理的工程师另一类是刚接触 CUDA/GPU 编程总觉得自己代码比别人慢好几倍的开发者。读完你会明白为什么有时加算力没用为什么显存带宽卖那么贵以及遇到性能问题时到底该从哪一本账开始查。1. 为什么说 GPU 性能是算力、带宽、搬运三笔账而不是单纯看算力1.1 用后厨理解 GPU计算单元只是锅想象一家大型餐厅灶台很多、火力很猛这是招牌但出一桌菜靠的是传菜员把食材从冷库搬到备料台切配好再送进灶台。GPU 内部就是这样一个分工极细的后厨。寄存器和共享内存是厨师手边的调料与案板L2 缓存是后厨备料区显存HBM/GDDR是冷库再往外的 CPU 内存就是供货商仓库磁盘和网络连仓库都算不上是远郊物流。CUDA Core 和 Tensor Core 是那几十口锅但每一口锅开工之前必须先把数据从全局内存搬到寄存器或者共享内存。每往上走一层存储层次延迟差一个数量级寄存器访问几个周期L1 缓存大约 30 个周期HBM 显存几百个周期PCIe 拷贝一次动辄几十微秒。如果一个 kernel 的每个线程都在等全局内存的数据时间可能 90% 耗在等冷库开门上那时候你算力再多也是空转。我见过太多项目第一版代码跑得慢就怪 GPU 不行结果用 profiler 一照GEMM 本身只占了 30% 时间剩下全在搬数据。换卡是最贵的偷懒方式。1.2 Compute Bound 与 Memory Bound 是两种病在加速计算里任何 kernel 都可以先归个类计算密集Compute Bound数据喂得很及时瓶颈在 ALU、FMA、Tensor Core 的吞吐纯粹是锅不够快。访存密集Memory Bound计算单元大部分时间在等数据瓶颈在搬运速度。判断方法很简单看一个叫算术强度Arithmetic Intensity的指标单位是 FLOP/Byte意思是每搬运 1 字节数据能摊到多少次浮点运算。把它跟这台 GPU 的临界值比临界算术强度 峰值算力FLOP/s÷ 峰值带宽Byte/s任务的算术强度低于这个临界值就是 Memory Bound优化方向必须是减少搬运高于它才是 Compute Bound才能靠堆算力解决。这就是 Roofline 模型的核心思想搞 GPU 优化的人天天在用。举个例子。RTX 4090 的 FP32 峰值约 82 TFLOPS显存带宽约 1 TB/s临界算术强度约 80 FLOP/Byte。而一个向量加法每个元素读两个数、写一个数每元素搬运 12 字节只做 1 次加法算术强度 0.08 FLOP/Byte跟临界值差了三个数量级。所以写 vector add 这种程序换再好的卡也跑不满利用率——不是 GPU 不行是任务本身就是个体力活。1.3 为什么大模型时代带宽被推上风口Transformer 时代把带宽这个词推到了前所未有的高度。训练时各层激活值来回搬运推理时问题更典型逐 token 生成每出一个 token理论上都得把模型全部权重读一遍。7B 参数的 FP16 模型权重 14GB。以 4090 约 1TB/s 的显存带宽算读完一遍权重约 14 毫秒单流推理的 token 率上限就在 70 tokens/s 附近。这个数字跟我实测在 4090 上跑 7B Qwen 的结果非常接近——不是巧合是物理。推理阶段的大模型是典型的 Memory Bound单 token 的计算量其实很小大头全在权重读取。这也解释了为什么量化到 INT8 甚至 INT4 之后token 率能翻好几倍。很多人以为量化是数学魔法其实它砍的首先是搬运的字节数跟算力关系不大。后文会专门展开这笔账。2. 算力账FLOPS 数字背后的精度、核心与估算2.1 FP64、FP32、FP16、INT8各有各的主场先理清精度的差别。浮点数精度越高硬件电路越复杂同样面积能塞下的计算单元越少峰值速度就越低。精度典型用途相对峰值算力FP64科学计算、有限元、分子动力学、气候模拟最低通常只有 FP32 的 1/2 到 1/64FP32传统深度学习训练、图形渲染基准FP16 / BF16大模型训练与推理的主流精度通常为 FP32 的 2~8 倍含 Tensor CoreINT8 / INT4推理加速、量化部署通常再翻一倍以上为什么 FP64 这么贵因为 64 位浮点需要维护更大的指数范围和尾数精度加法器和乘法器电路大幅变宽功耗和面积都翻着倍涨。很多 GPU 的 FP64 单元压根是残废级的比如 RTX 40 系消费卡FP64 算力被砍到几乎没法用。看厂商宣传的算力数字一定要看精度条件。H100 标称 FP16 Tensor Core 近 2000 TFLOPS含稀疏但它的 FP32 只有约 67 TFLOPSFP64 更是只有约 34 TFLOPS。你用一张宣称近 2000 TFLOPS的 AI 加速卡去跑双精度地球物理模拟可能比老款 A100 还慢。算力不是单一数字是一张按精度区分的价目表。2.2 CUDA Core 与 Tensor Core两条算力流水线再往深一层GPU 内部的算力分两条流水线。CUDA Core 是传统通用计算核心擅长标量运算、分支多、不规则的任务精度可控。Tensor Core 是专门做矩阵乘加D A×B C的电路大模型时代绝对主力。它通过 cuBLAS、cuDNN 等库暴露能力PyTorch 里写一行torch.matmul底层就在自动选择最优的 Tensor Core kernel。这带来一个关键认知你写 CUDA 代码如果不用wmma或mma.sync这类矩阵指令只靠普通 CUDA Core 硬算矩阵乘算力直接掉一个数量级。很多时候我的矩阵乘怎么比别人慢 10 倍不是优化技巧问题是根本没走对流水线。生产环境里能用库函数就用库函数实在要手写 GEMM再考虑分块、双缓冲、流水线这些高级手法。2.3 需要多少算力6NF 估算公式项目立项时最常被问的问题训练一个模型到底要多少算力有一个久经考验的估算公式训练总浮点运算量 ≈ 6 × 模型参数量 × 训练 token 数这就是所谓 6NF。GPT-3 是 1750 亿参数、3000 亿 token按公式算约 3.15e23 FLOPs。如果用 H100 的有效 FP16 算力跑考虑实际集群利用率大概只有 40%~60%大约需要 8000 卡日以上。这也是为什么大模型训练一开就是上万张卡——这不只是算力账还有通信和调度损耗。粗算之外更实际的评估方法是先采样 profile拿一小批数据跑几步记录单步时间和显存峰值再外推到总步数。对 LoRA 微调这类改造任务公式估算容易偏实测最靠谱。我一般会准备一个三行的脚本加载模型、跑 20 个 batch、打印平均步时和峰值显存这比任何理论计算都有说服力。3. 带宽账从显存到 PCIe、NVLink 与集群网络每一级都在收过路费3.1 显存带宽为什么是 GPU 最贵的部件之一一张显卡的成本构成里显存占比极高尤其是 HBM。H100 的 HBM3 提供约 3.35 TB/s 带宽而消费卡的 GDDR6/GDDR6X 普遍在几百 GB/s 到 1TB/s 之间。HBM 贵有贵的道理DRAM die 垂直堆叠靠硅通孔TSV互联工艺复杂度高、良率低但带宽密度是 GDDR 的好几倍。很多人买卡只看 TFLOPS可一旦跑小 Batch 推理显存带宽才是决定生死的参数。同一张卡跑大 Batch 训练时可以吃到 80%~90% 利用率切到在线推理的小 Batch 立刻掉到 30% 以下。硅片没变变量的就是算术强度——大 Batch 让数据复用率上去了小 Batch 则是搬一次用一次。3.2 内存带宽与 token 生成率之间的关系给个可以抄作业的公式单 token 生成的理论上限 ≈ 显存带宽 ÷ 权重字节数还是 7B FP16 模型权重 14GB4090 带宽约 1TB/s上限约 70 tokens/s。量化到 W4A16权重 INT4、激活 FP16权重降到 3.5GB理论上限变成约 280 tokens/s。这就是为什么低延迟部署几乎必做量化——核心不是算法魔法是省带宽。长上下文场景还要把 KV Cache 算进去。上下文越长KV Cache 的读取量越可能超过权重本身进一步加重带宽压力。现在推理引擎普遍对 KV Cache 做量化KV8/KV4逻辑一模一样砍字节保速度。所以内存带宽与 token能成为高频问题我一点也不意外。生成速度慢头号嫌疑往往是带宽不是算力。带宽高的卡跑小模型token 率明显涨但计算利用率照样低——这是两套衡量体系不能混在一起看。3.3 多卡之间的过路费PCIe、NVLink 与集群网络单卡内部账算完了还有卡与卡之间的账。PCIe 4.0 x16 理论双向约 32GB/s实际单方向约 25GB/sPCIe 5.0 x16 翻倍到约 64GB/s 双向NVLink 完全不在一个量级A100 约 600GB/sH100 约 900GB/s双向合计跨节点靠 RDMA 网络InfiniBand 或 RoCE常见 200Gbps/400Gbps折算单方向约 25~50GB/s跟 PCIe 差不多。多卡训练每步结束要做 AllReduce 同步梯度。以 7B 模型 FP16 梯度 14GB 为例4 卡用 PCIe 互联的 Ring AllReduce要搬约 2×(n-1)/n×14 ≈ 21GB 数据。即便跑满 25GB/s也接近 1 秒。如果单步计算只要 2 秒通信就占了 30%。4 卡一定比 1 卡快 4 倍从来是句玩笑话——通信是实打实的额外开销。判断该不该加卡关键看每步通信量和计算量的比值。通信占比高时加卡是负优化通信占比低加卡才划算。灶台扩了传菜窗口不跟着扩出菜速度还是卡在短板上。4. 搬运账数据移动的成本往往比计算更贵4.1 最隐蔽的坑Host 与 Device 之间的拷贝GPU 编程里最常见的初学者杀手是主循环里频繁做 CPU 和 GPU 之间的数据拷贝。PyTorch 训完一个 batch随手.cpu()取 loss再.item()打印每步来一次同步性能直接骨折。为什么这么伤PCIe 拷贝不仅要搬数据还涉及 CPU 侧内存分配、驱动 DMA 映射以及 kernel 启动的全局同步开销。每次.cpu()或.numpy()默认触发一次 device synchronization把 GPU 流水线整个排空。正确做法用固定内存pinned memory通过torch.cuda.Stream做异步拷贝loss 打印改成每隔几步采样一次别每步同步DataLoader打开pin_memoryTrue和合理num_workers让数据预处理在 CPU 侧并行拷贝和计算重叠。大模型微调尤其容易踩这个坑。很多人用 Trainer 或自定义训练循环loss.item()写在每步里。其实改成每 N 步记一次或者干脆每个 epoch 末尾从累积变量里取训练速度能提升 5%~10%代码改动几乎为零。4.2 Kernel 执行全流程与 CTA、Warp搬运的最小单位再往深处走一次 GPU 运算的执行是一套完整流水CPU 下发 kernel 启动命令带网格大小、块大小、参数GPU 前端调度器把 grid 里的 Block 分发到各 SMBlock 在编程模型里称为 CTACooperative Thread Array是线程可协作、可同步的基本单位CTA 内部按 32 个线程一组切成 WarpWarp 才是硬件真正调度的单位一条指令同时驱动 32 个线程SM 靠大量并发 Warp 来隐藏访存延迟。很多人搞混 CTA 和 WarpCTA 是软件视角的线程块Warp 是硬件调度视角的最小执行单位。一个 128 线程的 CTA 会被切成 4 个 Warp如果 CTA 只有 31 个线程照样按 32 个槽位调度空了一个浪费一个。这里有个极其关键的搬运思想Warp 访存时最好 32 个线程访问连续且 128 字节对齐的地址这叫合并访存coalesced access。如果 32 个线程跳到天南海北读数据这个 Warp 的访存会被拆成一堆事务有效带宽大幅下降。所以同一个算法把数据布局从 AoS结构体数组改成 SoA数组结构体性能差 5 倍以上都不奇怪——算力一点没变纯是搬运方式的差异。4.3 内层搬运共享内存、寄存器与 kernel 融合SM 内部的顶配搬运是共享内存和寄存器。矩阵乘法用分块tiling就是经典范例每个 CTA 先把数据从慢速 HBM 搬到自己的共享内存案板上反复使用而不是每算一次就回冷库取货。与此对应的是 kernel 融合。逻辑上先 A 再 B的两个算子分开写会产生两个 kernel中间结果先写回显存再读出来来回搬运一次。融合成一个 kernel 之后中间结果直接留在寄存器或共享内存里衔接下一步。我在优化 CUDA 程序时最爱做的一手就是把两三个小算子合并比如把 elementwise 归一化和量化合并把 LayerNorm 的前后半段合并。小算子场景下融合往往带来 30% 以上的收益因为它砍的是搬运不是计算。PyTorch 的torch.compile和 Inductor 之所以有效底层也是在做自动 kernel 融合和算子调度优化。如果还在手写 CUDA建议先用torch.compile看看它能优化到什么程度再决定要不要手工下场。5. 实战一例如何用工具判断是算、带、搬哪条腿瘸了5.1 nvidia-smi 与 Nsight怎么用看什么遇到性能问题我按这套顺序排查nvidia-smi # 看利用率、显存、功耗、温度 nvtop # 实时看负载和显存占用 nsys profile python train.py # 看时间轴找 kernel 之间的空隙和同步等待 ncu --set full ./your_kernel # 针对单个 kernel 看计算/内存吞吐指标重点看两个指标SM 利用率Compute Workload Analysis和 DRAM 吞吐Memory Workload Analysis。如果 DRAM 吞吐已经接近峰值90%以上而 SM 利用率很低毫无疑问是带宽问题反过来 SM 利用率高、DRAM 吞吐低说明计算已吃满优化空间有限。有个常见误区nvidia-smi显示的 GPU-Util 是这段时间内 GPU 有没有活干的占比不是计算单元有多忙的精确度量。一个 Memory Bound 的 kernel 可能显示利用率 90%但 SM 内部大部分线程都在等数据。要精确定位得上ncu这类细粒度剖析工具。5.2 一个典型的算力充足但效率低案例去年我帮一个团队排查推理服务一台 A100 跑 13B 模型单并发实测 token 率只有 25 tokens/s远低于理论上限。排查过程先用nsys看时间轴发现 CPU 线程在每次前向里反复触发权重张量的contiguous()拷贝产生大量 D2D显存到显存搬运。原因很朴素模型加载时做了量化权重布局不是连续内存每次推理框架都自动修一次 layout。kernel profile 显示大量时间花在memcpyD2D和 elementwise mask 操作上真正的 GEMM kernel 只占不到 30%。修复方案也朴素加载时手动contiguous()并固化布局保证权重进 GPU 前就是连续内存再打开torch.compile把 elementwise 小算子融合。改完以后token 率从 25 升到约 60 tokens/s。算力没动、显存没动、一颗螺丝没拧只是把搬运的无效往返砍掉了。这个案例我每次讲都觉得值性能问题绝大多数时候不是纸面参数不够而是账实不符。5.3 优化手段的优先级排序确认瓶颈方向后按这个顺序动手先砍数据搬运减少.cpu()/.item()同步消除隐性contiguous()拷贝固定数据布局再改访存模式合并访存、共享内存分块、AoS 改 SoA然后做 kernel 融合合并小算子减少中间结果回写显存最后才考虑调并行和数学库细节增大 batch、开 BF16/FP16 混精度、确保走 Tensor Core 路径。千万不要一上来就各种__launch_bounds__和#pragma unroll花式调优。没有方向的微观优化只是自我感动。方向对的时候一个最朴素的改动就能带来数量级的差异。6. 算力集群、GPU 虚拟化与租卡决策三角账的放大版6.1 AI 算力集群的典型架构单卡账算清集群账就是这张三角账的横向放大。一个现代训练集群通常三部分计算层GPU 服务器一般 8 卡一台机内用 NVLink/NVSwitch 互联网络层跨节点的 RDMA 网络InfiniBand 或 RoCE配合 NCCL 做梯度通信存储层并行文件系统或高性能对象存储扛数据集读取和 checkpoint 落盘。在大模型分布式训练里通信和存储往往是集群 MFU实际有效计算利用率的最大杀手。单堆 GPU 数量不配相应的网络带宽和存储带宽利用率线性下降。NVIDIA 给 DGX 平台配那么高规格的 InfiniBand 和 NVMe 存储不是炫技因为这三样东西在同一条链路上互为瓶颈。软件栈现在比较成熟容器加 Kubernetes 做编排通过 Device Plugin 把 GPU 暴露给 Pod调度时配合 MIG 或 HAMi 做切分多卡训练用 NCCL 加大模型并行框架监控走 Prometheus 加 DCGM 暴露的指标。6.2 GPU 虚拟化、分时与配额生产环境很少让一个 Pod 独占整卡虚拟化方案常见三类MIGMulti-Instance GPUA100/H100 的硬件级切分每个实例有独立显存和计算切片隔离强但实例数量有限HAMi原 KubeGPU社区方案在 Kubernetes 里实现更细粒度的显存和算力虚拟化可以按显存大小申请也能限制算力配额时间片time-slicing把一张卡分时给多个 Pod灵活但隔离弱容易互相干扰。选型看业务训练任务通常要整卡或 MIG 实例推理任务模型小、QPS 不高时HAMi 细粒度分配更划算。我见过不少团队在 K8s 里直接写下nvidia.com/gpu: 0.5这样的请求底层却没配任何虚拟化组件调度器直接拒绝。配置虚拟化之前先确认 Device Plugin 和你用的方案兼容。还有一个高频坑忘了设算力配额上限。几个 Pod 共享一张卡只限显存不限算力其中一个跑满其他任务延迟全崩。正确做法是显存和算力两档都限。6.3 自己买卡还是租卡算清楚这笔账最后聊大家最纠结的问题自建机房还是租算力云。自建的账一台 8 卡 H100 服务器不含机房成本就奔着七位数去了三年折旧之后残值很低还得自己搞定电力、散热、网络、驱动、容器、监控、故障恢复。GPU 驱动和 CUDA 版本不匹配的报错比如requires device capability (9, 0) but your GPU has capability (12, 0)实质是你的 CUDA 工具链编译出的二进制只支持到 sm_90H100 那代而 GPU 是 sm_120 的新架构RTX 50 系 Blackwell。解法是升级 CUDA toolkit别去绕配置文件——绕来绕去还是得升。租算力云的账比如 AutoDL 这类平台按卡时计费把电力、散热、运维打包外包了。好处是用完即走不用备货坏处是数据进出流量、节点间互联带宽、高峰期排队都得在项目规划里留 buffer。我个人的经验原型验证和调参阶段弹性租 GPU 非常香长期稳定跑训练且运维能力够自建或包年更划算。但关键还是回到三角账——你要评估的不只是租到的卡上算力还有平台给你的节点间互联带宽以及你自己代码的搬运效率。很多人租了 8 卡训练速度却和 4 卡没差别原因往往是节点不在同一网络域跨机互访带宽根本扛不住 NCCL 的 AllReduce 流量。签单之前先问清楚互联拓扑和带宽保障。最后再分享一个小习惯每次拿到新 GPU 机器我第一件事不是跑模型而是跑一套 micro-benchmark——测单卡矩阵乘 TFLOPS、显存拷贝带宽、PCIe/NVLink 双向带宽和卡间延迟把基线数字记下来。之后跑任何任务先拿这套基线对一遍很快就能判断瓶颈在哪本账上。算力、带宽、搬运这三角账说穿了就一句话别让数据在路上等算力也别让算力饿着肚子等数据。