GPU L2 Cache逆向分析：纯软件微基准测量容量、关联度与替换策略

1. 项目缘起与整体设计思路

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

GPU 性能调优做到一定程度，必然会撞上一堵墙：SM 内部的寄存器、共享内存、L1 这些你都能通过文档和微基准摸得比较清楚，但一旦数据出了 SM，落到 L2 Cache 这一层，很多行为就变得“黑盒”了。厂商的 whitepaper 只会告诉你 L2 有多大、带宽大概多少，但具体到 set 的映射方式、替换策略、sector 的粒度、跨 slice 的哈希规则，这些细节几乎不会公开。而恰恰是这些细节，决定了你的 kernel 在访存密集型场景下到底是跑满带宽还是频繁 miss。

这个项目的目标很明确：用一套纯软件、可复现的微基准方法，把一块 Ampere 架构 GPU 的 L2 Cache 关键参数和内部组织方式“测”出来。注意我用的是“测”而不是“读文档”，因为逆向的核心价值就在于——文档没写的，我用实验数据反推出来。整套方法不依赖任何厂商内部工具，只需要能跑 CUDA kernel、能读 clock 和计时就行。

适合谁来参考？一类是做高性能计算、想把手头 kernel 的访存效率再压榨一轮的工程师；另一类是做体系结构研究、需要真实硬件参数来校准模拟器的同学。哪怕你只是想搞清楚“为什么我的 kernel 换个访问顺序性能差一倍”，这套思路也能给你答案。

1.2 整体方案选型：为什么用计时法而不是性能计数器

逆向 Cache 有两条主流路线。第一条是读硬件性能计数器，比如通过厂商提供的 profiling 接口去读 L2 的 hit/miss、sector 访问数。这条路看起来直接，但问题在于：计数器本身是厂商定义好的聚合量，它不会告诉你“第 17 号 set 被映射到了哪个 slice”，而且很多底层事件在消费级卡上根本不可用。第二条路是纯计时法，也就是构造特定的访存模式，通过测量延迟和带宽的变化来反推结构。这条路慢、需要大量对照实验，但胜在通用、可控、不依赖任何特权接口。

我最终选的是计时法为主、计数器为辅的混合策略。核心逻辑是：Cache 的很多结构特征会体现在“访问时间随地址/步长的变化曲线”上。比如你要测 Cache 容量，就扫描一个数组，逐步增大工作集，当工作集超过容量时，平均访问延迟会出现一个明显的拐点，这个拐点对应的字节数就是有效容量。同理，测关联度就用“冲突地址”构造法，测 set 映射就用步长扫描找周期性。

这里有个关键设计决策：所有实验都必须在“冷 Cache”和“稳定时钟”两个前提下做。冷 Cache 靠的是每次实验前用一个大数组把 L2 冲刷干净；稳定时钟靠的是锁定 GPU 频率、关闭动态调频，否则你测出来的延迟波动能到 30% 以上，根本没法分析。这两点后面会展开讲，是整套方法能不能成立的地基。

1.3 实验平台与工具链的取舍

平台方面，我用的是一块 Ampere 架构的消费级显卡，具体型号不重要，因为方法本身是通用的。工具链上，CUDA C++ 写微基准 kernel，host 端用标准 chrono 计时配合 CUDA event 做精确测量。为什么不直接用 CPU 计时？因为 kernel 启动有开销，CPU 侧的 wall clock 会把启动延迟算进去，而 CUDA event 是记录在 GPU 时间线上的，精度高得多。

编译选项上有个坑：一定要关掉编译器的激进优化，尤其是那些会把你的“看似无用”的访存操作优化掉的 pass。我的做法是把测量结果写到一个 volatile 全局变量里，强制编译器保留访存。另外-O3可以开，但--use_fast_math这类会改变浮点行为的选项要慎重，因为我们的实验里如果用到浮点累加来防止优化，精度变化可能引入噪声。

提示：整个逆向过程建议在一台“干净”的机器上做，后台不要跑其他占 GPU 的任务。我试过一边跑实验一边开着浏览器硬件加速，结果 L2 的测量数据直接偏了 15%，排查了半天才发现是别的进程在抢 Cache。

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

