💻 computer science

Non-Uniform L2 Cache Latency Across the Streaming Multiprocessors of an NVIDIA L40

This paper reveals that the L2 cache hit latency on NVIDIA L40 and Blackwell GPUs is non-uniform and highly dependent on the specific physical Streaming Multiprocessor (SM) issuing the load, enabling a stable, user-level primitive for kernel self-localization, device fingerprinting, and optimized work distribution that reduces makespan by up to 11%.

Original authors: Faruk Alpay, Baris Basaran

Published 2026-06-23

📖 5 min read🧠 Deep dive

CC BY 4.0

Original authors: Faruk Alpay, Baris Basaran

Original paper licensed under CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ This is an AI-generated explanation of the paper below. It is not written or endorsed by the authors. For technical accuracy, refer to the original paper. Read full disclaimer

Imagine a massive, high-speed library (the GPU) where thousands of librarians (called SMs, or Streaming Multiprocessors) work together to fetch books (data) from a giant, shared storage room (the L2 Cache).

For years, computer scientists believed this library worked like a perfectly uniform system: no matter which librarian you asked, the time it took to grab a book from the storage room was exactly the same. They thought the storage room was a single, flat pool of information where distance didn't matter.

This paper, however, reveals that this belief is wrong. The authors discovered that on the NVIDIA L40 GPU, where a librarian stands physically determines how fast they can get a book.

Here is the breakdown of their findings using simple analogies:

1. The "Distance Matters" Discovery

Think of the storage room not as a flat floor, but as a large warehouse with aisles.

The Old View: Everyone thought every librarian stood exactly the same distance from the shelves. If you asked Librarian A or Librarian Z, the wait time was identical (about 279 "ticks" of the clock).
The New Reality: The authors measured the wait time for every single one of the 142 librarians. They found that some librarians are standing right next to the shelves and get books in 222 ticks. Others are standing at the far end of the warehouse and take 339 ticks.
The Result: That's a 52% difference in speed. It's like one person running to the front door while another has to run to the back exit just to get the same package.

2. The "Mirror Image" Pattern

When the authors mapped out who was fast and who was slow, they didn't see random noise. They saw a perfect pattern.

The 142 librarians are split into two identical halves (like two wings of a building).
Librarian #1 in the first wing has the exact same speed profile as Librarian #1 in the second wing.
This "mirror symmetry" matches the actual physical blueprint of the computer chip, proving this isn't a glitch in the software—it's a fact of the hardware's physical construction.

3. The "Fingerprint" Effect

Because every librarian has a unique speed based on their physical spot, the library has a secret identity.

Self-Location: If you are a librarian (a computer program) and you ask, "How long does it take me to get a book?", you can instantly figure out exactly which spot in the warehouse you are standing in. The authors built a tool that can identify your specific spot with 99% accuracy.
Device Fingerprinting: Even if you have two brand-new, identical L40 GPUs (two identical libraries), they are slightly different. One library might have Librarian #5 standing slightly closer to the shelves than the other library's Librarian #5 due to tiny manufacturing differences. The authors could tell the two identical machines apart 100% of the time just by measuring these tiny speed differences. It's like being able to tell two identical twins apart by the way they walk.

4. Is This a Security Risk?

The authors are careful to clarify what this means for security.

What it IS: A way for a program to know where it is inside the computer. It's like a person walking into a room and realizing, "Oh, I'm standing in the corner near the window."
What it IS NOT: It is not a way to steal secrets from other programs. The authors emphasize that this only measures the program's own speed. It cannot peek at what other programs are doing or steal their data. It is a "self-localization" tool, not a "spy" tool.

5. The Practical Benefit: Smarter Scheduling

Why does this matter for performance?

Imagine you have a list of 100 tasks to do.
The Old Way: You assign tasks randomly. Some go to the librarians standing far away (slow), and some go to those standing close (fast). The whole job waits for the slowest people to finish.
The New Way: Now that we have the map, we can assign the heavy lifting to the librarians standing closest to the shelves.
The Result: By doing this, the authors showed they could finish the work 11% faster for tasks that rely heavily on the cache. However, if the task requires fetching data from the main memory (a different, slower storage), this trick doesn't help.

6. It's Not Just One Chip

The authors tested this on a newer, different chip (the RTX 5090) as well. They found the same "distance matters" rule applies there too, though the pattern is slightly different. This proves that the old idea of a "uniform" cache is likely wrong for many modern high-end computers.

Summary

The paper shatters the illusion that all parts of a GPU's cache are equal. It reveals that physical location dictates speed. By mapping these hidden speed differences, we can:

Identify exactly where a program is running.
Distinguish between two identical machines.
Speed up specific tasks by assigning them to the fastest physical spots.

It turns out the computer chip isn't a flat, uniform field; it's a landscape with hills and valleys, and knowing the map makes you faster.

Technical Summary: Non-Uniform L2 Cache Latency on NVIDIA L40

