做AI基础设施这一行，我特别怕听到一句话：“这块卡标称多少TFLOPS？”问出这句话的人，通常已经把GPU默认成一台“算数特别快的机器”。但真正把训练和推理集群跑起来的人都知道，一张卡的落地性能，从来不是算力一个数字说了算，而是三条线同时结算的结果：算力能出多少活、带宽能喂多少料、搬运能把多少数据准时送到。所谓“加速计算”，说穿了就是这三角账怎么平衡的问题。
这篇内容就是《AI基础设施系列》的第一篇，把这三笔账拆开讲清楚：算力是什么、带宽卡在哪、搬运有多贵。适合刚接触GPU编程的朋友，也适合做大模型训练、部署、买卡、租云GPU，或者给集群做调度的同学。读完你起码能给自己手头的任务做一次“三角审计”，不再被一张规格表带走。
1. 三角账怎么来的：算力、带宽与搬运各占一角
1.1 把GPU当成一家工厂看
我习惯把GPU想象成一家工厂，这样很多性能问题就特别好解释。
- SM（流式多处理器）和CUDA Core是车间里的机床，负责真正的计算；
- 显存是原料仓库，模型权重、激活值、中间结果都堆在这；
- 显存带宽是车间到仓库之间的传送带，决定机床能多快拿到原料、多快把成品送回仓库；
- PCIe、NVLink、网卡是厂区外的货运网络，负责把数据从一个地方搬到另一个地方；
- Kernel启动和copy引擎则像调度员，每一次“发车”都有固定的手续和开销。
一次训练迭代，拆开看就是：数据从磁盘/网线搬进显存（物流），权重和梯度流过计算单元（传送带供料），中间结果在显存里反复读写（仓库内部周转），最后更新完的权重还要同步出去（再走一趟物流）。任何一个环节堵住，机床就得空转。
这三角各有各的计价单位：算力大体用TFLOPS量化，带宽用GB/s量化，搬运则要算“每一次搬多少字节、一共搬多少次”。把三本账加在一起，才是你真正感受到的训练时间或推理延迟。
1.2 只看FLOPs的翻车现场
FLOPs是峰值，不是实际产出。它描述的是“机床全速运转时的最大加工量”，但要达到这个数，工作负载的形状、数据排布、流水线调度都得配合到位。现实里绝大多数任务根本摸不到峰值。
我见过不少配置单，两张卡算力差了三倍，结果跑同一个小模型，耗时却差不多。原因往往是带宽和搬运已经把天花板焊死了。比如一个逐元素加法的任务，每读一个元素才做一次加法，算力再高也使不上劲，真正决定速度的是显存带宽。这种情况下，把算力从50TFLOPS换成100TFLOPS，收益可能只有百分之几。
还有一种更常见的翻车场景：数据在CPU和GPU之间来回搬，GPU大部分时间不是在算，而是在“等货”。很多新手把数据从Pandas里切出来再tensor.to('cuda')，每步都做全量拷贝，结果一看Nsight，memcpy占了四成时间，算力利用率低得可怜。
所以“三角账”的本质是：任何一个角成为短板，都会拖住整体。优化的目标不是把某一角拉满，而是让三角互相匹配。
2. 算力这一角：CUDA Core、Tensor Core与精度档位
2.1 warp、CTA与Tensor Core：算力是怎么组织起来的
先回答一个很多人第一次接触CUDA时会懵的问题：cooperative thread array（CTA）和warp到底是什么关系。
CTA是软件层面的概念，你可以把它理解成“一个班组”。一个CTA里的线程可以协作，能通过__syncthreads()互相等待，能共享一块显存。而warp是硬件调度的最小单位，固定32个线程，由GPU的调度器一次性打包发射。一个CTA里通常包含一个或多个warp，比如CTA=128线程，那它会被拆成4个warp来执行。
搞清楚这两层，你才能理解为什么GPU擅长矩阵乘法但又不擅长所有矩阵运算。矩阵乘法是大量同构运算：数据复用率高、分支少、指令开销低，非常适合SIMT模式批量处理。而跳跃访问多的逻辑、稀疏数据结构，GPU跑起来就特别别扭，因为“送料”比“加工”更费劲。
Tensor Core则是另一种“压路机”：它不是通用计算单元，而是专门吃矩阵乘加（GEMM）的专用电路，能一次完成一块小矩阵的乘加。训练和推理之所以这几年突飞猛进，很大程度就是靠它把矩阵乘的每瓦性能抬上去了。
2.2 精度换算力：FP32、FP16、BF16、INT8的性价比
Tensor Core支持的精度档位是分级的：FP32、TF32、FP16、BF16、FP8、INT8。大体规律是位宽每降一档，吞吐翻一番。拿常见卡举例（均指Tensor Core密集计算）：
| 卡 | FP32（TFLOPS） | FP16 Tensor（密集，TFLOPS） | 备注 |
|---|---|---|---|
| A100 SXM 80GB | 19.5 | 312 | FP16是FP32的约16倍 |
| H100 SXM 80GB | 67 | 495 | FP16是FP32的约7倍 |
| RTX 4090 | 82.6 | 330 | 消费卡也能吃到Tensor红利 |
为什么低了精度就快了？两方面：一是同样面积的电路可以塞更多MAC单元，二是数据位宽减半后，内存带宽的消耗也减半。大模型训练普遍用FP16/BF16混合精度，推理端再用INT8/FP8做量化，都是冲着这层性价比去的。
代价也很现实：低精度容易溢出、误差大。FP16的范围比FP32窄得多，训练里梯度可能直接变成0，所以需要loss scaling这类技巧；推理量化则要做校准，不然模型输出可能稀烂。精度不是越低越好，是“够用且能省则省”。
2.3 算力版本不匹配：capability (9,0)与(12,0)是怎么回事
最近热词里有一条特别典型：requires device with capability <= (9, 0) but your gpu has capability (12, 0)。翻译成人话就是：你运行的程序只带了算力9.0及以下GPU的程序产物，但你的显卡硬件是12.0的新架构，两者对不上。
这里的9.0、12.0叫compute capability，是GPU硬件代际编号。比如A100是8.0，H100是9.0，RTX 40系是8.9，RTX 50系（Blackwell）是12.0。CUDA程序在编译时会针对某个capability生成SASS（底层机器码），这份机器码老架构产品往往没法在新架构上直接执行，除非程序里同时打包了PTX（中间代码）由驱动JIT转译。老版本的PyTorch、TensorRT引擎、cuDNN，通常只带旧capability的产物，插上新Blackwell卡就会在加载阶段报这个错。
踩过坑之后我的处理顺序是：
- 先把NVIDIA驱动更新到支持新架构的版本；
- 换用支持新capability的CUDA运行库，比如PyTorch装cu128及以上的wheel；
- 再用
python -c "import torch; print(torch.cuda.get_device_capability())"确认识别到的能力是不是(12,0)； - 如果还跑自定义C++/CUDA扩展，就设置
TORCH_CUDA_ARCH_LIST="12.0"重新编译。
千万别去“改”capability，那是硬件属性，改不了。也别以为新显卡向下兼容所有老库，现实是二进制产物经常不兼容。
3. 带宽这一角：内存带宽才是大多数场景的真瓶颈
3.1 操作强度：算力和带宽之间的“汇率”
判断一个任务到底卡在算力还是卡在带宽，有个非常实用的指标叫“操作强度”（Arithmetic Intensity，单位FLOP/Byte），也就是“每读一个字节的数据，能顺便完成多少次浮点运算”。
举两个极端例子：
- 逐元素加法：读两个数、写一个数，每个元素才做2次浮点运算，操作强度大约0.2 FLOP/Byte；
- 大矩阵乘法：一个4096×4096的GEMM，算下来操作强度能到1300多 FLOP/Byte。
再把卡的“峰值算力÷峰值带宽”算出来，这是Roofline模型里的“脊点”。以H100为例：约495 TFLOPS的FP16密集算力除以3.35 TB/s的带宽，脊点大约148 FLOP/Byte。任务的操作强度高于148，算力才能吃饱；低于148，就算计算单元再快，也是带宽先到顶。
所以“提高算力”对带宽瓶颈任务几乎无效。机器再多也快不起来，因为传送带就那么大。这解释了很多迷惑现象：GPU跑矩阵乘法飞快，跑哈希、排序、稀疏操作却像蜗牛。
3.2 一张表看清主流GPU的算力与带宽
采购和调优前，先把规格表摊开看：
| GPU | FP32（TFLOPS） | FP16 Tensor（密集，TFLOPS） | 显存带宽（GB/s） | 显存 |
|---|---|---|---|---|
| RTX 4060 Laptop | 14.6 | 未单独列 | 256 | 8GB GDDR6 |
| RTX 4090 | 82.6 | 330 | 1008 | 24GB GDDR6X |
| A100 SXM 80GB | 19.5 | 312 | 2039 | 80GB HBM2e |
| H100 SXM 80GB | 67 | 495 | 3350 | 80GB HBM3 |
注意这里有个反常识点：RTX 4090的FP32算力比A100还高，但它的显存带宽只有A100的一半左右。真跑起大模型推理，A100往往更稳，因为它不被带宽卡得那么死。消费卡的浮点看起来很猛，但“供料”跟不上，很多任务根本发挥不出账面数字。
3.3 用带宽直接算大模型推理的token速度
大模型自回归推理有一个非常硬核的物理约束：每生成一个token，都要把模型全部权重从显存读一遍。所以理论上的单用户token速度大概等于：
token/s ≈ 显存带宽 ÷ 模型权重体积
7B模型用FP16存，大约14GB。那么：
- RTX 4090（带宽1008GB/s）大约72 token/s；
- A100（2039GB/s）大约145 token/s；
- H100（3350GB/s）大约239 token/s。
70B模型则是140GB，H100也只能跑24 token/s左右，H200（4.8TB/s）能到34 token/s。这就是为什么大模型推理天生是带宽生意，而不是算力生意。
理解了这一点，很多工程决策就顺理成章了：
- 为什么大家都往INT8、4bit量化上冲？因为权重体积减半再减半，token速度近乎线性翻倍；
- 为什么prefill（长prompt处理）和decode（逐个生成）要分开调度？因为prefill是计算密集，靠Tensor Core；decode才是带宽密集，靠显存带宽和缓存；
- 为什么单卡H100理论上也就能同时服务十几个用户？因为要满足每人每秒输出20个token左右，就要占掉每秒200多个token的带宽额度。
下次有人问“这卡能扛多少并发”，先别聊算力，拿带宽除一除，心里就有底了。
4. 搬运这一角：PCIe、NVLink与kernel启动的隐性账单
4.1 CPU与GPU之间：一次memcpy到底多贵
GPU自己再快，数据也得先到显存里。CPU和GPU之间的通道主要是PCIe。
单条PCIe Gen4 x16的单向理论带宽是32GB/s，实际传输效率大约24~28GB/s。听着不低，但算一笔账就心疼了：往显存里搬一个2GB的权重文件，大约要80毫秒。如果训练脚本每个step都做一次全量拷贝，1000步就是80秒纯等待。而这80秒里，GPU的计算单元是彻底闲着没事干的，功耗和成本一点没省。
有三个马上能做的事：
DataLoader(pin_memory=True)：固定页内存能绕过驱动中转，搬得快不少；- 用CUDA stream做异步拷贝，把搬运和kernel计算重叠起来；
- 认真审视代码里有没有“每步都to(device)”的无谓搬运，把数据预处理留在GPU上做。
很多“慢训练”查到最后，根本问题不是算力不够，而是CPU里用Pandas折腾半天再一次性搬上显卡，账单全记在搬运这一栏。
4.2 GPU与GPU之间：NVLink不是每张卡都有
单卡放不下模型，就要多卡跑。多卡之间梯度同步的本质也是搬运。
A100的NVLink总带宽约600GB/s，H100约900GB/s，比PCIe快一两个数量级。但消费级RTX 4090没有NVLink，多卡只能走PCIe，梯度同步慢得想哭。这也是为什么很多人组了几张4090做分布式训练，发现“加卡反而更慢”——通信账单把收益吃光了。
具体算一下：7B模型FP16梯度约14GB，8张卡做一次全量allreduce，ring算法下每个rank要搬运大约24.5GB。在NVLink 600GB/s上约41毫秒，在PCIe Gen4 x16上要近一秒。如果你的训练step本身就一两秒，那这笔通信开销占比相当可观。
跨节点就更夸张了。200Gbps的InfiniBand约25GB/s，千兆以太网只有0.125GB/s。算力再猛，数据在网络上搬不过去，一切都是白搭。这也顺带解释了一个现象：现在大家都在做东西，本质上是想办法让“搬运与算力”在多个任务之间匹配，而不是让一张大卡被一个任务独占后、其余时间空转到天荒地老。
4.3 Kernel启动开销：小算子如何拖垮大任务
除了显式搬运，还有一个隐藏账单：kernel启动。
每次在GPU上执行一个kernel，CPU要打包参数、下发给驱动、进队列、等待GPU调度执行、再通知CPU结果。这一来一回的固定开销大约5~10微秒。单个kernel看似不贵，但架不住数量多。我分析过一个朴素的训练脚本，一个step里跑了三千多个小算子，光是启动开销就占了三四十毫秒，而算子本身的运算可能只有几毫秒。
对策主要三条：
- kernel融合：把连续逐元素操作合并成一个kernel，减少启动也减少中间张量的显存读写；
- CUDA Graph：一次性捕获整个计算图，重复使用时启动开销几乎归零；
- torch.compile或融合库：它们自动帮你做上述合并，实测很多小模型能白捡一截速度。
另外提醒一句：小数据别迷信UVM和zero-copy。显卡访问零拷贝内存时可能触发缺页中断，数据被切成小片慢慢搬，延迟反而比一次性memcpy高。针对小张量，直接拷过去往往更省心。
5. 实战：给训练任务做一次三角审计
5.1 三件套工具：nvidia-smi、nsys、ncu
排查性能问题，我固定用三个工具，分工明确：
nvidia-smi：板卡级观测。看利用率、显存占用、功耗、温度。缺点是它只反映瞬时平均，小kernel之间的空隙会被抹平，所以“利用率低”不能直接说明算力有问题；nsys profile（Nsight Systems）：全局时间线。能直接看出哪段时间在算、哪段时间在memcpy、哪段时间CPU和GPU两头都在等。强烈建议先跑它，因为它能告诉你“问题在哪个环节”；ncu（Nsight Compute）：单kernel显微镜。ncu --set roofline能给出操作强度、SM利用率、内存吞吐，精准告诉你这个kernel是算力受限还是带宽受限。
常用命令就很直白：
nvidia-smi nsys profile -o myapp python train.py ncu --set roofline --target-processes all python train.py有的云GPU租用平台不让开profile，那也没关系，nvidia-smi加时间戳采样也能看出个大概，只是精度差点。
5.2 算账流程：判断负载卡在哪一角
别急着优化，先记账。我一般按四步走：
- 列出一次训练step（或一次推理请求）必须搬运的最小数据量：权重读取、梯度同步、数据批次、中间结果写回；
- 分别除以对应的带宽，得到理论耗时下限；
- 跑一次
nsys，看实际耗时分布； - 把实际分布和理论下限比对，短板就自己现形了。
给出一个速查表，方便对号入座：
| 现象 | 大概率瓶颈 | 手段 |
|---|---|---|
| GPU利用率高但全局吞吐上不去 | 搬运间隙、kernel启动排队 | 异步流水、CUDA Graph、融合 |
| 利用率低且显存带宽接近打满 | 带宽受限 | 减少访存、换高带宽卡 |
| 利用率低且带宽也没打满 | 算子太小、等待依赖 | 融合、重排loop、增大batch |
| memcpy占比超过30% | 搬运账太高 | pin_memory、缩短H2D链路、缓存 |
我自己常用的判断很简单：nvidia-smi里GPU利用率虽然高，但nsys里的kernel之间有大片空白，那就是“搬运/启动拖后腿”；如果利用率低但内存带宽接近上限，那就是“带宽天花板”。
5.3 三个立刻见效的优化动作
针对小模型微调这类场景，我实测下面三个动作最划算：
- 融合算子：用
torch.compile把逐元素操作合到一起，能明显减少中间张量读写； - 开启pin_memory并做异步搬运：把数据加载放到独立进程，用CUDA stream提前搬到显存，让GPU在算上一批的时候，下一批已经在路上；
- 用CUDA Graph收拢小kernel：固定shape的训练step可以整段捕获成图，之后每次重复执行，启动开销几乎消失。
代码层面大概长这个样子：
train_loader = DataLoader(ds, pin_memory=True, num_workers=4) s = torch.cuda.Stream() for epoch in range(epochs): for x, y in train_loader: # 让搬运流等当前计算流让出，再异步拷贝 s.wait_stream(torch.cuda.current_stream()) with torch.cuda.stream(s): x = x.cuda(non_blocking=True) y = y.cuda(non_blocking=True) torch.cuda.current_stream().wait_stream(s) loss = model(x, y) loss.backward() optimizer.step()注意：这些动作只是改变搬运节奏，不改变算法。真正更进一步的加速，还要靠FlashAttention这类访存优化算子，以及把整个计算图交给框架级编译器。但把三角账里的搬运账单先降下来，永远是性价比最高的第一步。
6. 常见问题速查：热词背后的事故现场
6.1 核显和独显并存：用错卡是很多诡异问题的源头
很多笔记本的设备管理器里同时挂着“Intel UHD Graphics”和“NVIDIA GeForce RTX 4060 Laptop GPU”。如果你跑AI程序发现GPU没用上，先别怪驱动，多半是应用默认落到了核显上。
排查顺序是这样：
- 终端执行
nvidia-smi，确认系统里能看到的独显和驱动版本； - 代码里设置
CUDA_VISIBLE_DEVICES=0，显式指定独显； - Windows的“设置→系统→屏幕→显卡”里，为Python、游戏或特定软件指定“高性能NVIDIA处理器”；
- NVIDIA控制面板→管理3D设置，把全局或特定程序的GPU选到独显。
有些软件即使指定了还是会出问题，比如Adobe Camera Raw 18.6里GPU选项灰掉、某些游戏报“unsupported GPU”，一般就是驱动版本太旧或者软件对新架构支持不全。先把驱动升级到最新，再把程序强制切到独显，八成能解决。
6.2 新卡被报unsupported GPU：游戏和软件排查顺序
热词里有条“天国拯救2 unsupported gpu”，其实这类报错在AI软件里一样常见，只是表现形式不同。
游戏判断GPU支不支持，通常看两件事：一是驱动版本，二是显卡是否具备它要求的功能特性（比如DX12 Ultimate、特定Vulkan版本、硬件光追）。老显卡或老驱动被报“不支持”太常见了；新显卡被报“不支持”，往往也是驱动太老，软件白名单还认不出新架构。
AI软件类似。比如ComfyUI装插件报冲突、Camera Raw勾不了GPU，根子上多半是驱动、CUDA runtime、三方库版本互相错配。我的经验是先把环境重建一遍：装最新厂商驱动→装支持当前GPU的PyTorch/CUDA版本→再加插件。别在一堆旧版本上打补丁，越打越乱。
6.3 云GPU和租卡怎么选：别被TFLOPS带偏
现在租卡渠道很多，常见的有按卡计费的云GPU平台、算力租赁站。挑卡时，请把规格表按四个列看，而不是只盯浮点算力：
- 显存容量：能不能放下模型和KV cache；
- 显存带宽：决定推理token速度和一部分训练效率；
- 卡间互联：多卡训练必须看有没有NVLink，别组廉价PCIe集群；
- 价格：把前两者折算成“每GB/s带宽多少钱”再比。
开发调试阶段，显存够用就行；正式训练优先看带宽和互联；部署推理服务直接拿带宽除模型体积来估并发。至于“用个人电脑共享算力出租”这种玩法，输出带宽和延迟是硬伤，跑推理服务的网络瓶颈会远大于本地算力，当个玩具可以，当生产主力不划算。
集群层面还有一个常被忽略的细节：共享带宽必须用调度机制管起来。GPU虚拟化和配额管理工具（比如HAMI这类开源方案）存在的意义，就是让多个任务共享一张卡时，谁也别把传送带独占，谁也别把带宽吃干榨净。没有这层调度，一个任务就能拖垮整机箱的邻居。
6.4 个人习惯：固定先做三角审计，再动手优化
最后分享一个我自己的固定动作。
每拿到一个新任务，我不看模型代码，先跑三行命令：nvidia-smi看全局、nsys profile看时间线、ncu --set roofline看关键kernel，然后把算力、带宽、搬运三笔账填到一张表里。多数时候，问题在填表的过程中就自己暴露了。
还有一个装机后必做的检查：
python -c "import torch; print(torch.cuda.get_device_capability(), torch.version.cuda)"这条命令能一次性确认三件事：GPU是否被正确识别、capability是哪个代际、CUDA runtime版本是否匹配。很多莫名其妙的“程序跑不起来”，检查完这一步就省掉大半天排查时间。
我自己踩过太多次“换卡解决一切”的坑。后来发现，真正值钱的不是换更贵的卡，而是先把钱花在刀刃上：算力不够就换算力，带宽不够就换显存带宽更高的型号或做量化，搬运太贵就改流水线和融合。以后不管是调优还是采购，先让这三笔账上桌，很多决定会变得异常清晰。