2.1 L2 Cache 的关键参数到底有哪些

在动手之前，得先明确我们要逆向的“目标参数”清单，否则实验会没有方向。对一块 GPU 的 L2 来说，真正影响性能的参数主要有这么几类：

• 容量（Capacity）：L2 的总大小，以及每个 slice 分到多少。注意 Ampere 的 L2 是分 slice 的，物理上是多块，逻辑上对 SM 呈现为一个统一地址空间。
• 关联度（Associativity）：每个 set 里有多少 way。关联度低会导致冲突 miss 增多，尤其是当你的访问步长和 set 数形成整数倍关系时。
• Set 数量与索引位（Index bits）：地址的哪几位用来选 set，这决定了什么样的地址会落到同一个 set。
• Sector 粒度：L2 和内存之间的传输粒度。Ampere 上常见的是 32 字节 sector，但实际行为需要验证。
• 替换策略：LRU、伪 LRU 还是随机？这直接影响冲突场景下的表现。
• 跨 slice 哈希：地址是怎么分布到不同 slice 上的，这决定了带宽能不能打满。

这几项里，容量和关联度相对好测，set 索引位和哈希规则最难，因为它们涉及地址位的具体映射，需要设计精巧的对照实验。

2.2 计时法的核心原理：延迟拐点与冲突构造

计时法的底层逻辑其实很朴素：Cache 命中快，Cache 缺失慢。L2 命中的延迟大概在 200 个时钟周期量级，而一旦 miss 到显存，延迟会跳到 400 到 600 周期甚至更高。这个差距足够大，大到可以通过统计平均延迟来区分。

测容量的经典做法是“指针追逐”（pointer chasing）。你构造一个链表，每个节点指向下一个节点，节点之间的地址间隔经过精心设计。然后你沿着链表走一圈，测量单次跳转的平均延迟。当链表的总工作集小于 L2 容量时，所有节点都在 L2 里，延迟低；当工作集超过容量，部分节点被挤出，延迟就会抬升。把工作集从 1MB 扫到 64MB，画一条延迟-容量曲线，拐点就是容量。

测关联度用的是“冲突构造”。假设你知道 set 数量是 S，那么你可以构造 S+1 个地址，让它们全部映射到同一个 set。如果关联度是 16，那么第 17 个地址就会把前面的挤出去，产生冲突 miss。通过逐步增加“同 set 地址”的数量，观察延迟在第几个地址开始飙升，就能反推关联度。

注意：指针追逐的链表构造必须保证节点之间的依赖是串行的，否则 GPU 的乱序执行和预取会把延迟隐藏掉，你测出来的就不是真实延迟了。我的做法是让每个节点的值作为下一个节点的地址偏移，形成真正的数据依赖链。

2.3 冷 Cache 与时钟锁定的实操细节

前面提到冷 Cache 和稳定时钟是地基，这里说具体怎么做。

冷 Cache 的标准操作是：在每次测量前，启动一个“冲刷 kernel”，它顺序读取一个远大于 L2 容量的数组（比如 256MB），把 L2 里的旧数据全部替换掉。这个冲刷 kernel 本身也要计时排除，不能让它污染测量。我通常会在冲刷后插入一个__syncthreads()和一段空转，确保冲刷的访存全部落地。

时钟锁定分两步。第一步是在 host 端用厂商提供的管理接口把 GPU 的图形时钟和内存时钟锁到固定值，关掉自动超频和动态调频。第二步是在 kernel 内部用clock64()读取 SM 时钟周期数来计时，而不是用 wall clock。为什么？因为即使你锁了频，不同 SM 之间的时钟也可能有微小漂移，用周期数计时能消除这个影响。clock64()返回的是 64 位周期计数，精度足够。

这里有个实测经验：Ampere 上如果不开时钟锁定，L2 延迟的测量标准差能到 ±40 周期；锁定之后能压到 ±5 周期以内。这个改善对后续分析拐点是决定性的。

2.4 地址生成与对齐的坑

