内存层次：global / L2 / shared / register 的带宽与延迟

上一篇 讲了线程怎么执行。这一篇讲数据住在哪。GPU 的算力很强，但喂数据的速度跟不上——大多数算子的真实瓶颈是访存，而访存有快有慢，差几个数量级。理解存储金字塔，才能回答”数据该放哪一级”这个主导性能的问题。

一、存储金字塔：五级

从快到慢、从小到大：

延迟数字是 CUDA 架构的常见量级（来自 NVIDIA 编程指南与公开微基准），不同代际略有差异；本文用实测重点说明带宽的层级差距。

flowchart TD
  R["寄存器<br/>每线程私有 · 最快"] --> S["shared / L1<br/>每 block 共享 · 100KB/SM"]
  S --> L2["L2 cache<br/>全 GPU · 3MB"]
  L2 --> G["global memory<br/>8GB GDDR6 · 448 GB/s 理论"]
  G --> H["host memory<br/>过 PCIe"]

优化访存的核心思路就一句话：把要反复使用的数据尽量留在更高的层级。GEMM 的 tiling、FlashAttention 的分块，本质都是把数据从 global 搬到 shared / 寄存器后反复复用。

二、寄存器与 shared memory：片上存储

• 寄存器是线程私有的最快存储，编译器自动分配。kernel 里的局部标量变量通常就在寄存器里。寄存器用量直接影响 occupancy（第 02 篇 已提到，每 SM 65536 个），用太多会溢出到 local memory（实际在显存里，很慢）。

• shared memory 是同一个 block 内所有线程共享的可编程片上存储，用 __shared__ 声明。它和 L1 cache 共用同一块物理 SRAM——在 RTX 3060 Ti 上每 SM 共 128 KB，可配置最多 100 KB 给 shared。shared memory 是手动管理的”暂存区”：你显式地把 global 数据搬进来、__syncthreads() 同步、然后让 block 内线程反复读。它的带宽远高于 global，延迟也低得多，但容量小、生命周期只在 block 内。

shared memory 还有一个性能陷阱叫 bank conflict（访问模式不当会串行化），这是 访存优化篇 的重点。

三、L2 与 global：片外存储，实测带宽差近一个数量级

global memory 就是显卡上的 GDDR6 显存，所有 block 可见，也是 host 数据拷入拷出的落点。它容量大（8 GB）但带宽相对低、延迟高。L2 cache 夹在 SM 和显存之间，全 GPU 共享 3 MB，自动缓存 global 的访问。

L2 和 global 的带宽差距有多大？用一个重复读实验测量：每个 block 把一个工作集反复读若干遍累加。工作集小到能放进 3 MB L2 时，第一遍之后都是 L2 命中；工作集远大于 L2 时，几乎每次都打到 DRAM。RTX 3060 Ti 实测（grid-stride 读，256 线程/block，304 个 block，CUDA event 取中位数）：

L2 命中时约 3.4 TB/s，是 DRAM（约 400 GB/s）的 8 倍以上。注意 4 MB 略超 3 MB L2 仍然很高，是部分驻留的结果；一旦工作集明显超过 L2（16 MB 起），带宽就回落到 DRAM 水平，并逼近该卡理论带宽 448 GB/s。

这组数字解释了一个常见现象：同一个 kernel，数据规模跨过 L2 容量时性能会突然掉一截。算子调优时，估算工作集和 L2 的关系，往往比调线程数更有用。

DRAM 这一档（约 400 GB/s）本身也不是随便就能跑满的，要靠合并访问。前一篇提过的向量加法实测 413 GB/s，已经接近理论上限——这是访问模式良好的结果。访问模式不当会让有效带宽腰斩，下一篇专门讲。

四、把层次用起来：复用率决定上限

一个算子能跑多快，先看它对每个从 global 读进来的字节”用了几次”。这个比值越高，越能把工作压在片上存储，越不受 global 带宽限制。

• 向量加法：每个 float 读一次用一次，复用率约等于 1，必然受 global 带宽限制，无优化空间。
• 朴素 GEMM：每个元素被反复从 global 读取，复用率本可以很高，但朴素写法没把数据留在片上，于是反复打 global，浪费带宽。
• tiled GEMM：把子块搬进 shared，让 block 内线程反复复用，复用率提升到块大小量级，于是从受 global 带宽限制转为受算力限制。

这条主线贯穿后面的 GEMM 篇 和 FlashAttention 篇：优化访存密集算子，本质是改造数据流，让复用发生在更高的存储层级。怎么判断一个算子是受算力还是受带宽限制，是 Roofline 篇 的核心方法。

五、小结与下一步

• GPU 存储分五级，带宽与延迟逐级递减/递增；优化访存就是把复用数据留在更高层级。
• 寄存器和 shared memory 是片上快存储，shared 需手动管理且容量小（本卡 100 KB/SM）。
• L2 命中与 DRAM 的实测带宽差 8 倍以上（约 3.4 TB/s vs 约 400 GB/s），工作集是否落进 L2 会让性能突变。
• 算子的复用率决定它能否摆脱 global 带宽限制。

下一篇先把抽象落地——写第一个 CUDA kernel，把索引、同步和启动配置讲清楚，为后面的访存与计算优化打基础。