GPU L2 Cache逆向工程：Ampere微基准测试与性能调优实战-开发者社区

1. 项目缘起与整体思路拆解

1.1 为什么要对 L2 Cache 做逆向分析

GPU 性能调优做到一定程度，必然会撞上一堵墙：SM 内部的寄存器、共享内存、L1 这些资源，厂商文档给得还算清楚，但一旦跨出 SM，进入 L2 Cache 和显存子系统，公开资料就变得非常稀薄。NVIDIA 从 Volta 之后基本不再公开 L2 的替换策略、分区映射、sector 粒度这些底层细节，而这些东西恰恰决定了你的 kernel 在访存密集型场景下到底能跑多快。

我最初做这个项目的动机很直接：手上有一个访存模式比较特殊的 kernel，profiler 显示 L2 命中率忽高忽低，同样的数据量、同样的访问顺序，换个 block 配置命中率能差出 30%。官方文档只告诉你 L2 有多大、带宽多少，但完全不解释“为什么同样的访问会得到不同的命中率”。这种黑盒状态没法做优化，只能靠试。试到后来我意识到，与其盲目调参，不如把 L2 的行为规律系统性地测出来——这就是所谓“逆向工程”的起点。

需要先明确一点：这里的逆向工程不是去拆芯片、看版图，而是通过精心设计的微基准测试（microbenchmark），从外部可观测的时序和带宽数据反推内部结构参数。这是一种黑盒建模方法，在体系结构研究里很常见。Ampere 这一代（以 GA10x 和 GA100 为代表）的 L2 相比 Turing 有比较大的改动，尤其是 GA100 把 L2 切成了两个 partition，这个结构差异会直接影响跨 partition 访问的延迟，值得单独测。

1.2 逆向分析能拿到哪些关键参数

在动手之前，得先想清楚要测什么。L2 逆向工程的目标参数大致可以分成几类：

容量与关联度：L2 总容量、组相联的路数（ways）、cache line 大小。这些决定了冲突缺失（conflict miss）的边界。
替换策略：LRU、伪 LRU、还是某种自适应策略。这决定了你能否通过访问顺序“骗过”缓存。
分区与切片（slice）映射：物理地址如何映射到不同的 L2 slice，slice 之间是否独立。这直接影响 bank 冲突和跨 slice 带宽。
sector 粒度：Ampere 上 L2 的访问粒度是 32B sector 还是 64B line，部分填充怎么处理。
延迟与带宽特性：命中延迟、缺失延迟、不同访问模式下的有效带宽。
一致性相关行为：写分配（write-allocate）策略、写回（write-back）时机。

这些参数里，容量和关联度相对好测，替换策略和 slice 映射最难，因为它们不会直接暴露在任何一个计数器里，只能靠访问模式的“指纹”去推断。

1.3 测试方法论的整体框架

我采用的框架是经典的pointer-chasing + 访问模式扫描组合。核心思路是：

用 pointer-chasing 测延迟。构造一个链表，每个节点指向下一个节点，节点之间的地址间隔可控。通过改变工作集大小和步长，观察延迟曲线的拐点，拐点位置就对应容量边界。
用 streaming 访问测带宽。连续访问大块内存，改变访问粒度和并发度，观察有效带宽的变化，推断 sector 粒度和 slice 并行度。
用冲突构造测关联度。固定工作集大小，改变访问地址的 stride，当 stride 导致大量地址映射到同一个 set 时，命中率会骤降，通过扫描 stride 找到冲突点，反推 set 数量和 ways。

这套方法的好处是不依赖任何厂商内部工具，纯靠 CUDA 就能跑，可复现性强。坏处是测量噪声大，需要大量重复实验和统计处理，后面会专门讲怎么降噪。

提示：做这类微基准测试，一定要固定 GPU 时钟频率（用nvidia-smi -lgc锁频），否则动态调频会让延迟数据完全不可用。这是踩过的第一个大坑。

2. 核心细节解析与实操要点

2.1 测试环境搭建与时钟锁定

环境这块，我用的是单卡 Ampere 消费级和计算级各一张做对比，驱动和 CUDA 版本尽量新，因为老驱动在某些计数器行为上有差异。系统层面要做几件事：