构造测试地址时，对齐是个容易被忽视但影响巨大的因素。GPU 的访存是按 sector 对齐的，如果你的测试地址没有按 128 字节对齐，一次逻辑访问可能触发两个 sector 的传输，延迟直接翻倍，把 Cache 行为的信号淹没掉。

我的做法是：所有测试数组都用cudaMalloc分配，它返回的地址天然是 256 字节对齐的。然后在数组内部，所有节点的偏移都取 128 字节的整数倍。这样每次访问都精确落在一个 sector 内，信号干净。

另一个坑是页表。GPU 的虚拟内存到物理内存的映射是以大页为单位的，如果你分配的数组跨越了页边界，物理地址可能不连续，导致你精心设计的“地址间隔”在物理层面失效。规避方法是：分配一个比实际需要大得多的数组，然后只用其中一段，确保这段落在同一个大页内。具体多大？我一般分配 2 倍于实验所需，然后从中间开始用。

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

3.1 容量测量的完整流程与参数计算

容量测量是整个逆向的第一步，也是最容易出结果的一步。完整流程如下。

第一步，确定扫描范围。Ampere 消费级卡的 L2 容量通常在 4MB 到 6MB 之间，专业卡可能到 40MB 以上。所以扫描范围我设为 512KB 到 64MB，按 2 的幂次取点，每个点测三次取中位数。

第二步，构造指针追逐链表。链表节点结构体只放一个uint64_t next，指向下一个节点的地址。节点之间的步长设为 128 字节，这样每个节点独占一个 sector，避免预取干扰。链表长度从 4096 个节点（512KB）开始，逐步翻倍到 524288 个节点（64MB）。

第三步，kernel 实现。每个线程独立走一条链表，走 N 圈，用clock64()记录总周期数，除以总跳转次数得到平均延迟。这里要注意：线程数不能太多，否则多个线程的访存会互相干扰 L2 状态。我一般用 32 个线程，每个线程走不同的链表。

第四步，数据分析。把“工作集大小”作为横轴，“平均延迟”作为纵轴画图。你会看到一条先平后升的曲线。平的那段是 L2 命中，升的那段是开始 miss。拐点的确定方法：对曲线做分段线性拟合，找到斜率突变最大的点。实测下来，这个拐点对应的容量和真实 L2 容量的误差能控制在 5% 以内。

参数计算上有个细节：链表节点数乘以 128 字节就是工作集大小，但实际有效容量要减去链表指针本身的开销吗？不用，因为每个节点就是一个 128 字节的 sector，指针就在这个 sector 里，没有额外开销。

3.2 关联度测量的冲突地址构造法

关联度测量比容量测量更考验设计。核心思路是构造一组“必然冲突”的地址。

假设我们已经从容量测量中推断出 set 数量 S（容量除以关联度再除以 sector 大小，但关联度未知，所以这里有个循环依赖，实际做法是先假设一个关联度范围，比如 8 到 32，然后逐一验证）。

具体操作：选取一个基准地址 base，然后构造 K 个地址，每个地址都是 base 加上 i 倍的 (S × 128) 字节，i 从 0 到 K-1。这样这 K 个地址的 set 索引位完全相同，必然落到同一个 set。然后让一个线程依次访问这 K 个地址，测量平均延迟。当 K 小于等于关联度时，所有地址都能待在 Cache 里，延迟低；当 K 等于关联度加一时，就会发生冲突，延迟跳升。

实测中，我会把 K 从 1 扫到 64，每个 K 值测 1000 次取平均。延迟曲线会在某个 K 值处出现台阶，那个台阶的位置就是关联度。Ampere 上我测出来的结果是 16 way，和很多公开的体系结构分析一致。

提示：这个实验对地址对齐极其敏感。如果 (S × 128) 这个步长没有正确对齐到 set 索引的边界，冲突就不会发生，你会误以为关联度很高。建议先用一个已知关联度的简单场景（比如 L1）验证你的地址生成逻辑是否正确。

3.3 Set 索引位的定位：步长扫描法

Set 索引位是地址的哪几位，这个最难测，但也不是无迹可寻。方法是“步长扫描”。