Problem Statement
The NVIDIA L40 GPU, based on the full AD102 implementation, exposes a 96 MiB L2 cache that is conventionally modeled in performance tools and academic literature as a uniform pool with a single, device-wide hit latency. This abstraction assumes that any Streaming Multiprocessor (SM) accessing the L2 cache incurs the same cost. However, the physical reality of the L2 cache involves multiple slices attached to memory controllers, accessed via an on-chip network (NoC). While prior work on V100, A100, and H100 architectures established that NoC latency is non-uniform, those studies emphasized that per-core average latency was approximately uniform, with non-uniformity appearing primarily as discrete steps across partition boundaries. It remained an open question whether a single-die, single-partition Ada Lovelace data-center GPU (like the L40) exhibits a comparable structure, and if so, whether the per-SM average latency is uniform or placement-dependent.

Methodology
The authors developed a user-level microbenchmarking methodology to measure L2-hit latency from every physical SM without requiring elevated privileges or hardware performance counters.

Probe Design: The measurement uses a single-thread, single-warp dependent load chain (ld.global.cg) to ensure a single round-trip through the cache fabric. Two buffer shapes were employed: a "single-line" shape to probe specific 128 B slices and a "chain" shape (2 MiB random permutation) to measure average distance to all slices.
SM Isolation: To eliminate on-chip contention, blocks are serialized using a global turn counter. Each block claims a turn via atomicAdd, waits, performs warm-up, and times its dependent loads in isolation.
Coverage: The probe covers all 142 physical SMs on the L40 and 256 slice probes. A similar probe was executed on an NVIDIA RTX 5090 (Blackwell GB202) to test generalization.
Modeling: The latency data is analyzed using an additive Non-Uniform Cache Access (NUCA) model: $L = \mu + a(sm) + b(slice)$ , where $a(sm)$ represents the SM placement term and $b(slice)$ represents the slice placement term.

Key Results

Significant Latency Variation: Contrary to the uniform model, L2-hit latency on the L40 is not a constant. It spans 222.5 to 339.2 cycles (a 52% range), with a mean variation of 57.2 cycles across SMs. The noise floor is negligible (<0.01 cycles), confirming the variation is physical.
Two-Fold Symmetry: The SM placement term exhibits a strong two-fold symmetry. Splitting the 142 SMs at identifier 72 reveals two halves (72 and 70 SMs) with nearly identical latency profiles (correlation $r=0.999$ ). This aligns with the AD102 floorplan of two symmetric six-GPC clusters.
Continuous Gradient vs. Discrete Steps: Unlike the A100/H100, which show a discrete latency step at the partition boundary, the L40 exhibits a continuous gradient of latency across SMs with no bimodal gap.
Additive Model Fit: The additive model explains 87% of the latency variation ( $R^2=0.87$ ). Adding a single rank-1 interaction term increases this to 0.98, indicating the latency map is a low-dimensional geometry following the GPC layout.
Generalization and Specificity:
- Cross-Architecture: The RTX 5090 also shows non-uniformity (46% spread), but with a weaker symmetry ( $r=0.80$ ) and slower absolute latency.
- Device Fingerprinting: Two physically distinct, identical-model L40s are distinguishable at 100% accuracy based on their per-SM maps, despite having nearly identical mean latencies (0.28-cycle difference). The per-SM map acts as a stable hardware identity.
Stability: The per-SM latency map remains unchanged after one hour of 100% utilization on both devices, confirming the topology is a property of the silicon rather than a transient state.

Contributions and Significance
The paper claims the following contributions:

Topology Mapping: It provides the first per-SM L2-latency map for the L40, revealing a 52% spread that invalidates the uniform cache assumption for latency-bound workloads.
Self-Localization Primitive: The authors demonstrate that a single L2-latency probe allows a kernel to identify its own physical SM placement with 99.2% accuracy (or 92% across a mixed device set) and identify the specific device die with 100% accuracy. This is framed as a "self-localization" primitive, not an attack on a victim kernel, as it extracts no secret data and relies on the kernel's own placement.
Performance Implications: By distributing latency-bound work according to the per-SM map (placing work on faster SMs), the authors achieved a 10.9% reduction in makespan for L2-resident workloads. This benefit vanishes in DRAM-bound regimes, highlighting the regime-specific nature of the topology's impact.
Reinterpretation of Residency: The paper reinterprets existing single-thread residency controls (capacity boundaries, line-tag scaling, and persisting-L2 results) as averages over this non-uniform distribution, rather than constants.

Scope and Limitations
The authors explicitly limit the scope of their security claims. The technique enables self-localization (a kernel reading its own placement) and device fingerprinting. It does not infer the placement of a separate, co-resident "victim" kernel, nor does it extract secret data; such capabilities would require a contention or interference channel not measured in this study. The study is limited to two L40s and one RTX 5090, and the mapping of %smid to the physical floorplan is inferred from symmetry and GPC counts rather than independently confirmed by die photography. The absolute cycle counts are specific to the microbenchmark design and do not directly translate to application throughput where independent requests and L1 usage occur.