1. 项目缘起与整体思路拆解1.1 为什么要对 L2 Cache 做逆向分析GPU 性能调优做到一定程度必然会撞上一堵墙SM 内部的寄存器、共享内存、L1 这些资源厂商文档给得还算清楚但一旦跨出 SM进入 L2 Cache 和显存子系统公开资料就变得非常稀薄。NVIDIA 从 Volta 之后基本不再公开 L2 的替换策略、分区映射、sector 粒度这些底层细节而这些东西恰恰决定了你的 kernel 在访存密集型场景下到底能跑多快。我最初做这个项目的动机很直接手上有一个访存模式比较特殊的 kernelprofiler 显示 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 sliceslice 之间是否独立。这直接影响 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 cycleL2 命中约 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::vectorint perm(N); std::iota(perm.begin(), perm.end(), 0); std::shuffle(perm.begin(), perm.end(), std::mt19937(42)); std::vectorNode 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 和随机替换在“访问 N1 个同 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 配置和地址布局命中率提升明显。整个过程最深的体会是黑盒不可怕可怕的是没有系统性的测量方法。只要实验设计得当外部可观测的时序数据足以反推出足够多的内部结构信息支撑起实际的性能优化决策。 
长河君 · 2026/10/18 17:31:30 · 阅读 1.2k GPU L2 Cache逆向工程：Ampere微基准测试与性能调优实战
1. 项目缘起与整体思路拆解1.1 为什么要对 L2 Cache 做逆向分析GPU 性能调优做到一定程度必然会撞上一堵墙SM 内部的寄存器、共享内存、L1 这些资源厂商文档给得还算清楚但一旦跨出 SM进入 L2 Cache 和显存子系统公开资料就变得非常稀薄。NVIDIA 从 Volta 之后基本不再公开 L2 的替换策略、分区映射、sector 粒度这些底层细节而这些东西恰恰决定了你的 kernel 在访存密集型场景下到底能跑多快。我最初做这个项目的动机很直接手上有一个访存模式比较特殊的 kernelprofiler 显示 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 sliceslice 之间是否独立。这直接影响 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 cycleL2 命中约 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::vectorint perm(N); std::iota(perm.begin(), perm.end(), 0); std::shuffle(perm.begin(), perm.end(), std::mt19937(42)); std::vectorNode 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 和随机替换在“访问 N1 个同 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 配置和地址布局命中率提升明显。整个过程最深的体会是黑盒不可怕可怕的是没有系统性的测量方法。只要实验设计得当外部可观测的时序数据足以反推出足够多的内部结构信息支撑起实际的性能优化决策。