关闭所有图形界面和后台进程，避免其他任务抢占 L2。
用nvidia-smi -pm 1开启持久模式。
用nvidia-smi -lgc <freq>锁定 SM 和显存频率到固定值。这一步极其关键，我试过不锁频直接测，同一段代码两次运行的延迟能差 40%，数据完全没法用。
用cudaDeviceSetLimit把 L1 和共享内存的 carveout 调到最小，尽量让 L1 不干扰 L2 的测量。不过要注意，L1 和 L2 是包含关系还是非包含关系，Ampere 上是非包含的（non-inclusive），所以 L1 的配置对 L2 命中率影响有限，但会影响你观测到的延迟。

时钟锁定之后，还要做一次“热身”，先跑几轮空转 kernel 把 GPU 拉到稳定温度，避免热降频。这些准备工作看起来琐碎，但直接决定了后面数据的可信度。

2.2 Pointer-Chasing 延迟测量的实现细节

Pointer-chasing 的原理很简单：一个链表，节点在内存里随机分布，每个节点存下一个节点的地址。kernel 从头部开始，逐个跳转，因为每次跳转都依赖上一次的结果，所以无法并行，测出来的就是单次访存的真实延迟。

实现上有几个坑：

第一，链表必须在 kernel 启动前就在 device 内存里构造好，不能在 kernel 里现构造，否则构造过程本身会污染 cache。我一般用 host 端生成随机排列，再cudaMemcpy过去。

第二，节点大小要控制。如果节点是 64B 对齐的，那每次跳转至少拉一个 sector；如果节点跨 line，会触发两次访问。为了测准，我把节点做成 128B 对齐，每个节点只放一个 8B 的指针，其余填充，确保每次跳转只触碰一个 line。

第三，要防止编译器优化。用volatile指针或者内联 PTX 的ld.global.cg（绕过 L1）来强制真实访存。我实测下来，用ld.global.cg最干净，因为它直接打到 L2，不经过 L1，测出来的就是 L2 的延迟。

第四，循环次数要够多。单次测量噪声大，我一般让每个链表跑 10000 次跳转，取平均，再重复 100 次取中位数。中位数比平均值稳，因为偶尔会有中断或调度抖动。

2.3 工作集大小与延迟曲线的解读

这是整个项目里信息量最大的一步。固定 stride（比如 128B），逐步增大工作集大小，从 4KB 一直加到 64MB，记录每个大小下的平均延迟。你会看到一条典型的阶梯曲线：

工作集小于 L1 容量时，延迟最低，大概几十个 cycle。
跨过 L1 边界，延迟跳升到 L2 命中延迟。
跨过 L2 边界，延迟再跳升到显存延迟。

Ampere 上我测到的典型值（锁频后）：L1 命中约 30 cycle，L2 命中约 200 cycle，显存约 400-600 cycle。注意这些数字随频率变化，所以一定要在锁频条件下比较。

关键在拐点。L2 的容量拐点不是一条锐利的线，而是一个渐变区间，因为关联度有限，工作集接近容量时就开始出现冲突缺失。我的做法是：对每个工作集大小，测出延迟分布，找到延迟开始显著上升的那个点，再结合冲突测试交叉验证。

2.4 关联度与替换策略的探测方法

关联度测试用 stride 扫描。固定工作集大小（略小于 L2 容量），让访问地址按不同 stride 递增。当 stride 使得大量地址映射到同一个 set 时，即使工作集没超容量，也会因为路数不够而频繁缺失，延迟飙升。

具体操作：stride 从 128B 开始，按 2 的幂次递增到几 MB，每个 stride 测一遍延迟。如果 L2 是 16 路组相联、line 64B、容量 4MB，那么 set 数 = 4MB / (64B × 16) = 4096。当 stride 等于 set 数乘以 line 大小（即 4096 × 64B = 256KB）时，所有访问会落到同一个 set，冲突最严重。实测中你会看到在 256KB 及其倍数附近出现延迟尖峰，这就验证了 set 数和 ways。

替换策略的探测更微妙。我用的方法是“访问序列指纹”：构造一个刚好超过 ways 的访问序列，比如访问 17 个映射到同一 set 的地址（假设 16 路），然后回头再访问第一个地址。如果是纯 LRU，第一个地址已经被踢出，会缺失；如果是伪 LRU 或随机替换，行为会不同。通过对比不同序列下的命中模式，可以区分 LRU 和近似 LRU。Ampere 上我测到的结果更接近一种自适应/伪 LRU 策略，不是严格 LRU，这也解释了为什么某些访问顺序下命中率会反常。

注意：替换策略的探测对噪声极其敏感，一定要在锁频、空载、多次重复的条件下做，单次结果不可信。

