状态传导图
HBM stack 与 GPU die 分离;在 H100 等产品中二者共封装。图中容量与距离只表达相对概念。
观察点 01
完整权重和 KV Cache 主要住在显存大仓库
模型加载后,权重通常常驻 HBM。它容量大、带宽高,但相对片上存储更远;每生成 token 仍要读取许多层的权重与历史状态。
- 可容纳数据量最大
- 离计算单元距离较远
“模型在 GPU 上”不等于数据已经在 Tensor Core 手边。
容量回答“放不放得下”,带宽回答“每秒能搬多少”,延迟回答“一次访问从发出到数据可用要等多久”。拖动数据层级,观察工作集越靠近计算单元,容量怎样缩小、复用要求怎样升高。
状态传导图
HBM stack 与 GPU die 分离;在 H100 等产品中二者共封装。图中容量与距离只表达相对概念。
观察点 01
模型加载后,权重通常常驻 HBM。它容量大、带宽高,但相对片上存储更远;每生成 token 仍要读取许多层的权重与历史状态。
“模型在 GPU 上”不等于数据已经在 Tensor Core 手边。
状态传导图
L1 cache 与 Shared Memory 逻辑语义不同;后者不是每个字节的必经站。
观察点 02
矩阵太大,kernel 按 tile 发起 global-memory load。Cacheable 请求先查 L1/TEX;L1 miss 或 bypass 才请求 L2,L2 miss 才到 HBM。需要线程块协作时,再显式把当前 tile 放进 Shared Memory 复用。
SRAM 的价值不是装下整层,而是让当前 tile 被高效复用。
状态传导图
REG / SMEM / TMEM 的选择依架构与指令而变;load → compute → store 比例仍决定瓶颈。
观察点 03
传统 MMA kernel 常把操作数与累加值放在寄存器;但 Hopper WGMMA 可让操作数来自 Shared Memory,Blackwell tcgen05 还使用 Tensor Memory(TMEM)保存累加器等数据。寄存器压力仍可能引发 spill。
最快的实现要让数据按目标 ISA 及时到位,不能把一张层级图当成固定指令合同。
状态传导图
连续批处理可让一次权重流式读取服务更多 token,提高复用;它不等于权重只被读取一次。
观察点 04
低 batch Decode 的矩阵往往“瘦”,计算单元难以完全吃满;若每步主要在等 HBM 搬权重与 KV,增加理论算力帮助有限,量化和批处理可能更直接。
Decode 常见而非必然 memory-bound;batch、上下文、量化、缓存与 kernel 都会移动瓶颈。
你现在应该能解释:Decode 常见而非必然 memory-bound;batch、上下文、量化、缓存与 kernel 都会移动瓶颈。
看 GPU 参数时,很多人只盯 TFLOPS。真正做大模型时,经常先撞上的却是三堵墙:放不下、搬不动、等太久。这篇把 HBM、SRAM、CPU 内存、NVLink、PCIe、KV Cache、优化器状态串成一张“数据搬运地图”。
| 容量 | 放不放得下:权重、KV Cache、激活、梯度、优化器状态。 |
| 带宽 | 搬不搬得快:HBM、SRAM、NVLink、PCIe 每秒能搬多少数据。 |
| 延迟 | 等不等得起:一次远距离访问、一次跨机通信、一次 offload 往返要多久。 |
越靠近计算单元的工作集通常越小、访问越快;越远的存储通常容量更大、访问代价更高。这里先画“数据复用视角”,但要注意:register、shared memory、cache 是不同的编程语义;在现代 NVIDIA GPU 上,L1 data cache、texture cache 与 shared memory 又会共享一部分片上 SRAM 资源。
| 层级 | 你可以怎么理解 | 大模型里常见用途 |
|---|---|---|
| Registers | 每线程分配在 SM register file 的操作数与累加空间;溢出会产生 local-memory 流量 | kernel 临时变量、传统 MMA 操作数与累加器 |
| L1 / Shared | 逻辑语义不同:L1 由硬件缓存,shared memory 由线程块显式寻址;Hopper/Blackwell 将它们放在统一片上 SRAM 资源中并支持 carveout | FlashAttention tile、矩阵乘分块、归约、局部缓存 |
| L2 Cache | 芯片级共享缓存;L1 miss/bypass 的 global-memory 请求会到 L2,L2 miss 才访问 HBM | 复用权重块、KV 块、跨 SM 数据 |
| HBM | HBM stack 与 GPU die 分离;在 H100 等产品中与 GPU 芯片共封装。容量远大于片上 SRAM,但带宽更低、访问延迟更高 | 模型权重、激活、KV Cache、梯度、优化器状态 |
| CPU RAM / NVMe | 更远的仓库 | offload、检查点、数据加载;能救容量,但会伤延迟和吞吐 |
tcgen05 还引入 Tensor Memory(TMEM)承载累加器等数据。优化必须以目标架构和具体指令为准。显存讨论里最容易乱的是:容量、带宽、延迟是三件不同的事。它们分别对应“装不装得下”“每秒能搬多少”“一次访问从发出到数据可用要等多久”。应用层 TTFT/TPOT 还会叠加排队、调度与计算,不能直接等同于一次内存访问延迟。
仓库有多大。决定模型权重、KV Cache、batch、上下文长度能不能塞进显存。
货梯每秒能搬多少货。decode 阶段常卡在“每生成一个 token 都要读权重/KV”。
叫一次货要等多久。CPU↔GPU、跨机器通信、随机小访问,常被延迟支配。
| 硬件 / 通道 | 容量或带宽 | 说明 |
|---|---|---|
| H100 SXM | 80GB HBM3,HBM 峰值 3.35TB/s;NVLink 900GB/s | NVIDIA 官方峰值;900GB/s 是每 GPU 的 18 条第四代 NVLink 双向汇总。 |
| H200 SXM | 141GB HBM3e,HBM 峰值 4.8TB/s;NVLink 900GB/s | Hopper 的容量与 HBM 带宽升级;NVLink 同样是每 GPU 双向汇总峰值。 |
| HGX B200 SXM | 180GB HBM3e,HBM 峰值最高 8TB/s;第五代 NVLink 1.8TB/s | 当前 NVIDIA HGX 架构页口径;1.8TB/s 是每 GPU 双向。具体 Blackwell 产品的容量会随形态而异。 |
| HGX B300 SXM | 288GB HBM3e,HBM 峰值最高 8TB/s;第五代 NVLink 1.8TB/s | 当前 HGX B300 官方架构口径;NVLink 是每 GPU 双向汇总峰值。 |
| PCIe 5.0 x16 | 编码后理论值约 63.0GB/s/方向,约 126.0GB/s 全双工汇总 | 尚未扣 TLP/DLLP、流控、DMA 与软件开销;不是应用可持续吞吐。 |
32 GT/s × 16 lanes × 128/130 ÷ 8 = 63.015 GB/s,这是单方向、只扣 128b/130b 编码后的理论数据率;全双工两方向相加约 126.03GB/s。厂商常把它四舍五入写作 64GB/s/方向或 128GB/s aggregate,不能拿汇总值与单向实测混比。“模型多少 B”只是权重规模。训练和推理还会额外占激活、梯度、优化器状态、KV Cache。先把账拆开,后面的 量化、ZeRO、PagedAttention 才讲得通。
| 模型 | FP16/BF16 权重 | INT8 权重 | INT4 权重 |
|---|---|---|---|
| 7B | 约 14GB | 约 7GB | 约 3.5-4GB |
| 70B | 约 140GB | 约 70GB | 约 35-40GB |
| 671B MoE | 约 1.34TB(若全 BF16 存权重) | 约 671GB | 约 335GB+ |
表中按十进制 GB/TB 粗算;INT4 的“参数数 × 0.5 byte”只是紧凑存储下界,真实格式还可能包含 scale、zero-point、分组元数据、对齐与运行时 buffer。这只是权重,不含 KV Cache 和激活。MoE 每 token 只激活少数专家,省的是计算量;若所有专家常驻显存,权重容量账仍然要付。
B=当前副本上同时驻留的序列数,L=缓存 KV 的层数,S=每条序列已缓存长度,H_kv=KV 头数,D=每头维度,乘 2 是 K 和 V。这是单个模型副本、分片前的逻辑 payload;真实物理占用还受 TP/CP 分片、分页末块、对齐、prefix sharing、滑动/局部注意力层、latent cache 与 KV 量化影响。
B=1,设 32 层、4096 token、8 个 KV 头、每头 128 维、BF16:32×4096×8×128×2×2 = 536,870,912 bytes,即每条序列约 512MiB。若改成 32 个 KV 头的 MHA,其余不变,就是约 2GiB,正好 4 倍。| 项目 | 为什么占显存 | 常见优化 |
|---|---|---|
| 参数 weights | 模型本体 | FSDP / ZeRO、低精度、张量并行 |
| 梯度 gradients | 反向传播要给每个可训练参数存梯度 | ZeRO/FSDP 梯度分片、冻结参数或 PEFT;部分系统也支持低精度梯度 |
| 优化器状态 | AdamW 通常要一阶/二阶矩,混精还可能有 FP32 master weight | ZeRO、8-bit optimizer;也可给适合的参数组选择 Adafactor、Muon 等不同状态结构 |
| 激活 activations | 反向传播要用前向中间结果;序列越长越贵 | 减小 micro-batch / 梯度累积、gradient checkpointing、FlashAttention、sequence parallel |
梯度累积只有在把单次 micro-batch 调小时,才会降低该步保存的激活;它不会自动缩小“每个可训练参数一个梯度”这组张量。一个常见但非通用的混精 Adam 粗算是 BF16 参数 2B + BF16 梯度 2B + FP32 master weight 4B + FP32 一阶/二阶矩 8B = 16 bytes/param;框架可能复用 buffer、改变梯度/状态精度或不保留 master weight。再叠加激活、临时 workspace 与通信 buffer,训练 70B 不能按“权重 140GB”简单估算。
大模型不是一个矩阵乘完就结束。一次训练或推理,会在 HBM、片上 SRAM、多张 GPU、CPU 内存之间来回搬数据。要按真实拓扑分边界:NVLink Domain 可能只是一个 HGX 主机,也可能像 GB200/GB300 NVL72 那样横跨多个 compute tray / node;跨 NVLink Domain 或跨机架才通常进入 NIC 与网络 Fabric。
| 路径 | 典型瓶颈 | 对应技术 |
|---|---|---|
| HBM ↔ SRAM ↔ Tensor Core | HBM 读写太多,Tensor Core 等数据 | FlashAttention、fused kernel、tiling、重计算 |
| GPU ↔ GPU(NVLink Domain 内) | NVLink/NVSwitch 拓扑、带宽与延迟限制扩展效率;Domain 可能跨节点 | 张量并行、专家并行、all-reduce / all-to-all 优化 |
| GPU ↔ CPU RAM | 普通服务器中 PCIe 相比 HBM/片上路径带宽低得多,offload 往返会拉高延迟 | ZeRO-Offload、paged optimizer、CPU KV offload |
| 跨 NVLink Domain / 跨机架 | NIC 与 InfiniBand/Ethernet Fabric 通信遮不住,训练吞吐下降 | 流水并行、通信计算重叠、拓扑感知调度 |
Roofline 模型的思想很简单:如果从所选内存层级每搬 1 byte 能做很多计算,就更可能接近算力上限;如果搬很多数据只做一点计算,就更可能卡在该层级的带宽。Operational intensity 会受缓存复用、数据布局和 kernel 实现影响,并非模型名称对应的固定常数。
图中两个点只表达典型相对位置,不是任何模型或 GPU 的实测。斜线斜率取决于所选内存层级带宽,水平屋顶取决于匹配的数值精度、dense/sparse 口径与硬件峰值。
100/2=50 FLOP/byte。强度 10 FLOP/byte 的 kernel 上限约 20TFLOP/s;强度升到 100 时,公式给出 200TFLOP/s,但最终仍被 100TFLOP/s 的算力屋顶截住。整段 prompt 可并行计算,大矩阵乘较多;但很长的注意力、较小 batch 或低效 kernel 也可能受内存 I/O 限制。FlashAttention 等实现会影响 TTFT。
低 batch 时,每步只生成少量 token,却要读权重和历史 KV,算力可能闲着。batch、量化、上下文长度与 kernel 都可能把瓶颈位置推走。
把这套语言掌握后,很多看似分散的技术会突然连起来:FlashAttention、量化、PagedAttention、ZeRO、MoE,本质都在重新安排数据。
| 技术 | 它在优化什么 | 一句话机制 |
|---|---|---|
| FlashAttention | HBM ↔ 片上 SRAM 的中间矩阵读写 | 用 tiling 与在线 softmax 避免在 HBM 物化完整 N×N 注意力矩阵,同时保持精确注意力。 |
| 量化 | 权重/激活/KV 的容量与带宽 | 更少 bit 可减少存储和流量;端到端加速仍取决于硬件、kernel 与反量化开销。 |
| PagedAttention | KV Cache 碎片和显存利用率 | 借鉴虚拟内存分页,以非连续块管理 KV,减少碎片和冗余复制;仍会有末块及元数据开销。 |
| Continuous Batching | decode 阶段的权重读取摊销 | 动态把活跃请求组成批次,让一次权重流式读取服务更多 token,而非字面上只搬一次权重。 |
| ZeRO / FSDP | 训练时参数、梯度、优化器状态容量 | 把状态切到多卡,不让每张卡都复制全套。 |
| MoE + Expert Parallel | 计算量 vs 通信量的交换 | 每 token 少算专家,但 token 要跨卡路由,all-to-all 成新瓶颈。 |
| Prefill/Decode 分离 | 两类资源特征不同的 workload 混跑 | 把“读题”和“吐字”拆到不同资源池分别优化,但会新增 KV 传输、调度和排队成本。 |
不一定。decode、embedding lookup、小 batch、跨卡通信,都可能让 Tensor Core 吃不饱。
容量大只是装得下。带宽不够、通信慢、kernel 不好,吞吐仍然上不去。
MoE 主要省每 token 计算量。专家权重是否全驻留、是否分布在多卡,决定它到底省不省显存。
要看阶段与并发:长 prefill 可能受注意力计算和 I/O 影响;高并发 decode 常先撞上 KV Cache 容量与带宽。
两者峰值带宽差距明显,但实际影响还取决于拓扑与通信频率。张量并行、专家并行、all-reduce/all-to-all 通常尤其敏感。
CPU/NVMe offload 能救容量,但数据走远路,延迟和吞吐会付代价。
| 显存容量 | 决定能不能装下权重、KV Cache、激活和优化器状态。 |
| 显存带宽 | 决定每秒能从 HBM 搬多少数据;decode 阶段常被它卡住。 |
| 延迟 | 决定每次访问要等多久;小批量、跨设备、随机访问最敏感。 |
| 训练大户 | 权重 + 梯度 + 优化器状态 + 激活。 |
| 推理大户 | 权重 + KV Cache;上下文越长、并发越高,KV 越致命。 |
| 算法工程主线 | 少存、少搬、搬得更顺,最后让 Tensor Core 不空等。 |
核查日期:2026-07-14。厂商 GB/TB 按十进制规格口径;文中 KV 手算用二进制 MiB/GiB。硬件带宽均为理论峰值或官方标称值;NVLink 写每 GPU 双向汇总峰值,PCIe 同时列单方向与全双工汇总,二者都不能直接视为单 peer 或应用实测吞吐。