GPU显存层次 L1/L2/HBM 存储体系

可视化演示

展示系统处理流水线、实时性能指标曲线及多方案对比分析，包含数据流转、吞吐延迟监控与方案评估。

存储层级总览

后厨高峰到来前，主厨会把盐罐放在灶台边，切好的配料放在操作台，半成品放进冰箱，大宗食材留在冷库。每一层“取货距离”不同，出菜速度也因此不同。GPU 的存储体系正是这样一层层“按距离定价”：离计算单元越近，速度越快、容量越小、造价越贵。

这里的“离得近”指的是物理上更靠近 SM（Streaming Multiprocessor，可以想象成一组灶台）。寄存器在 SM 内部，L1 缓存也在 SM 内；L2 缓存是多个 SM 共用的“后厨冰箱”；HBM 则是封装在 GPU 旁的片外显存。为了让数据访问尽量“近水楼台”，设计者把最常访问的数据放在最快的层，这就是局部性原理：程序倾向于反复访问同一批数据和时间上邻近的数据。

可以用平均内存访问时间（AMAT）来统一衡量“取一趟食材到底要多久”：命中时的时间，加上未命中时额外付出的代价。

主厨会把最常用的调料放手边，GPU 也把最常访问的数据放 L1。一旦放错位置，每次做菜都要跑一趟冷库，AMAT 就会飙升，整条流水线都慢下来。

寄存器与L1缓存

盐罐是主厨的“私人领地”，别人不会顺手拿走；切好的葱姜蒜放在操作台上，虽然也近，但属于本灶台共享的备料。寄存器就是线程的盐罐，L1 缓存就是操作台。寄存器是每个线程私有的，只有当前线程能读写，因此延迟极低、带宽极大，通常一个时钟周期就能取到。

L1 缓存是 SM 内部的高速缓存，缓存该 SM 最近访问过的全局内存数据。当线程需要某个地址的数据时，先查 L1；命中则快速返回，未命中才去 L2 和 HBM。L1 容量通常是几十 KB，以 128 字节的缓存行为单位加载。编译器会将局部变量优先分配到寄存器，但如果变量太多或生命周期太长，就会“溢出”到局部内存，速度骤降。

寄存器的数量非常有限，典型一个线程最多 255 个；L1 要与共享内存共享同一块片上 SRAM，配置需要权衡。这就像操作台就那么大，摆满了盐罐就放不下备菜。正因如此，写 CUDA 时要注意减少每个线程的寄存器占用，同时提高 L1 命中率，让数据尽量“留在手边”。

L2缓存机制

后厨里有一台公用冰箱，几位灶台师傅都会从里面取腌肉，也会往里放做好的半成品。由于大家都用它，必须有个规则保证“不会拿错”。GPU 的 L2 缓存就是这台公用冰箱，它是多个 SM 共享的片上缓存，也是全局内存访问的一致性点。

L1 未命中后，请求会来到 L2。L2 容量比 L1 大得多，常见几 MB 到几十 MB，带宽也非常可观。更重要的是，L2 要对所有 SM 的读写请求做协调：当一个 SM 修改了某地址，其他 SM 必须看到最新值，否则就像两个厨师拿着同一块肉，一个腌了一个没腌，菜品就不一致了。GPU 用硬件一致性协议维护这份“台账”，让每个 SM 看到的全局内存视图保持一致。

L2 的存在让跨 SM 的数据共享更高效。比如两个线程块都要读同一份权重，第一次从 HBM 取走后，第二次就能从 L2 命中，省掉一次“跑冷库”。不过 L2 也有延迟代价，约为 L1 的好几倍，所以不能指望它替代 L1。设计 Kernel 时，应尽量让数据在 L1 和寄存器中完成复用，把 L2 当作第二道缓冲网。

HBM高带宽显存

冷库是后厨最大的仓库，猪牛羊肉、整袋面粉都放在里面，但每次取货都要坐货梯，来回一趟耗时远高于操作台取菜。HBM（High Bandwidth Memory）就是 GPU 的冷库：容量最大，通常几十 GB，用来存放模型权重、激活值、中间结果等大块数据。

HBM 是堆叠式 DRAM，利用 3D 堆叠和 TSV（硅通孔）把多颗 DRAM 芯片垂直连在一起，再通过硅中介层与 GPU 封装在一块基板上。相比传统 GDDR 显存，HBM 用更小的面积和功耗换来了极高的带宽，例如 HBM2e 单颗即可提供超过 1 TB/s 的带宽。之所以能这么高，除了堆叠带来的短距离连线，还因为它的总线位宽非常宽。

理论带宽可以由频率和位宽估算：