3. 实操过程与核心环节实现

3.1 链表构造与 kernel 骨架

先看链表构造。host 端生成一个 0 到 N-1 的随机排列，然后按排列顺序把节点串起来。每个节点 128B 对齐，只存下一个节点的 device 地址。

// host 端构造 struct Node { unsigned long long next; char pad[120]; }; std::vector<int> perm(N); std::iota(perm.begin(), perm.end(), 0); std::shuffle(perm.begin(), perm.end(), std::mt19937(42)); std::vector<Node> h_nodes(N); for (int i = 0; i < N; i++) { int cur = perm[i]; int nxt = perm[(i + 1) % N]; h_nodes[cur].next = (unsigned long long)(d_nodes + nxt); } cudaMemcpy(d_nodes, h_nodes.data(), N * sizeof(Node), cudaMemcpyHostToDevice);

kernel 端就是纯粹的指针追逐，用ld.global.cg绕过 L1：

__global__ void chase(unsigned long long* head, unsigned long long* out, int iters) { unsigned long long p = *head; for (int i = 0; i < iters; i++) { asm volatile("ld.global.cg.u64 %0, [%1];" : "=l"(p) : "l"(p)); } *out = p; }

这里ld.global.cg是关键，它表示“cache global”，在 Ampere 上会绕过 L1 直接访问 L2。如果你用普通的ld.global，数据可能被 L1 缓存，测出来的是 L1 延迟，就失去意义了。

计时用clock64()读 SM 时钟，在 kernel 内部记录起止，避免 host 端 launch 开销干扰。每个线程独立跑一条链表，最后取所有线程的中位数。

3.2 容量拐点的精确定位

跑完一轮扫描后，我拿到的是“工作集大小 vs 延迟”的散点。定位拐点不能靠肉眼，我用的是分段线性拟合：假设曲线由几段斜率不同的直线组成，用最小二乘找出斜率突变点。具体做法是对相邻点求差分，差分值超过阈值的地方就是拐点候选。

实测 Ampere 消费级卡的 L2 拐点大约在 4MB 附近（具体数值因型号而异），计算级卡因为 L2 更大，拐点在几十 MB。这里有个细节：拐点位置会随 stride 变化，因为 stride 影响有效工作集。所以容量测试要固定 stride，我一般用 128B（一个 line），保证每次访问都拉满一个 line。

另一个细节是 TLB 的影响。工作集大到一定程度，TLB 缺失会混进来，让延迟曲线多一个台阶。为了排除 TLB 干扰，我用大页（2MB page）分配内存，cudaMallocManaged配合cudaMemAdvise设置大页属性。这样在几十 MB 范围内基本不会有 TLB 缺失。

3.3 Slice 映射与跨分区延迟测量

GA100 的 L2 分成两个 partition，每个 partition 有自己的 slice。物理地址到 slice 的映射是 hash 过的，不是简单的取模。要测这个，我用的是“地址对延迟”方法：固定一个基准地址，然后扫描另一个地址，测两个地址交替访问的延迟。如果两个地址映射到同一个 slice，会有 bank 冲突，延迟偏高；映射到不同 slice，延迟低。

扫描下来，延迟随地址差的分布呈现周期性，周期对应 slice 的映射粒度。Ampere 上这个粒度我测到的是 256B 左右，也就是说相邻 256B 的地址块会映射到不同 slice。这个结论对优化很有用：如果你能让并发访问的地址分散到不同 slice，就能吃满 L2 带宽。

跨 partition 访问的延迟明显高于 partition 内访问，这也是为什么 GA100 上要尽量让数据和处理它的 SM 在同一个 partition 附近。不过这个优化在 CUDA 层面能做的有限，更多是给底层库开发者参考。

3.4 带宽测试与 sector 粒度推断

带宽测试用 streaming kernel：每个线程连续读一大块内存，用向量化加载（float4或ulonglong2）提高单次事务的数据量。改变并发线程数和访问粒度，观察有效带宽。

sector 粒度的推断靠“部分访问”实验：只访问一个 64B line 里的前 32B，看带宽是否减半。如果 L2 以 32B sector 为单位传输，那只访问 32B 应该能拿到接近满带宽的一半；如果以 64B 为单位，访问 32B 和 64B 的带宽应该差不多。实测 Ampere 上更接近 32B sector 行为，这意味着部分填充的 line 不会浪费整个 line 的带宽，对稀疏访问模式是利好。

