GPU 内存层级：寄存器、shared memory、L2 到 HBM 的容量带宽延迟金字塔
更新时间：2026-09-25。前置阅读：GPU 架构：warp 调度器。上一篇讲了 GPU 靠成千上万个 warp 互相切换来隐藏停顿，但 warp 切换只能掩盖延迟，不能消灭延迟——数据从哪来、离计算单元有多近，决定了一个 kernel 到底能跑多快。本篇沿着寄存器 → shared memory/L1 → L2 → HBM/GDDR 这条链，把每一层的容量、带宽、延迟和作用域摊开来讲。
本文要回答的问题
- GPU 的内存层级从上到下分哪几层，各自的容量、带宽、延迟是什么数量级？
- 寄存器和 shared memory 的作用域有什么不同，为什么说 shared memory 是「软件可管理」的？
- L2 和全局显存分别为谁服务，HBM 带宽那么高为什么还说它「慢」？
- 为什么 GPU 优化的核心动作是把数据搬进 shared memory 反复用？
- 这套层级和 CPU 的 cache 金字塔看起来相似，本质区别在哪？
一、四层金字塔：越靠近计算单元，越快、越小、作用域越窄
GPU 上的存储按离 SM（Streaming Multiprocessor，流多处理器）计算单元的远近分成四层。从上往下，容量指数级变大，单位延迟指数级变长，数据可见的作用域也从「一个线程」扩大到「整张卡」。
这张图是全篇的骨架。注意三个词：容量、延迟、作用域。寄存器只对所属线程可见，shared memory 只对同一个 block（线程块，CUDA 里称 CTA）内的线程可见，L2 和显存则全卡共享。越往下共享范围越大，但每次访问付出的周期代价越高。下面这张表给出当代数据中心级 GPU（A100/H100 一档）的典型数量级，具体数字随架构变化，量级关系不变。
| 层级 | 典型容量 | 带宽量级 | 延迟（时钟周期） | 作用域 | 谁管理内容 |
|---|---|---|---|---|---|
| 寄存器 | 每线程最多 255 个 32 位寄存器；每 SM 约 256 KB | 最高，单周期可读多个操作数 | 约 1 周期 | 线程私有 | 编译器分配 |
| shared memory / L1 | 每 SM 最多约 100~164 KB（可配置切分） | 每 SM 每周期约百余字节，全卡聚合十余 TB/s | 约 20~30 周期 | block 内共享 | 软件显式搬移 |
| L2 | 40~50 MB | 数 TB/s | 约 200 周期 | 全卡共享 | 硬件自动缓存 |
| HBM/GDDR 显存 | 40~192 GB | 约 2~3.35 TB/s（HBM3 代际继续提高） | 约 400~800 周期 | 全卡共享 + DMA | 软件分配，硬件取数 |
读这张表的正确方式不是记数字，而是算比值：访问一次显存的延迟大约是访问一次寄存器的五百倍，而 shared memory 比显存快约二十倍、容量却小三个数量级。所有 GPU 访存优化，本质上都是在这张表的比值之间做搬运调度。
二、寄存器：线程私有的最高层
寄存器是离计算单元最近的存储，ALU 的操作数直接从寄存器取，单周期可用。每个 SM 配备巨量寄存器（A100 一档约 256 KB，即 65536 个 32 位寄存器），但这些寄存器要分给该 SM 上同时驻留的所有 warp 的所有线程。一个 warp 含 32 个线程，每个线程最多可用 255 个 32 位寄存器。
寄存器是线程私有的：线程 A 看不到线程 B 的寄存器，同 warp 内也不行。线程间交换数据必须往下走 shared memory。寄存器的分配由编译器完成，程序里的局部标量变量、循环索引、中间计算结果通常都落在寄存器里，程序员不需要显式管理，但需要关心用量。
关键约束在于寄存器总量决定 occupancy（占用率，即 SM 上能同时驻留多少 warp）。kernel 里每线程用的寄存器越多，能同时驻留的线程块越少，调度器手里可切换的 warp 就越少，长延迟越藏不住。当寄存器不够用时，编译器会把部分变量溢出（spill）到 local memory——名字叫 local，物理上却在显存上，一次溢出访问就是几百周期。所以寄存器压力过大会以两种方式伤人：occupancy 下降和 spilling，二者常常同时出现。
具体算笔账：SM 上约 65536 个寄存器要同时供给所有驻留线程，若 kernel 每线程用 32 个寄存器，一个 256 线程的 block 占 8192 个，SM 上可同时驻留多个 block；若每线程用到 128 个，可驻留的 block 数立刻缩到四分之一。除了寄存器，shared memory 容量、每 SM 最大 block 数、最大 warp 数也都是硬上限，实际 occupancy 取这些约束的最小值。因此编译产物里的每线程寄存器用量和占用率分析，是调一个 kernel 时首先要看的两份数据。
三、shared memory/L1：block 内共享、软件可管的片上 SRAM
往下一层是 SM 片上的 SRAM，在同一物理存储上切出两种用途：一部分做 L1 data cache，一部分做 shared memory（CUDA 术语；OpenCL 称 local memory），切分比例可由应用配置。A100 上每 SM 这块存储合计 192 KB，shared memory 最多可配到约 164 KB；H100 一代提高到 228 KB。
shared memory 有三个本质特征。
**其一，作用域是 block。**同一个线程块内的所有线程都能读写同一块 shared memory，块外线程不可见。GPU 没有跨 block 直接共享片上数据的机制，block 之间交换数据必须经 L2 或显存。
**其二，内容由软件管理。**它不会自动缓存显存内容，程序必须显式地把数据从显存 load 进来、用完再写回去。这让它更像一块程序员可见的 scratchpad（便笺存储），而不是传统 cache。软件可控意味着延迟可预期、复用可安排，这与 MCU 上 TCM 紧耦合内存的哲学一致，只是共享范围从「一个核」变成了「一个 block 的成百上千线程」。
**其三，内部按 bank 组织。**shared memory 被分成 32 个等宽 bank，每个 bank 每周期可独立服务一个地址，宽度通常 4 字节。当一个 warp 的 32 个线程访问的地址恰好落在不同 bank 时，所有访问一个周期并行完成；若多个线程访问同一 bank 的不同地址，则发生 bank conflict，访问被串行化，最坏情况下吞吐降到 1/32。好消息是访问同一地址会走广播，不算冲突。bank 的划分思路与片上 存储分块与交织完全同源，具体的冲突情形与绕开手法见 shared memory bank conflict。
编程上 shared memory 用 __shared__ 关键字声明，分静态（定长数组）和动态（启动 kernel 时按字节数申请）两种。block 协作把数据加载进来之后，必须调用 __syncthreads() 同步栅栏：它保证块内所有线程都走到同一位置，且此前对 shared memory 的写入对全块可见，否则先跑到后面的线程可能拿着未初始化的旧数据参与计算。同步点也不是越多越好——栅栏让先到的线程空等，切得过碎的同步会直接拉低有效吞吐。
延迟方面，shared memory 访问约 20~30 周期，比 L1 hit 略高或相当，但带宽极高——每周期可向寄存器馈送上百字节，这正是它能喂饱密集计算的原因。
四、L2：全卡共享的最后一道片上防线
所有 SM 共享一块 L2 cache，A100 为 40 MB，H100 为 50 MB。它是片上最后一级缓存，SM 访问显存的所有请求（除显式绕过的流式访问外）都经过它，跨 block 的数据复用、kernel 之间的数据驻留主要靠它承接。
L2 的内容由硬件自动管理，按地址缓存、按替换策略淘汰，软件不能直接往里面放指定数据，只能通过访问模式间接影响命中率。它的延迟约 200 周期，带宽数 TB/s——比显存快、比 shared memory 慢一个量级。
要理解 L2 的价值，得对比作用域：shared memory 只在 block 内共享，而一个 kernel 的 block 之间无法直接通信；相邻 kernel 启动时 shared memory 内容更是早已不保证存在。L2 恰好填补这个空档：一个 block 写出、另一个 block 读入的 tile，或者同一数据被前后两个 kernel 相继使用，只要热度够高、没被挤出，就能在 L2 命中，省掉往返显存的几百周期。现代架构还允许通过访问策略窗口（persisting access policy 之类）提示某类数据优先驻留 L2，但这只是提示，存放决定权仍在硬件。
物理上 L2 被划成若干分片挂在片上网络（NoC）上，就近服务不同区域的 SM，跨分片的请求要经网络转发。40~50 MB 相对显存仍是零头，它只装得下「最近、最热」的那部分数据；一个工作集明显超过 L2 的 kernel 反复重读同一批数据时，命中率会很难看，这类场景只能靠分块把工作集压小，不能指望 L2 自动兜住。
五、HBM/GDDR 全局显存：大、带宽高，但一次访问几百周期
最底层是全局显存，物理上是挂在 GPU 芯片外的 HBM 或 GDDR 颗粒，通过宽总线与封装内互连连接（物理实现见 HBM 存储接口）。数据中心卡用 HBM：A100 80GB 版约 2.0 TB/s，H100 约 3.35 TB/s；消费级显卡用 GDDR，如 RTX 4090 的 24GB GDDR6X 约 1 TB/s。显存技术的代际演进见 GPU 显存 HBM 与 GDDR。
这里有个初学者常绕不过来的矛盾：**2 TB/s 以上的带宽明明是世界上最快的商用存储系统之一，为什么还说显存「慢」？**因为带宽和延迟是两件事。HBM 靠数千根数据线的极宽接口（一次并发搬运数百字节）和极高频率拿到惊人的总带宽，但一次单独访问从发出请求到数据返回，要经过片上 NoC、内存控制器、DRAM 阵列激活和原路返回，物理延迟约 400~800 周期，对单个线程而言是漫长的停顿。
GPU 对这一矛盾的解法分两半：延迟靠 warp 切换隐藏（上一篇 warp 调度器的主题），带宽靠合并访问榨取。一个 warp 32 个线程若访问连续地址，32 次访问会合并成一次宽事务；若地址离散，则退化成多次窄事务，峰值带宽用不满。合并的规则与坑见 内存合并访问。
显存也是 GPU 上唯一的大容量存储：模型权重、帧缓冲、数据集、kernel 代码全部装在这里，它的容量直接决定「多大的问题能一次放下」。HBM 靠堆叠 DRAM 和极宽接口换带宽，GDDR 则在常规封装上拼高频，二者单次访问延迟相近，带宽与价格差距明显。对程序员来说最该记住的是：几百周期的单请求延迟无法靠单线程消化，必须让大量请求同时在途，再用调度和合并把它们变成流水。
六、编程核心：把数据搬进 shared memory 反复用
理解了四层的比值，GPU 优化的主线动作就很直白：让从显存来之不易的每条数据，在片上被复用尽可能多次，再放它走。
典型手法是分块（tiling）。以矩阵乘法 C = A × B 为例，朴素写法中计算一个输出块需要反复按行、按列读 A 和 B，同一元素被不同输出点重复从显存读取几十上百次。分块做法是：整个 block 协作，先把 A、B 当前要用到的子块从显存合并加载到 shared memory，设置同步点保证块内所有线程都拿到数据，然后计算只从 shared memory 取数；一块算完再加载下一块。
一次显存读取在 shared memory 里被复用几十次，意味着顶层那条几百周期的慢路径被摊薄到原来的几十分之一；同时，块内线程协作加载时天然排成连续访问，把显存带宽也用满了。这就是为什么卷积、矩阵乘、模板计算、归约这类高复用度的 kernel 能逼近 roofline 上限，而没有分块的版本常只有峰值的几个百分点。归约（reduction）是另一个典型：朴素写法里参与求和的线程各自直接去显存读输入，总流量随线程数膨胀；先把输入分段搬进 shared memory、在片上做完树形求和，再让每段写回一个部分和，显存流量能压低一两个数量级。复用倍数其实可以定量估算：矩阵分块中加载一块 k 维子块后，A 的列和 B 的行被块内线程在算内积时各读取约 BLOCK 次，复用率就近似分块边长。
选择 tile 尺寸时要同时满足三个约束：tile 放得进每 SM 的 shared memory、访问排布不撞 bank、每线程寄存器用量不溢出，三者相互牵制。还要注意 shared memory 与 L1 共享同一块 SRAM：shared 占多了 L1 就小，普通全局变量缓存和溢出路径都会受影响；而每块申请的 shared memory 过多，又会直接压低可驻留 block 数。所以 tile 大小是在复用率、bank 冲突和 occupancy 之间找平衡，不是一味求大。反过来说，若数据只用一次，搬进 shared memory 只是多一次拷贝，不如直接经 L1 流向寄存器。
七、与 CPU cache 的区别：相似的金字塔，相反的管理权
GPU 层级和 CPU 的 L1/L2/L3/内存金字塔外形相似，常被混为一谈，但二者的设计前提和管理边界有本质区别。
| 对比维度 | CPU cache 层级 | GPU 内存层级 |
|---|---|---|
| 片上高速层管理 | L1/L2/L3 全部硬件自动缓存，软件基本透明 | 寄存器编译器管，shared memory 软件显式管理，只有 L1/L2 硬件自动 |
| 延迟隐藏方式 | 乱序执行、预取、深流水线 | 海量 warp 零成本切换 + 软件分块 |
| 片上容量取向 | 相对小而精，靠命中率服务少数核 | 每 SM 大寄存器堆 + 大 shared memory，面向数千线程 |
| 优化责任划分 | 硬件负责数据放置，程序员写顺序代码即可 | 程序员负责数据复用与搬运，硬件只保证带宽和调度 |
| 一致性 | 硬件维护多核 cache 一致性（MESI 协议） | block 内靠同步指令，跨 block 经 L2/显存，无透明一致性 |
一句话概括差异：CPU 的思路是「硬件替你猜数据放哪，猜对了就快」，cache 组织与替换见 Cache 组织结构；GPU 的思路是「你比硬件更知道数据接下来怎么用，那搬哪、何时复用由你显式安排」。GPU 保留 L1/L2 自动缓存兜住不可预见的访问，同时把带宽收益最大、复用模式最规则的那层（shared memory）开放给软件。这不是偶然——GPU 面对的是数千线程规则访问已知 tile 的场景，软件分块的命中率远高于硬件猜测；代价是性能责任前移到了程序员身上，写 CUDA 本质上是在为这四层金字塔做数据调度。
这种责任划分也解释了一个常见现象：同一份算法，朴素实现和精心分块的实现在 GPU 上能差出十倍甚至几十倍，而在 CPU 上同类代码的差距通常小得多——CPU 的硬件 cache 会替所有人兜底，GPU 上程序是否贴合层级则几乎一比一反映到耗时上。新手常有的失落感正源于此：结果正确、线程开得够多，速度却上不去，瓶颈几乎总是数据没有在正确的层上复用，而不是算力不足。profiler 报告的各级命中率、带宽利用率和 stall 原因，本质上就是在告诉你这四层各自被用得怎么样。
这一篇在知识体系里的位置
本篇是 GPU 硬件与编程模型的第二篇，上承 GPU 架构：warp 调度器（谁在计算、延迟怎么藏），下接三个具体机制专题：内存合并访问解决显存带宽怎么用满，shared memory bank conflict解决片上吞吐怎么不被串行化，GPU 显存 HBM 与 GDDR展开最底层的显存技术。往芯片物理方向走，HBM 的封装与接口实现见 HBM 存储接口（expert）；想对照 CPU/SoC 的存储设计，可读 Cache 组织结构、TCM 紧耦合内存与 存储分块与交织。
总结
GPU 内存层级是寄存器（线程私有、约 1 周期、256 KB/SM）→ shared memory/L1（block 内共享、软件显式搬移、32 个 bank、约 20~30 周期、百余 KB/SM）→ L2（全卡共享、硬件自动管理、40~50 MB、约 200 周期）→ HBM/GDDR 全局显存（数十 GB、2~3.35 TB/s 但单次访问 400~800 周期）构成的金字塔：越往下容量越大、作用域越广、延迟越长；GPU 编程的核心因此不是「算」而是「搬」——用分块把显存数据合并加载到 shared memory、在寄存器里反复复用几十次，把几百周期的慢路径摊薄；它与 CPU 全自动 cache 的根本区别在于 shared memory 这一层由软件掌控，硬件只负责 L1/L2、带宽和 warp 切换，常见坑是寄存器 spilling 压垮 occupancy、bank conflict 串行化片上访问、非合并访问浪费显存带宽，以及数据只用一次却多此一举地搬进 shared memory。