冷库传送带再宽，也不可能比操作台取菜更快。HBM 延迟远高于 L1/L2，所以即便拥有高带宽，程序员依然要在意“去冷库的次数”，否则再宽的传送带也会排队。

访存模式优化

备菜师傅会把同一道菜需要的配料码在一个托盘里，一次端到灶台；而不是每取一根葱就跑一趟。GPU 的访存也一样：一个线程束（32 个线程）如果访问连续地址，就可以合并成少数几个内存事务，这就是合并访存（coalesced access）。

具体来说，GPU 以 32 字节为最小内存事务单位。32 个线程各读一个 4 字节整数，如果地址连续，这 128 字节只落在 4 个事务内；如果地址跳着访问，比如每个线程间隔 64 字节，就可能触发 32 个事务，带宽浪费 8 倍。因此，写内核时第一原则是“相邻线程访问相邻地址”，同时尽量做到 16 字节对齐。

除了合并访存，还要注意数据复用：一个数据被多次使用时，先把它从 HBM 请进共享内存，再反复读。另一个常见坑是共享内存的 bank 冲突——共享内存被分成多个 bank，若同一波线程访问同一 bank 的不同地址，就会串行化，等于操作台排队取菜。通过填充或改变映射方式，可以让访问均匀散开。

数据驻留策略

开业前，主厨会把最受欢迎的食材从冷库搬到后厨冰箱，并安排专人负责补货。数据驻留策略的本质，就是人工决定“什么数据放在哪一层、何时搬、何时清”，从而减少对慢速层的访问。

CUDA 中，这类策略通常由程序员显式控制。例如，用 __ldg 指令走只读数据缓存，用 cudaMemcpyAsync 在后台异步把数据从 HBM 拷贝到共享内存，或者使用 cp.async 指令实现异步预取。更进阶的还有 L2 驻留控制 API，允许把某些数据“钉”在 L2 中，避免被其他访问挤出去。

一个经典案例是分块矩阵乘法：把 A 和 B 的小块先搬进共享内存，让同一线程块的多次计算反复复用，从而把访存量从 O(N³) 降到 O(N³/B)，B 为分块边长。这相当于把一整袋土豆搬上操作台，而不是每次炒菜都去冷库挖一个土豆。数据驻留策略的核心就是提高复用率，让每次 HBM 访问都“物有所值”。

带宽计算与调优

冷库货梯的运力是固定的，订单再多也只能一趟一趟送。判断一个 Kernel 是“被计算拖累”还是“被访存拖累”，常用算术强度（Arithmetic Intensity）这个指标：

调优时，先用 Profiler（如 NVIDIA Nsight Compute）查看内存吞吐利用率、L2 命中率、HBM 带宽利用率等指标。如果 HBM 带宽利用率接近 100%，说明传送带已经满载，这时应该用合并访存、共享内存复用、数据驻留等手段减少搬运量；如果计算单元空闲等待，则要增加并行度或减少分支。这就像先看清到底是货梯不够宽还是后厨人手不够，再对症下药。

AI 应用场景

Python 演示

import math def count_transactions(addresses, transaction_size=32):    """统计一组访存地址触发的缓存事务数"""    blocks = set()    for addr in addresses:        blocks.add(addr // transaction_size)    return len(blocks)  def coalesced_access(thread_id):    """相邻线程访问连续4字节地址"""    return thread_id * 4  def strided_access(thread_id):    """相邻线程访问间隔64字节地址"""    return thread_id * 64  def main():    nthreads = 32    addr_coal = [coalesced_access(t) for t in range(nthreads)]    addr_strided = [strided_access(t) for t in range(nthreads)]    co = count_transactions(addr_coal)    st = count_transactions(addr_strided)    print(f"Coalesced   : {co} transactions")    print(f"Strided     : {st} transactions")    print(f"Amplification: {st / co:.1f}x")  if __name__ == "__main__":    main()

1. def count_transactions(addresses, transaction_size=32):用集合记录不同事务块，块数越少越好。
2. blocks.add(addr // transaction_size)把地址映射到所在的 32 字节事务块，同一块内的请求可以合并。
3. addr_coal = [coalesced_access(t) for t in range(nthreads)]构造连续地址序列，模拟合并访存。
4. addr_strided = [strided_access(t) for t in range(nthreads)]构造跨步地址序列，模拟糟糕的非合并访存。
5. co = count_transactions(addr_coal)统计连续访问事务数，结果应为 4。
6. st = count_transactions(addr_strided)统计跨步访问事务数，结果应为 32。
7. print(f"Amplification: {st / co:.1f}x")输出事务放大倍数，说明跨步访存浪费了多少带宽。

常见误区

面试题