这个结论直接影响优化策略：如果你的 kernel 访问是稀疏的、每次只碰几个字节，那在 Ampere 上不用太担心“拉整条 line 浪费带宽”的问题，sector 机制会帮你省。

4. 常见问题与排查技巧实录

4.1 数据噪声大、不可复现怎么办

这是最常见的问题，几乎每个人第一次做都会遇到。我的排查顺序是：

先确认时钟锁了没有。没锁频的话，后面所有排查都是白费。
确认没有其他进程占用 GPU。用nvidia-smi看有没有残留进程，必要时cudaDeviceReset。
增加重复次数，用中位数而非平均值。异常值往往来自偶发的调度抖动。
检查链表是否真的随机分布。如果构造时随机性不够，访问模式会退化，延迟数据会失真。
检查是否被 L1 缓存。用ld.global.cg强制绕过 L1，或者用cudaFuncSetAttribute把 L1 carveout 调到最小。

我踩过最深的坑是第 5 条：早期用普通ld.global，测出来的“L2 延迟”其实是 L1 延迟，因为链表节点少的时候全被 L1 缓存了。后来统一改成ld.global.cg才拿到真实数据。

4.2 拐点模糊、找不到清晰边界

容量拐点模糊通常有两个原因：一是关联度有限导致渐变，二是测量粒度太粗。解决办法：

在工作集接近预期容量时，加密采样点，比如每 64KB 测一次，而不是每 1MB。
用冲突测试交叉验证。如果 stride 扫描在某个 stride 下延迟骤升，那个点对应的 set 数乘以 ways 乘以 line 大小就是容量，比单纯看延迟曲线更准。
排除 TLB 干扰，用大页。

4.3 替换策略测不准的应对

替换策略是最难测的，因为它是“软”行为，没有硬边界。我的经验是：

不要指望单次实验得出结论，要设计多组对照实验。比如 LRU 和随机替换在“访问 N+1 个同 set 地址”这个场景下行为不同，但在“访问 N 个地址循环”场景下可能相似。多设计几个场景，找行为差异。
用统计方法。同一个序列跑几百次，看命中率的分布。LRU 的命中率是确定的，随机替换会有方差。
接受“近似”结论。工业界很多缓存用的就是伪 LRU，你测出来“接近 LRU 但不完全一致”是正常的，不必强求精确匹配某个理论模型。

4.4 常见问题速查表

问题现象	可能原因	排查方向
延迟数据波动超过 30%	未锁频或后台干扰	锁频、清空后台进程
测出的延迟远低于预期	被 L1 缓存	改用 ld.global.cg
容量拐点找不到	采样太粗或 TLB 干扰	加密采样、用大页
stride 扫描无尖峰	stride 未覆盖 set 周期	扩大 stride 范围到 MB 级
带宽远低于理论值	并发度不足或访问不连续	增加线程数、用向量化加载
跨 slice 延迟测不出差异	地址映射 hash 均匀	换更大地址范围扫描

4.5 几个独家避坑心得

第一，先测延迟再测带宽。延迟测试对噪声更敏感，要在 GPU 状态最好的时候先做。带宽测试相对鲁棒，可以后面做。

第二，链表节点数要远大于预期容量。如果你想测 4MB 的 L2，链表至少要覆盖 64MB，否则工作集上不去，测不到缺失区。

第三，不要忽略写操作。读延迟和写延迟、写分配策略是两回事。我后来补做了写测试，发现 Ampere 的写分配行为和读有明显差异，写操作会先分配 line 再写，这会影响写密集 kernel 的优化策略。

第四，记录原始数据。逆向工程的结论依赖大量实验，中间任何一步的原始数据都可能在后面对比时用上。我习惯把每次实验的配置、原始时序、统计结果都存成 CSV，方便回溯。

第五，交叉验证。任何一个参数，至少用两种独立方法测。比如容量，既用延迟拐点测，也用 stride 冲突测，两者吻合才敢下结论。单一方法的结论很容易被系统性误差带偏。

这套方法跑下来，我对 Ampere L2 的行为有了比较完整的画像：容量、关联度、sector 粒度、slice 映射、替换策略倾向，都有了可复现的数据支撑。这些结论后来直接指导了我优化几个访存密集 kernel 的 block 配置和地址布局，命中率提升明显。整个过程最深的体会是：黑盒不可怕，可怕的是没有系统性的测量方法。只要实验设计得当，外部可观测的时序数据足以反推出足够多的内部结构信息，支撑起实际的性能优化决策。