原理是这样的：如果你以某个固定步长 stride 访问数组，当 stride 恰好等于 set 数量乘以 sector 大小时，所有访问都会落到同一个 set，冲突最严重，延迟最高。所以你可以扫描 stride，找到延迟峰值对应的 stride 值，这个值就等于 S × 128。知道了 S × 128，再结合容量和关联度，就能反推 set 索引位的位置。

实操中，stride 从 128 字节扫到 1MB，按 128 字节递增。每个 stride 下，访问一个固定大小的数组（比如 8MB，确保超过 L2），测量平均延迟。你会看到延迟曲线在几个特定的 stride 值处出现尖峰，这些尖峰对应的就是冲突步长。

我实测下来，Ampere 的 L2 set 索引并不是简单的连续位，而是经过了一层哈希。这意味着单纯的步长扫描只能找到“主冲突步长”，还有一些次级冲突步长对应哈希函数的周期性。要完全还原哈希规则，需要结合多组不同基地址的实验数据做拟合，这部分工作量很大，但如果你只是想做性能调优，找到主冲突步长通常就够了。

3.4 替换策略的判定实验

替换策略决定了 Cache 在冲突时“踢掉谁”。LRU 会踢最久未使用的，随机策略则不可预测。判定方法如下。

构造一个 set 里能放下的最大地址集合（数量等于关联度），然后按顺序访问它们，再重复访问第一个地址。如果是 LRU，第一个地址会被保留（因为它是最近最少使用的反面——它刚被访问过？不对，这里要仔细）。正确的做法是：先访问地址 A、B、C...直到填满 set，然后访问一个新的冲突地址 X，把某个 way 挤出去。接着立即访问 A，如果 A 还在，说明被踢的是别人；如果 A 不在了，说明 LRU 踢了最久未用的 A。

通过设计不同的访问序列，可以区分 LRU、FIFO 和随机。我测下来 Ampere 的 L2 更接近伪 LRU，也就是用一个近似位来模拟 LRU，而不是精确维护每个 way 的时间戳。这个结论对性能调优的意义在于：伪 LRU 在某些访问模式下会比真 LRU 更容易发生冲突，所以你的 kernel 如果对某个 set 反复访问，可能需要手动做地址偏移来打散。

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

4.1 测量数据波动大怎么办

这是最常见的问题。表现是同一组参数测三次，延迟能差 20% 以上。排查顺序如下。

先查时钟。用nvidia-smi之类的工具确认 GPU 时钟是否真的锁定了。很多时候你以为锁了，实际上驱动在后台又把它调回去了。我遇到过一种情况：锁频命令执行成功，但因为没有关掉 persistence mode，GPU 在空闲时降频，一跑 kernel 又升频，导致测量期间频率不稳定。

再查干扰。用nvidia-smi看有没有其他进程占着 GPU。哪怕是一个小的显示进程，只要它在刷新，就会周期性访问显存，污染 L2。最彻底的办法是切到纯命令行模式，关掉图形界面。

最后查代码。确认你的计时区间里没有混入其他操作。我见过有人在计时循环里加了printf调试，结果 IO 开销把信号全盖住了。计时区间里只放访存和clock64()，别的什么都不要放。

4.2 拐点不明显、找不到容量边界

如果延迟-容量曲线很平滑，没有明显拐点，通常是两个原因。

一是工作集步长太粗。如果你从 1MB 直接跳到 2MB，而真实容量是 1.5MB，拐点就被跳过去了。解决办法是把扫描步长细化到 128KB 甚至 64KB，尤其是在预期容量附近加密采样。

二是预取器在捣乱。GPU 的 L2 有硬件预取，如果你的访问模式太规律，预取器会提前把数据拉进来，掩盖真实的 miss 行为。破解方法是把指针追逐的步长随机化，或者用多个链表交替访问，让预取器无法预测。我一般会在链表构造时加入一个伪随机偏移，每个节点的步长在 128 到 256 字节之间随机，这样预取器就失效了。

4.3 冲突实验测不出关联度

冲突实验测不出台阶，延迟一直很平稳，说明你的“冲突地址”根本没有冲突。检查两点。

