NVIDIA GPU 中的 L2 Cache
在 NVIDIA GPU 中，L2 Cache 是位于 HBM（高带宽内存）之前的最后一级片上缓存（Last Level Cache, LLC）。如果熟悉缓存一致性协议的话，对这个cache 的工作机制会比较容易理解。接下来介绍一下其在 GPU 存储器层次结构中的位置和作用。
1. GPU 存储层次结构（从计算核心向外）
寄存器 (Registers)：最快，每个线程私有。
L1 Cache / Shared Memory：位于每个 SM（流多处理器）内部。速度快，容量小，由 SM 内的 CUDA 核心共享。
L2 Cache (LLC)：位于所有 SM 之外，但在 HBM/DRAM 内存之外。它是一个全局的、统一的缓存，为整个 GPU 上的所有 SM 服务。
HBM (High Bandwidth Memory) / DRAM：最慢，容量最大，是板载的物理内存（显存）。
数据流路径：
CUDA Core -> L1/Shared Memory -> (互联网络) -> L2 Cache -> (内存控制器) -> HBM/DRAM
2. L2 Cache 作为 LLC 的关键特性和作用
全局性 (Global)：
不同于 L1 缓存是每个 SM 独有的，L2 缓存是所有 SM 共享的资源。它缓存来自所有 SM 的内存请求数据，无论这些请求是来自全局内存、常量内存还是本地内存。统一性 (Unified)：
它是一个“统一缓存”，意味着指令、数据和纹理等都可以存储在其中，而不像某些 CPU 架构那样有独立的指令缓存和数据缓存。高速缓冲区 (High-Speed Buffer)：
其主要目的是减少对慢速的 HBM 显存的访问次数。当多个 SM 中的线程访问全局内存中的同一数据时（例如，一个只读的查找表），该数据可以被缓存到 L2 中。后续的访问如果命中 L2，就可以极大地降低延迟并节省 HBM 的带宽。实现内存访问优化：
L2 缓存对于实现某些内存访问模式的优化至关重要。例如：持久化访问 (Persistent Access)：在计算能力 6.0+ 的 GPU 中，可以使用__ldg()函数或const __restrict__关键字来提示编译器将数据持久化在 L2 缓存中，供整个网格中的所有线程块重复使用；全局内存的合并访问：虽然主要由 L1 处理，但 L2 作为后备；与架构强相关：
L2 缓存的大小因 GPU 架构而异，并且是衡量 GPU 性能的一个重要指标。例如：NVIDIA H100 (Hopper): 50 MB；NVIDIA A100 (Ampere): 40 MB；NVIDIA Tesla V100 (Volta): 6 MB 可以看到，随着架构迭代，L2 缓存的容量在显著增加，以应对日益增长的计算能力和对内存带宽的需求。
3. 小结一下
| 特性 | 描述 |
|---|---|
| 位置 | 在所有 SM 和内存控制器之间 |
| 角色 | 最后一级缓存 (LLC)，位于 HBM 显存之前 |
| 范围 | 全局共享，为所有 SM 服务 |
| 类型 | 统一缓存（存储指令、数据等） |
| 主要目的 | 减少对 HBM 的访问，降低延迟，提高有效带宽 |
因此， NVIDIA GPU 的缓存中，L2 Cache 是那个最外侧的、位于高带宽 HBM 内存之前的最后一级大型片上缓存（LLC）。它的效率和命中率对GPU的整体性能有着至关重要的影响。其他公司的 gpu 中的 cache 构成也类似。