跳转至
计算机基础手册 体系结构：从 CPU 到 GPU

GPU 内存系统：从寄存器到 HBM

decode 每生成一个 token，都要把权重和 KV Cache 从显存里读一遍，所以推理的速度首先是显存带宽的速度。这一章把 GPU 的内存系统从下往上走一遍：HBM 的带宽是怎么来的，为什么它比 GDDR 宽得多；显存容量和带宽之比给 decode 定下了什么上限；各级存储的容量、带宽、延迟差多少；为了跑满带宽要有多少数据同时在路上，这又怎样逼出了 cp.async 和 TMA；最后是 GPU 自己的虚拟内存和 CPU、GPU 之间的内存通路。

自测：能答上来就可以跳过本章

  1. H100 SXM 的 3.35 TB/s 显存带宽可以拆成哪两个数相乘？HBM 为什么能做到 GDDR 做不到的带宽？
  2. 显存装满时，把整块显存读一遍要多久？这对 decode 意味着什么？B300 的显存比 B200 大了一半、带宽一样，意味着什么？
  3. 一张 H100 上，片上存储（寄存器、共享内存、L2）加起来有多大？和显存差多少倍？
  4. 要跑满 H100 的显存带宽，同时在途的数据要有多少？分到每个 SM、每个线程是多少？为什么说这逼出了 TMA？
  5. 访存的最小粒度是多少字节？DRAM 内部为什么怕随机的小块访问？
自测参考答案（先自己答，再展开对照）
  1. 总线宽度 × 每个引脚的速率：5 个 HBM3 堆叠、每个 1024 位，共 5120 位，每个引脚约 5.2 Gb/s，5120 × 5.23 / 8 ≈ 3350 GB/s。HBM 把 DRAM 芯片垂直堆叠，和 GPU 一起放在硅中介层上，连线又短又密，每个堆叠就有 1024 位宽的接口；GDDR 芯片焊在电路板上，走线长，一块卡的总线只有 384～512 位，只能靠提高每个引脚的速率（20 多 Gb/s）来弥补。
  2. 容量 ÷ 带宽：H100 80 GB / 3.35 TB/s ≈ 24 ms。decode 的每一步要读全部权重和这一批请求的全部 KV，显存装满时，一步至少要这么久，每个请求每秒最多生成约 40 个 token。B300 的 288 GB 能多放一半的 KV，批次更大、吞吐更高，但装满时读一遍要 36 ms，每个请求反而更慢：容量换来的是吞吐，带宽决定的是速度。
  3. 寄存器 132 × 256 KB ≈ 33 MB，L1 / 共享内存同样约 33 MB，L2 50 MB，加起来 100 多 MB；显存 80 GB，差了约 700 倍。权重放不进片上，每一步 decode 都要从显存把它流一遍。
  4. 带宽 × 延迟：按 600 ns 的访存延迟估算，3.35 TB/s × 600 ns ≈ 2 MB，每个 SM 约 15 KB，每 SM 2048 个线程全满时每个线程要有约 7.4 字节在路上，B200 上约 16 字节。普通的读取要占用寄存器来接收数据，线程数又有上限，靠"每个线程多读几个"已经很吃力；cp.async 让数据绕过寄存器直接进共享内存，TMA 更进一步，一个线程一条指令就能发起一整块（几 KB 到几十 KB）的搬运，在途的数据量不再受线程数和寄存器限制。
  5. 最小粒度是 32 字节的扇区（L2 的缓存行 128 字节，分成 4 个扇区）。DRAM 按"行"组织，读一行要先把它激活到行缓冲区里，同一行里连续的数据很快，换一行要先关掉当前行、再激活新行，要额外的几十纳秒。随机的小块访问几乎每次都要换行，有效带宽大打折扣。

HBM：把内存堆在 GPU 旁边

显存带宽 = 总线宽度 × 每个引脚的速率。两种显存走了两条路：

  • GDDR（消费级卡、L40S 等）：一颗颗显存芯片焊在电路板上，围着 GPU 一圈，每颗 32 位。走线长，总线做不宽（384～512 位），只能把每个引脚的速率拉到 20 多 Gb/s；
  • HBM（数据中心卡）：把 8～12 层 DRAM 芯片垂直堆叠起来，层与层之间用硅通孔连接，每个堆叠有 1024 位的接口；堆叠和 GPU 一起放在一块硅中介层上（台积电的 CoWoS 封装），连线又短又密。几个堆叠就是几千位的总线，每个引脚的速率反而不用那么高，每传一位的能耗也更低。

