缓存体系详解：L1/L2缓存架构、缓存行、缓存命中与未命中的代价

GPU 的存储层次结构是决定访存性能的基石。与 CPU 依赖大容量缓存不同，GPU 依靠 高带宽显存 与 多级缓存 的配合来隐藏延迟。理解 L1/L2 缓存的组织方式、缓存行的行为以及命中/未命中的代价差异，是编写高效 GPU 内核的前提。本节将从硬件视角拆解缓存体系，并给出可操作的优化依据。

4.1 核心原理：GPU 缓存层次与数据流

现代 NVIDIA GPU（Volta 及以后架构）的缓存体系分为三个关键层级：寄存器文件（每个线程私有）、L1 缓存/共享内存（每个 SM 私有）以及 L2 缓存（全局共享，所有 SM 可见）。L1 与共享内存在物理上位于同一块 SRAM 上，可通过配置分割容量。L2 缓存是 GPU 与显存之间的最后一级缓存，所有全局内存访问（包括纹理、表面）都必须经过 L2。

下图展示了从线程访存指令到数据返回的完整路径，以及各级缓存的角色与典型延迟。

关键点在于：L1 命中 时延迟极低（约 30 周期），而 L2 命中 约 200 周期，显存访问 则高达 400~800 周期。因此，优化目标就是让尽可能多的访问在 L1 或 L2 中完成。

4.2 缓存行与内存合并：命中率的基石

GPU 缓存与 CPU 缓存最大的不同在于 合并访问（Coalescing） 机制。当线程束中的 32 个线程同时访问全局内存时，硬件会检查这些地址是否落在连续的 128 字节缓存行内。如果完全连续，则只需 1 次事务 即可完成；如果分散，则可能产生 多次事务，造成带宽浪费。

考虑以下两种访问模式：

// 模式 A：合并访问（推荐）
int idx = blockIdx.x * blockDim.x + threadIdx.x;
float val = array[idx];  // 线程 0~31 访问连续地址

// 模式 B：非合并访问（跨步访问）
int idx = (blockIdx.x * blockDim.x + threadIdx.x) * 2;
float val = array[idx];  // 线程访问地址间隔 8 字节，浪费一半带宽

模式 A 中，32 个线程访问 32 个连续的 4 字节浮点数，恰好覆盖 128 字节，只需 1 个缓存行事务。模式 B 中，线程访问间隔 8 字节，32 个线程覆盖 256 字节范围，需要 2 个缓存行，但实际只使用其中一半数据，有效带宽减半。

4.3 案例：矩阵转置的缓存命中分析

矩阵转置是展示缓存行为差异的经典案例。假设一个 1024×1024 的浮点矩阵，按行主序存储。朴素转置内核如下：

__global__ void transpose_naive(const float* in, float* out, int n) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < n && col < n) {
        out[col * n + row] = in[row * n + col];
    }
}

该内核的访存特征为：

• 读 in[row * n + col]：同一行内线程访问连续地址，合并良好，L1 命中率高。
• 写 out[col * n + row]：线程访问列方向，地址间隔 n×4 字节，完全非合并，每个 128B 缓存行只写入 4 字节，造成严重的写放大。

通过 共享内存分块 可以解决该问题：先将数据块读入共享内存（合并读），再从共享内存转置后写回（合并写）。

__global__ void transpose_tiled(const float* in, float* out, int n) {
    __shared__ float tile[TILE][TILE];
    int x = blockIdx.x * TILE + threadIdx.x;
    int y = blockIdx.y * TILE + threadIdx.y;
    if (x < n && y < n)
        tile[threadIdx.y][threadIdx.x] = in[y * n + x];
    __syncthreads();
    x = blockIdx.y * TILE + threadIdx.x;
    y = blockIdx.x * TILE + threadIdx.y;
    if (x < n && y < n)
        out[y * n + x] = tile[threadIdx.x][threadIdx.y];
}

分块后，读写均变为合并访问，L1 缓存命中率显著提升。实测在 A100 上，分块版本比朴素版本快 3~5 倍，主要收益来自缓存行利用率的提升。

4.4 缓存未命中的代价量化

为了直观理解未命中的代价，下表对比了不同访问位置在典型 GPU 上的延迟与带宽消耗：

访问目标	典型延迟	每线程束事务数	有效带宽
寄存器	0 周期	0	无限
L1 命中	~30 周期	1	理论峰值
L2 命中	~200 周期	1	理论峰值
显存（L2 未命中）	~400-800 周期	1（合并）/ 多（非合并）	取决于合并度
非合并显存访问	~800 周期	4~32	降至 1/4 ~ 1/32

从表中可见，非合并访问的代价不仅是延迟增加，更严重的是事务数倍增。一个线程束的非合并访问可能产生 32 个独立事务，每个事务只使用 4 字节，导致实际带宽利用率仅为理论值的 3%。

4.5 误区与复盘：常见缓存优化陷阱