第一，set 索引的计算是否正确。如果你假设的 S 是错的，那么 (S × 128) 这个步长就不会让地址落到同一个 set。建议先用一个你确定知道 set 数量的 Cache（比如 L1）来验证你的地址生成公式。

第二，地址是否真的在物理上冲突。前面说过大页的问题，如果两个虚拟地址在同一个 set，但映射到了不同的物理页，物理层面可能并不冲突。解决办法是确保所有测试地址在同一个大页内，或者用cudaMemAdvise之类的接口固定物理页。

4.4 常见问题速查表

4.5 几个只有踩过才知道的坑

第一个坑：clock64()的读取本身有开销，大概几个周期。如果你测的是很短的访问序列，这个开销会占比很大。解决办法是让每次测量至少走几千个周期，把固定开销摊薄。

第二个坑：GPU 的 L2 在 Ampere 上是分区的，不同 SM 访问同一个地址可能走不同的路径。如果你用多个 SM 同时测，测出来的延迟是混合值。建议所有测量都固定在一个 SM 上做，用__smid之类的寄存器确认。

第三个坑：显存的刷新和 ECC 会影响延迟。如果开了 ECC，每次 miss 到显存都要做校验，延迟会增加。做逆向时最好关掉 ECC，或者至少在报告结果时注明 ECC 状态。

第四个坑：驱动版本会影响 L2 的行为。我对比过两个相邻版本的驱动，同一个 kernel 的 L2 命中率差了 8%。所以如果你的实验结果和别人的对不上，先确认驱动版本是否一致。

5. 逆向结果的实际应用与性能调优建议

5.1 把测出来的参数用起来

测完这些参数，最终目的是指导 kernel 优化。举几个直接的应用场景。

如果你测出 L2 是 16 way、set 索引有哈希，那么在设计数据布局时，就要避免让热点数据集中在少数几个 set 上。具体做法是给数组的每一行加一个 padding，让行与行之间的起始地址错开，不要形成固定的 stride 冲突。这个 padding 的大小可以根据你测出的主冲突步长来定，通常是 128 字节的奇数倍。

如果你测出替换策略是伪 LRU，那么在循环访问多个数组时，要注意访问顺序。伪 LRU 对“循环扫描”这种模式特别不友好，因为它会把整个 set 都标记为“最近使用过”，导致真正需要保留的数据被踢掉。解决办法是分块处理，让每个块的工作集能完整放进 L2，而不是一遍遍扫描整个大数组。

5.2 带宽与延迟的权衡

L2 逆向的另一个产出是你能更准确地估算带宽瓶颈。知道了 sector 大小和关联度，你就能算出在给定访问模式下，L2 能提供的有效带宽上限。如果实测带宽远低于这个上限，说明瓶颈不在 L2，而在更下游的显存或者上游的 SM 发射。

我一般会做一个简单的 roofline 分析：横轴是计算强度，纵轴是性能。把 L2 带宽作为一条水平线画上去，如果你的 kernel 落在 L2 带宽线以下，那优化方向就是减少 L2 访问次数，比如通过共享内存做数据复用。

5.3 后续可以扩展的方向

这套方法不只适用于 L2。同样的思路可以往上扩展到 L1 和共享内存，往下扩展到显存控制器和 DRAM 的 bank 结构。如果你有多个不同架构的 GPU，还可以做横向对比，看看不同代际之间 L2 的设计差异。

另一个方向是把逆向出来的参数写进模拟器，比如一些开源的 GPU 模拟器，用真实参数校准之后，模拟结果的参考价值会高很多。这个工作我做过一部分，校准后的模拟器在访存密集型负载上的预测误差从 30% 降到了 10% 以内。

我个人在实际操作中的体会是，逆向 Cache 这件事，最难的不是写代码，而是设计出能“问出正确问题”的实验。一个好的实验设计，能让硬件自己把答案告诉你；而一个粗糙的实验，只会给你一堆无法解释的噪声。每次动手前，先想清楚你要测什么、预期看到什么、如果看不到可能是什么原因，这三问能帮你省掉大量返工时间。