用规格反推每个引脚的速率，再算一个推理里很有用的数——把整块显存读一遍要多久：

hbm.py
# 显存带宽 = 总线宽度 × 每个引脚的速率。用规格反推每个引脚的速率，再算"把显存读一遍"要多久
gpus = [
    # 名称, 显存类型, 总线宽度（位）, 容量 GB, 规格带宽 GB/s
    ("A100 80GB", "HBM2e x5", 5 * 1024, 80, 2039),
    ("H100 SXM", "HBM3 x5", 5 * 1024, 80, 3350),
    ("H200", "HBM3e x6", 6 * 1024, 141, 4800),
    ("B200", "HBM3e x8", 8 * 1024, 180, 8000),
    ("B300", "HBM3e x8", 8 * 1024, 288, 8000),
    ("MI300X", "HBM3 x8", 8 * 1024, 192, 5300),
    ("RTX 4090", "GDDR6X", 384, 24, 1008),
    ("RTX 5090", "GDDR7", 512, 32, 1792),
]
print("GPU        显存        总线宽度  每引脚速率    带宽    读满一遍")
for name, kind, bits, cap, bw in gpus:
    pin = bw * 8 / bits                              # GB/s × 8 / 位宽 = 每个引脚每秒多少 Gb
    print(f"{name:10s} {kind:10s} {bits:7d} 位 {pin:6.2f} Gb/s {bw / 1000:6.2f} TB/s {cap / bw * 1000:6.0f} ms")
输出
GPU        显存        总线宽度  每引脚速率    带宽    读满一遍
A100 80GB  HBM2e x5      5120 位   3.19 Gb/s   2.04 TB/s     39 ms
H100 SXM   HBM3 x5       5120 位   5.23 Gb/s   3.35 TB/s     24 ms
H200       HBM3e x6      6144 位   6.25 Gb/s   4.80 TB/s     29 ms
B200       HBM3e x8      8192 位   7.81 Gb/s   8.00 TB/s     22 ms
B300       HBM3e x8      8192 位   7.81 Gb/s   8.00 TB/s     36 ms
MI300X     HBM3 x8       8192 位   5.18 Gb/s   5.30 TB/s     36 ms
RTX 4090   GDDR6X         384 位  21.00 Gb/s   1.01 TB/s     24 ms
RTX 5090   GDDR7          512 位  28.00 Gb/s   1.79 TB/s     18 ms

（H200、B200 的可用容量比堆叠的物理容量略少，表里用的是可用容量。）

"读满一遍"的时间是 decode 的一条硬边界。decode 每一步都要读全部权重，以及这一批请求的全部 KV Cache。显存越满（批次越大、上下文越长），每一步读的就越多；显存装满时，一步至少要读满一遍的时间，每个请求每秒最多生成 1 / 24 ms ≈ 40 个 token（H100）。几代数据中心卡的这个数都在 20～40 ms 之间，这是硬件设计时就在容量和带宽之间做的平衡。

B200 到 B300 是个好例子：堆叠从 8 层换成 12 层，容量从 180 GB 涨到 288 GB，带宽还是 8 TB/s。多出来的容量可以放更多请求的 KV，批次更大，吞吐更高；但把显存用满时，读一遍要 36 ms，每个请求的生成速度反而下降。容量换吞吐，带宽定速度，部署时按延迟要求决定"用满"到什么程度（见推理手册的压测、SLO 与容量规划）。

HBM 这几年的演进，就是堆叠层数、每个引脚的速率和堆叠个数一起往上涨：HBM2e → HBM3 → HBM3e，下一代 HBM4 把每个堆叠的接口从 1024 位加宽到 2048 位。

DRAM 内部：为什么怕随机访问

每一片 DRAM 分成很多个 bank，每个 bank 是一个由行和列组成的阵列。读数据要三步：激活（把一整行，通常 1 KB 左右，读进这个 bank 的行缓冲区）、读列（从行缓冲区里取出要的那一段）、预充电（关掉这一行，准备激活下一行）。同一行里接着读，只需要第二步；换到另一行，三步都要走一遍，多花几十纳秒。

所以 DRAM 喜欢大块、连续的访问：内存控制器会把请求重新排序，尽量把落在同一行的请求放在一起；GPU 以 32 字节的扇区为最小单位访问显存，一个 warp 的 32 个线程读连续的地址时，只要 4 个扇区（128 字节），这就是合并访问。随机的小块访问（比如按 token 零散地读 KV、查一张巨大的嵌入表）既浪费扇区里没用到的字节，又让 DRAM 频繁换行，有效带宽可能只剩峰值的一小部分。PagedAttention 的块不能太小，这是原因之一。

各级存储：容量、带宽与延迟

寄存器 每个 SM 256 KB 每个线程私有，约 1 个周期 共享内存 / L1 每个 SM 最多 228 KB 同一个 block 共享，二三十个周期 L2 缓存 50 MB 所有 SM 共享 HBM（显存） 80 GB，3.35 TB/s 几百个周期 越往上越快、越小
图：GPU 的存储层次——寄存器、共享内存、L2、HBM

从 H100 看整个层次：

层级 容量（H100 SXM 整卡） 带宽 延迟量级 由谁管理
寄存器 每 SM 256 KB，共约 33 MB 最高 1 个周期级 编译器
L1 / 共享内存 每 SM 256 KB，共约 33 MB 每 SM 每周期 128 字节，全卡约 30 TB/s 几十个周期 共享内存由程序员管，L1 由硬件管
L2 50 MB，分成两个分区 显存带宽的几倍 约 200 个周期，远端分区更慢 硬件；可以划出一部分做持久化
HBM 显存 80 GB 3.35 TB/s 约 400～800 个周期 cudaMalloc、PyTorch 的缓存分配器
主机内存（PCIe 5.0 x16） 几百 GB 到几 TB 单向约 64 GB/s 微秒级 锁页内存 + DMA
其他 GPU（NVLink 4） — 单向 450 GB/s 微秒级 点对点访问、NCCL

几个比例决定了推理 kernel 的写法：

  • 显存比片上存储大约 700 倍。一个 8B 模型的 BF16 权重就有 16 GB，片上放不下，decode 的每一步都要把权重从显存流一遍。每个权重读进来之后，要在片上被尽量多地用几次：批次里有 B 个请求，每个权重就被用 B 次，算术强度随 B 线性增长——这就是批处理对 decode 如此重要的原因，也是屋脊点的物理含义。
  • 共享内存的带宽约是显存的 10 倍，寄存器更快。所以高性能 kernel 的套路永远是：从显存读一次，放进共享内存和寄存器，在片上反复使用。
  • L2 装不下权重，但装得下很多别的东西：一个 decode 步里各层的激活、小的查找表、正在被多个 block 同时读的 KV 块。A100 开始可以用"访问策略窗口"（cudaAccessPolicyWindow）把一段数据标成"持久化"，让它优先留在 L2 里。L2 分成两个分区，每个 SM 就近访问一个分区，访问另一个分区要经过交叉开关，延迟更高；数据在 L2 里的位置，普通程序无法控制。
  • 离开这张卡，带宽就掉一个数量级：NVLink 比显存慢约 7 倍，PCIe 又比 NVLink 慢约 7 倍。KV 卸载到主机内存、从主机加载权重，都要按这个带宽来估算。

在途的数据：Little 定律与 TMA

显存带宽要"跑满"，需要足够多的访存请求同时在路上。按 Little 定律，在途的数据量 = 带宽 × 延迟。访存延迟按 600 ns 估算：

拨一拨带宽和延迟：

in_flight.py
# Little 定律：要维持带宽 B，在途（已发出、还没返回）的数据量 = B × 延迟。访存延迟按 600 ns 估算（量级）
LATENCY = 600e-9
gpus = [("A100", 2039e9, 108), ("H100 SXM", 3350e9, 132), ("B200", 8000e9, 148)]
print("GPU        整卡在途   每个 SM   每个线程（每 SM 2048 个线程全满）")
for name, bw, sms in gpus:
    total = bw * LATENCY
    per_sm = total / sms
    print(f"{name:9s} {total / 1e6:6.2f} MB {per_sm / 1024:7.1f} KB {per_sm / 2048:9.1f} 字节")
输出
GPU        整卡在途   每个 SM   每个线程（每 SM 2048 个线程全满）
A100        1.22 MB    11.1 KB       5.5 字节
H100 SXM    2.01 MB    14.9 KB       7.4 字节
B200        4.80 MB    31.7 KB      15.8 字节

A100 上，每个 SM 要保持约 11 KB 的数据在路上，一个 SM 驻留 2048 个线程时，每个线程约 5.5 字节，多于一个 float，所以每个线程只读一个 float 的 kernel 在 A100 上就跑不满带宽（CUDA 手册的延迟掩盖一节算过这笔账）。到了 B200，带宽涨了近 4 倍，SM 数只多了 40%，每个线程要有约 16 字节在路上，相当于占用率拉满、每个线程还得时刻挂着一个 float4 的读。

普通的读取还有一个隐藏的成本：数据要落进寄存器，在途的每个字节都占着一个寄存器，而寄存器还要用来计算。硬件的应对是让数据绕开寄存器：

  • Ampere 的 cp.async：从显存直接异步拷贝到共享内存，不经过寄存器，发出之后线程可以去做别的事；
  • Hopper 的 TMA：一个线程发一条指令，搬一整块多维的数据（比如 128 × 64 的 BF16 块，16 KB），地址计算、越界处理都由硬件完成。两三条这样的指令就够一个 SM 维持几十 KB 的在途数据，完全不受线程数和寄存器的限制。

这就是为什么 Hopper、Blackwell 上的高性能 kernel 都变成了"一两个 warp 专门用 TMA 搬数据、其他 warp 专门计算"的结构（见 CUDA 手册的异步拷贝与 TMA）。带宽涨得比 SM 数快，是这些新硬件单元出现的根本原因。

GPU 的虚拟内存

GPU 也有自己的 MMU、页表和 TLB。每个 CUDA 上下文有自己的虚拟地址空间，cudaMalloc 返回的是虚拟地址，GPU 访存时同样要做地址翻译。显存一般按 2 MB 的大页映射，TLB 能覆盖的范围比 CPU 上 4 KB 的页大得多；但访问跨度大到几十 GB 的随机访问（比如在巨大的 KV 池里零散地读），TLB 缺失同样会出现，拖慢访存。

和 CPU 一样，虚拟内存让"分配地址"和"分配物理内存"可以分开，这在推理系统里有直接的用处：CUDA 的虚拟内存管理接口（cuMemAddressReserve、cuMemCreate、cuMemMap）可以先保留一大段连续的虚拟地址，再按需以 2 MB 为粒度映射物理显存。这和 PagedAttention 是两层不同的东西：PagedAttention 是软件维护的块表，注意力 kernel 自己查表找 KV 块；虚拟内存管理接口改的是硬件页表，kernel 看到的就是连续的地址。两者的对比见虚拟内存、页表与大页。

统一内存（cudaMallocManaged）让 CPU 和 GPU 共用一个地址空间，访问到不在本地的页时触发缺页，由驱动把页迁移过来。它写起来方便，但缺页和迁移都很慢，推理系统的热路径上一般不用。

CPU 与 GPU 之间的通路

PCIe 是 CPU 和 GPU 之间最常见的通路：PCIe 5.0 x16 单向约 64 GB/s，比显存慢 50 倍。这决定了 KV Cache 卸载到主机内存能有多快、冷启动时加载权重要多久（估算见锁页内存、DMA 与 NUMA）。

NVIDIA 的 Grace Hopper（GH200）和 Grace Blackwell（GB200）把 CPU 和 GPU 用 NVLink-C2C 连起来，带宽 900 GB/s（双向合计），是 PCIe 5.0 的 7 倍左右，并且硬件保证 CPU 内存和显存之间的缓存一致：GPU 可以直接、高效地读写 CPU 那边几百 GB 的 LPDDR5X 内存。对推理来说，这相当于多了一层又大又不太慢的"二级显存"，KV Cache 卸载、超长上下文、大模型的部分权重都可以放在那里。

怎么讲清楚

讲显存与带宽：带宽 = 总线宽度 × 引脚速率，HBM 靠堆叠和硅中介层做出每堆 1024 位的宽接口，H100 是 5 堆、5120 位、每引脚约 5.2 Gb/s，合 3.35 TB/s。容量 ÷ 带宽是 decode 的硬边界：H100 读满一遍约 24 ms，显存越满、每一步越慢；B300 容量涨一半、带宽不变，是用速度换吞吐。层次上，寄存器、共享内存、L2 加起来 100 多 MB，显存 80 GB，差约 700 倍，所以 kernel 要从显存读一次、在片上反复用，decode 靠批处理提高权重的复用。跑满带宽要满足 Little 定律：在途数据 = 带宽 × 延迟，H100 约 2 MB、每 SM 15 KB，B200 翻倍；带宽涨得比 SM 数快，寄存器接不住，于是有了绕开寄存器的 cp.async 和一条指令搬一整块的 TMA。DRAM 按行组织，怕随机小块访问，GPU 以 32 字节扇区为粒度，要合并访问。

练习

1. 反推与估算。 H20 的显存是 96 GB、带宽 4.0 TB/s，用的是 HBM3。如果每个堆叠 1024 位、每个引脚约 5.2 Gb/s，它大概有几个堆叠？读满一遍要多久？在它上面部署一个 BF16 权重 16 GB 的模型，其余显存全部用来放 KV，显存用满时一个 decode 步至少要多久？

参考答案

4000 GB/s × 8 / 5.2 Gb/s ≈ 6150 位，约 6 个堆叠（6 × 16 GB = 96 GB，容量也对得上）。读满一遍 96 / 4000 ≈ 24 ms。显存用满时，一步要读 16 GB 权重加约 80 GB 的 KV（减掉激活和运行时的少量占用），接近读满一遍，也就是 20 多毫秒。这一步里 batch 有多大，这 20 多毫秒就摊给多少个 token，所以吞吐可以很高，但每个请求每秒最多 40 个 token 左右。

2. 片上存储能帮多少。 H100 的 L2 有 50 MB。一个 decode 步里，一个 8B 模型每层的权重有几百 MB（16 GB 分给 32 层）。L2 能不能让下一步 decode 少读显存？为什么？L2 在 decode 里能帮上什么忙？

参考答案

不能。每一步要读 16 GB 权重，L2 只有 50 MB，按顺序流过一遍之后，L2 里留下的只是最后读到的一点点，下一步从第一层开始读时早就被替换掉了，下一步照样要从显存全部读一遍。L2 能帮忙的是一步之内的复用：同一块权重被多个 block 同时读（比如矩阵乘按 N 方向切给多个 block，它们读的是同一块输入激活），以及同一个 KV 块被多个查询头读（GQA 里一组查询头共用一个 KV 头），这些重复读取可以在 L2 里命中，不必每次都去显存。

3. 在途的数据。 一个 H100 上的逐元素 kernel，每个 SM 驻留 1024 个线程（占用率 50%），每个线程每次读一个 BF16。按本章的估算（每 SM 需要约 15 KB 在途），它能跑到峰值带宽的多少？改成每个线程一次读 16 字节（8 个 BF16）呢？

参考答案

每 SM 在途约 1024 × 2 = 2 KB，是所需 15 KB 的约 13%，带宽最多跑到峰值的 13% 左右（Little 定律：带宽 = 在途数据 / 延迟）。每个线程一次读 16 字节时，在途 16 KB，超过了 15 KB，理论上可以跑满（实际能到峰值的八九成）。这就是逐元素 kernel 要"向量化读取"的定量理由；也说明占用率只有 50% 并不妨碍跑满带宽，只要每个线程的在途数据够多。

4. 选通路。 一个请求的 KV Cache 有 4 GB，需要从别处搬进一张 H100 的显存。分别估算它从本机主机内存（PCIe 5.0 x16）、从 GH200 上 CPU 那侧的内存（NVLink-C2C）、从同机另一张 GPU（NVLink 4）搬过来的时间。

参考答案

按单向带宽估算：PCIe 5.0 x16 约 64 GB/s（实际约 50 GB/s），4 GB 要 60～80 ms；NVLink-C2C 900 GB/s 是双向合计，单向约 450 GB/s，约 9 ms；NVLink 4 单向 450 GB/s，也约 9 ms。同一个请求的 KV，放在 PCIe 那头和放在 NVLink 那头，恢复时间差了一个数量级，这决定了 KV 卸载值不值得做：重算这个请求的 prefill 要多久，就是卸载能接受的上限。

小结

  • 带宽 = 总线宽度 × 引脚速率：HBM 靠堆叠和硅中介层做出每堆 1024 位的宽接口，GDDR 靠高引脚速率；HBM4 接口加宽到 2048 位。
  • 容量 ÷ 带宽是 decode 的硬边界（H100 约 24 ms）；容量换吞吐，带宽定速度。
  • DRAM 按行组织，怕随机小块访问；GPU 以 32 字节扇区为粒度，要合并访问。
  • 片上存储约 100 MB，显存是它的 700 倍：从显存读一次、在片上反复用；decode 靠批处理提高权重复用。
  • 在途数据 = 带宽 × 延迟，每 SM 需要十几到几十 KB；带宽涨得比 SM 数快，于是有了 cp.async 和 TMA。
  • GPU 有自己的页表和 TLB，按 2 MB 大页映射；虚拟内存管理接口改硬件页表，PagedAttention 是软件块表。PCIe 比显存慢约 50 倍，NVLink-C2C 让 CPU 内存成为又大又不太慢的二级存储。

发现错误或有更好的写法？在 GitHub 上编辑此页提交问题