Non-Uniform L2 Cache Latency Across the Streaming Multiprocessors of an NVIDIA L40 Faruk Alpay*
Barış Başaran
arXiv:2606.22588v1 [cs.AR] 21 Jun 2026
Department of Computer Engineering, Bahçeşehir University Istanbul, Turkey {faruk.alpay, baris.basaran}@bahcesehir.edu.tr
Abstract The NVIDIA L40 exposes a 96 MiB L2 cache usually modeled as one uniform pool with a single hit latency. We show this is wrong at the granularity a kernel sees: L2-hit latency depends strongly and reproducibly on which physical streaming multiprocessor (SM) issues the load. A turn-serialized, %smid-resolved probe maps the hit latency across all 142 SMs in one launch; it is not a constant near 279 cycles but spans 222–339 cycles (a 52% range), with per-repetition noise below 0.01 cycles. An additive model L = µ + a(sm) + b(slice) explains R2 = 0.87 (0.98 with one rank-1 term), and the SM term is two-fold symmetric—two halves of 72 SMs at correlation r = 0.999—following the AD102 GPC layout. Independent access patterns agree per SM at r = 1.000, so the effect is physical. The same probe on a Blackwell RTX 5090 shows it generalizes, while the per-die pattern is device-specific. Read as a fingerprint, a single user-level probe identifies the SM within a device at 92%, and two physically identical L40s are separated at 100% despite near-identical mean latency (per-SM map r = 0.63)—a per-die hardware identity, not a clock artifact. This is a self-localization and fingerprinting primitive: a kernel reads its own placement and device, not a victim’s, and extracts no secret data. The map is stable, unchanged after an hour at full utilization on both devices. As a consequence, distributing latency-bound work by the map cuts makespan by up to 11%. Single-thread capacity, line-tag, prefetch-modifier, and persisting-L2 results appear as controls. The artifact contains seeds, raw observations, the trained model, and regeneration scripts.
1
Introduction
The NVIDIA L40 is a data-center Ada Lovelace GPU built from the full AD102 implementation. CUDA reports a 96 MiB L2 cache and a single device-wide L2; performance models and teaching material routinely treat that cache as a uniform pool with one access latency. The interface gives no reason to expect otherwise: there is one l2CacheSize, one persisting set-aside, and one nominal hit cost. This abstraction hides a physical fact. The L2 is not a monolith; it is a set of slices attached to memory controllers, reached from the SMs through an on-chip network (NoC). Distance through that fabric is not free. Prior measurement work on real GPUs established that NoC latency between cores and memory partitions is non-uniform and depends on physical placement, with up to roughly 70% variation across V100, A100, and H100 [4]. That study did not cover any Ada GPU, and it emphasized that on those parts the per-core average latency is approximately uniform while the non-uniformity appears only between specific core–slice pairs, with an additional discrete step when traffic crosses the A100/H100 dual-partition ∗
Corresponding author: [email protected]
1
boundary. Whether a single-die, single-partition Ada data-center GPU exhibits a comparable structure—and what form it takes without the partition split—was open. We answer this for the L40 by measuring L2-hit latency from every physical SM. The central question is: is the application-visible L2-hit latency on this GPU a single number, or a function of physical SM placement, and if so what is its structure? The answer is that placement governs the hit latency. We map the full SM×slice latency space, fit a compact additive model, recover a two-fold symmetry that matches the AD102 floorplan, confirm the effect against an independent access pattern, and quantify the resulting timing leakage. The methodology follows established GPU microbenchmarking [7, 3, 6, 4]; the contribution is the L40-specific topology and its consequences, including a reinterpretation of single-thread residency numbers that look like clean uniform-cache results until placement is taken into account. The contributions are: 1. A per-SM L2-hit latency map of one L40 covering all 142 SMs and 256 slice probes, showing a 222.5–339.2-cycle (52%) spread where a uniform model predicts a constant. 2. A compact additive model L(sm, slice) = µ + a(sm) + b(slice) (R2 = 0.87) that separates an SM-placement term (57.2 cycles) from a slice-placement term (39.5 cycles), and a near-low-rank refinement (one rank-1 interaction reaches R2 = 0.98) showing the map is a two-coordinate geometry following the GPC layout. 3. A two-fold symmetry of the SM term: two halves of 72 SMs with profiles correlated at r = 0.999, consistent with the AD102 two-cluster GPC layout, together with evidence that the L40 lacks the discrete inter-partition latency step of A100/H100 but, unlike those parts, has a non-uniform per-SM average. 4. A robustness argument: single-line slice probes and a 2 MiB random chain agree at r = 1.000 per SM, and the symmetry pairs rule out a measurement-order confound; per-rep noise is below 0.01 cycles. 5. A side-channel quantification and a trained placement oracle: a single L2-latency probe localizes the executing SM to one of roughly 70–120 classes (≈6–7 bits), and a published classifier identifies the exact SM (one of 142) on 99.2% of held-out fingerprints. 6. Evidence that the effect generalizes: the same probe on a Blackwell RTX 5090 finds a 46% spread and an absolutely slower L2, and shows the learned model is device-specific (chance accuracy across devices). 7. Same-model device fingerprinting: two physically distinct, identical-model L40s are separated at 100% from their per-SM signature despite a 0.28-cycle mean-latency difference (per-SM map r = 0.63), and the per-SM oracle does not transfer between dies (0% vs 98.6% native)—each die has a stable hardware identity. 8. A user-level self-localization and physical-location primitive: from one fingerprint a kernel identifies its own SM within a device at 92% and its device, using no privileges or counters. It is self-localization, not victim or co-resident inference, and extracts no secret data. 9. A stability result: the per-SM map is unchanged after an hour at 100% utilization on both devices (correlation 1.000, drift below 0.4 cycles), so the topology is a property of the silicon. 10. A measured consequence: scheduling latency-bound work by the per-SM map reduces makespan by up to 11%, with no benefit in the DRAM-bound regime.
2
11. Single-thread residency controls—capacity boundary, line-tag scaling, a PTX prefetchmodifier null result, and a persisting-L2 boundary—now read as averages over the nonuniform latency they each assume to be a constant.
2
Method: Per-SM L2 Latency Probing
Table 1 summarizes the tested instance. The GPU was accessed through SSH on a Vast.ai CUDA 13.0 base image. The container runs without elevated privileges, so hardware performance counters were not available through Nsight Compute on this host [10]; the artifact relies on timing, CUDA device properties, nvidia-smi telemetry, source hashes, command lines, and repeated raw observations. Table 1: System under test. Values are recorded in the run manifests.
Component
Observed property
GPU SMs Global memory L2 cache Persisting L2 maximum Access-policy max window Shared memory per SM Memory interface Reported SM clock ECC Driver / runtime Host CPU
NVIDIA L40, compute capability 8.9 142 47 665 709 056 B visible (44.39 GiB) 100 663 296 B (96 MiB) 69 206 016 B (66 MiB) 134 213 632 B (127.996 MiB) 102 400 B (100 KiB) 384 bit GDDR6, 9001 MHz 2490 MHz Enabled 580.142 / CUDA 13.0 Intel Xeon Gold 5418Y, 48 visible cores
Probe. Each measurement times a single-thread, single-warp dependent load chain, following the standard one-request-in-flight design that removes coalescing and queueing so the timed value is one round trip through the cache fabric [4, 7]. The load is ld.global.cg, which uses the global cache path and bypasses the private L1, so a hit is an L2 hit. Two buffer shapes are used. The single-line shape fills a buffer with buf[i] = i, so any start word loops forever on its own 128 B line; repeated loads of that one line are L2 hits to the single slice that owns the address, and sweeping the byte offset moves the probe across slices. The chain shape is a random permutation pointer chain over a 2 MiB footprint that touches the whole slice ensemble, so its per-SM latency is the average distance to all slices. For footprint F and stride S in the residency controls (Section 9), the number of distinct 128 B line tags touched is T (F, S) = F/128 for S < 128 and F/S otherwise, with tag-equivalent footprint Ftag = 128 T = F min(1, 128/S). SM attribution and isolation. Each block reads its physical SM through the %smid special register and writes its measurement indexed by that identifier. One block is forced resident per SM by requesting more than half of the 100 KiB shared memory per SM as dynamic shared memory; the occupancy calculator then reports one block per SM, so a launch of 142 blocks covers every SM exactly once. To remove on-chip contention during timing, the blocks are serialized by a global turn counter: a block claims a turn with atomicAdd, waits with __nanosleep back-off until its turn is served, performs warm-up and then times its dependent loads alone, and releases the turn. Only one block touches memory in its timed region, so the measured cycles isolate that SM’s path. The reported latency is 3
end cycles − begin cycles , A with A timed loads per block; the slice campaign uses A = 8192, four repetitions, and 256 consecutive 128 B probes; the chain control uses A = 8192 and 16 repetitions. L(sm, slice) =
Model. We summarize the map with an additive decomposition, in the spirit of non-uniform cache access (NUCA) models for wire-delay-dominated caches [5]. Definition 1 (Additive NUCA model). For measured latencies L(sm, slice) with grand mean µ, b = µ + a(sm) + let a(sm) = L(sm, ·) − µ and b(slice) = L(·, slice) − µ. The additive model is L b is the SM×slice interaction. b(slice), and the residual L − L The fraction of variation explained, the ranges of a and b, and the residual standard deviation are reported in Section 3; the interaction is the part of the latency that depends jointly on which SM reaches which slice.
3
The L2 Latency Topology
Figure 1 maps the mean L2-hit latency of every physical SM. Panel (a) plots latency against the SM identifier. The value is not constant: it ranges from 249.8 to 307.0 cycles when averaged over slices, and the curve has visible internal structure rather than noise. Panel (b) shows that this structure is two-fold symmetric. Splitting the SM-placement term a(sm) at identifier 72 yields two halves of 72 and 70 SMs whose profiles are nearly identical, with correlation r = 0.999 and a mean absolute difference of 0.99 cycles; the six fastest SMs occur in exact identifier pairs (49, 121), (59, 131), (61, 133) separated by 72, and the slowest pair as (4, 76) and (16, 88). Within each half the latency varies with position (correlation −0.53 with index), and across the chip the per-SM averages form a continuous distribution with no bimodal gap. The period of 72 matches six of the twelve AD102 GPCs at twelve SMs each, so the most direct reading is that the %smid enumeration walks two symmetric six-GPC clusters and that an SM’s latency is set by its position within a cluster.
two halves of 72 SMs (profile correlation r=0.999)
300
Per-SM latency term a(sm) (cycles)
Mean L2-hit latency (cycles)
(a) Latency by physical SM
290 280 270 260 250 0
20
40
60 80 100 Physical SM identifier
120
140
(b) Two-fold symmetry of the SM term
30
half 1 (SM 0..71) half 2 (SM 72..141)
20 10 0 10 20 30
0
10
20
30 40 50 Position within half
60
70
Figure 1: Per-SM L2-hit latency on one L40. (a) Mean latency against physical SM identifier; the dashed line marks the split into two halves of 72 SMs. (b) The SM-placement term a(sm) for the two halves overlaid; the profiles agree at correlation r = 0.999, a two-fold symmetry consistent with the AD102 two-cluster GPC layout.
Figure 2 shows the full latency map over the SM×slice space, with SMs ordered by mean latency and slice probes ordered by slice term. The surface is smooth in both directions: latency 4
rises from 222.5 cycles at the near SM / near slice corner to 339.2 cycles at the far corner, a 116.6-cycle (52%) total range, equivalently 89.4 ns to 136.2 ns at the reported 2.49 GHz. The vertical striping is the slice structure: consecutive 128 B probes alternate among slices, and the slice-term autocorrelation has its first strong period at four lines (512 B), the granularity at which the L2 interleaves addresses across slices on this part. The two visible axes are exactly the two additive terms; the fine texture on top of them is the interaction.
L2-hit latency L(sm, slice) on one L40
120
320
100
300
Latency (cycles)
Physical SM (sorted by mean latency)
140
80
280
60
260
40
240
20 0
0
50 100 150 200 L2 slice probe (sorted by slice latency term)
250
Figure 2: L2-hit latency L(sm, slice) over all 142 SMs and 256 slice probes. SMs are ordered by mean latency and probes by slice term. Latency rises smoothly from 222.5 to 339.2 cycles; the vertical striping is the per-slice term with a 512 B interleave period.
Figure 3 quantifies the decomposition. Panel (a) plots measured latency against the additive prediction µ + a + b; the model explains R2 = 0.87 of the variation, with an SM term spanning 57.2 cycles, a slice term spanning 39.5 cycles, and a residual standard deviation of 6.85 cycles for the interaction. Panel (b) shows the per-SM latency spectrum: the 142 SMs spread continuously over 57 cycles with no partition gap. This is the first point of contrast with A100/H100, where the dominant non-uniformity is a discrete step at the partition boundary and the per-core average is approximately uniform [4]. On the single-partition L40 there is no such step, but the per-SM average is itself non-uniform by 23% (the chain control independently measures a 57.1-cycle per-SM spread), so placement matters here even for an access pattern that touches the whole cache evenly.
5
(a) Additive NUCA model (R 2 = 0.87)
(b) Per-SM latency spectrum
300 280 260 240 220
142 SMs span 57 cycles; 73 levels separable at 0.5 cyc
10
320 Number of SMs
Measured latency (cycles)
340
8 6 4 2
220
240 260 280 300 320 Additive model + a(sm) + b(slice) (cycles)
0
340
250
260 270 280 290 300 Mean L2-hit latency per SM (cycles)
Figure 3: (a) Additive model µ + a(sm) + b(slice) against measured latency (R2 = 0.87); the diagonal banding is the SM×slice interaction. (b) Per-SM latency spectrum: a continuous 57-cycle spread with no discrete partition step.
Two finer structures sharpen this picture. First, the map is nearly low-rank. Adding a single rank-1 interaction term c u(sm) v(slice) to the additive model raises R2 from 0.87 to 0.98 on the L40 (and from 0.83 to 0.99 on the RTX 5090 of Section 5). The interaction’s SM coordinate u is essentially uncorrelated with the additive term a (|r| = 0.06), so it is a second, independent placement axis rather than a rescaling of the first; two coordinates per SM and per slice thus reproduce the latency map almost completely. Second, the SM term is hierarchical: its autocorrelation peaks at a period of 12 SMs on the L40 (and 10 on the RTX 5090), matching the per-GPC SM counts of the two floorplans, nested inside the two-fold macro symmetry of Figure 1. The non-uniformity is therefore not noise on a flat cache but a structured, low-dimensional geometry that follows the GPC layout. The effect is physical, not procedural. Three checks rule out artifacts. First, perrepetition noise is small: the median standard deviation of the per-access latency across repetitions is 0.006 cycles, so a difference of tens of cycles between SMs is four orders of magnitude above the noise floor. Second, the two-fold symmetry decouples placement from measurement order: SMs k and k + 72 are timed roughly 72 turns apart yet differ by 0.99 cycles on average, whereas an order- or temperature-driven drift would scale with turn separation. Third, the single-line slice probe and the 2 MiB random chain are different access patterns measured in independent launches with independent turn orders, and their per-SM averages correlate at r = 1.000; within a fixed SM, latency is uncorrelated with turn across repetitions (mean correlation −0.13). The per-SM latency is therefore a stable function of the physical SM.
4
Timing Leakage of Physical Placement
Because the per-SM latency is stable and the measurement noise is small, the latency itself identifies the SM. We state this as a separability bound. Proposition 1 (SM separability). Let the per-SM mean latencies {L(sm)} have single-probe standard deviation σ. The number of SM classes distinguishable at confidence level kσ is the largest C such that the sorted per-SM means admit C − 1 consecutive gaps exceeding kσ. With σ ≤ 0.01 cycles and the measured spread of 57.2 cycles, C ≥ 118 at k = 5. The bound follows directly from ordering the per-SM means and counting gaps wider than kσ; with the measured σ almost every pair of SMs is separable. A coarser, more conservative count that bins the means at 0.5-cycle resolution still yields 73 distinct levels. Either way, a single 6
L2-latency probe localizes the executing SM to one of roughly 70–120 classes, about 6–7 bits of its own physical placement. This is the L40 analogue of the core-location timing channel that motivates real-GPU NoC analysis [4] and connects to practical GPU side channels [8]: a user-level kernel can infer where it is physically scheduled, and a scheduler that wanted to hide placement would have to defeat a 57-cycle signal that sits far above the noise. What is measured is self -localization—a kernel reading its own placement; inferring the placement of a separate, co-resident victim kernel would require a contention or interference channel that we do not measure, and no secret data are recovered. This is a localization primitive—a stepping stone toward de-anonymization or co-residency detection [1], not a complete exploit. The same signal is usable constructively, as a cheap placement oracle for locality-aware scheduling.
4.1
A trained placement oracle
To make the bound concrete and reusable, we train a model that predicts the physical SM from a measured fingerprint. The fingerprint is a vector of L2-hit latencies to 32 fixed 128 B lines spaced 256 B apart, each timed with A dependent loads; a single launch of the serialized probe yields one fingerprint per SM, so a “shot” is one fingerprint for each of the 142 SMs. We collect 500 shots at each of A ∈ {32, 64, 128, 256} loads, label each fingerprint with its %smid, and split train/test by shot so the test fingerprints come from shots never seen in training. A random-forest classifier is trained per condition; the published model uses A = 256 and all 32 probes. With a 256 load fingerprint the oracle identifies the exact SM, one of 142, on 99.2% of held-out shots, with the correct SM in its top five every time and the two-cluster label recovered at the same 99.2% against a 0.7% chance rate; even a fast 32 load fingerprint reaches 96.3%. A nearest-centroid baseline already reaches 98.9%, so the signal, not the model, carries the leakage. The leakage does not need a long fingerprint either: a single probe localizes the SM to 75.6% exact and 90.6% cluster accuracy, rising to 99.2% with all 32 probes. The remaining errors are almost entirely the symmetric twins: on the held-out sample the only misclassification maps an SM to its k ↔ k + 72 partner, which Section 3 showed are near-identical. The oracle therefore fails exactly where the fabric is symmetric, and its error mode re-derives the two-fold structure. We publish the trained model, its probe-bank metadata, and a held-out sample so the oracle runs offline without a GPU. The oracle performs self -localization: it reports the physical SM of the kernel that took the fingerprint. It is not an attack on a separate victim—reading another kernel’s placement would need a contention channel we do not measure—and it extracts no secret data. The same primitive is usable constructively, as a cheap placement estimate for locality-aware scheduling (Section 7).
5
Cross-Architecture Generalization
To test whether the non-uniformity is specific to the L40, we ran the identical probe on an NVIDIA RTX 5090 (Blackwell GB202, 170 SMs, the same nominal 96 MiB L2). Table 2 compares the two devices. The phenomenon generalizes: the 5090 L2-hit latency also depends on physical SM placement, spanning 288.2–419.5 cycles (46%), with an additive model at R2 = 0.83 and the same 512 B slice-interleave period. Two architectural differences stand out. First, the 5090 L2 is markedly slower in absolute time—119.7 ns–174.3 ns versus the L40’s 89.4 ns–136.2 ns—despite the identical nominal capacity, which is consistent with the larger die and higher SM count lengthening the average SM-to-slice path. Second, the L40’s two-fold symmetry is near-perfect (r = 0.999) while the 5090’s is weaker (r = 0.80 at a split of 88 SMs), indicating a less symmetric arrangement of SMs around the fabric on GB202.
7
Table 2: The per-SM L2 latency topology on two architectures, measured with the identical probe. The SM/slice terms are the additive-model term ranges in cycles. Device
SMs
L40 (Ada AD102) RTX 5090 (Blackwell GB202)
142 89.4–136.2 170 119.7–174.3
(a) Per-SM L2 latency, two architectures
160 150 140
Number of SMs
Mean L2-hit latency (ns)
L2-hit (ns)
L40 (Ada, 142 SMs) RTX 5090 (Blackwell, 170 SMs)
130 120 110 100
0.0
0.2 0.4 0.6 0.8 Physical SM (normalized index)
1.0
R2
SM/slice (cyc)
2-fold r (split)
52% 0.87 46% 0.83
57 / 40 64 / 46
0.999 (72) 0.80 (88)
Spread
16 14 12 10 8 6 4 2 0
(b) Per-SM latency spectra L40 (Ada) RTX 5090
100
110
120 130 140 150 Mean L2-hit latency per SM (ns)
160
Figure 4: Per-SM L2-hit latency on the Ada L40 and the Blackwell RTX 5090, measured with the same probe. (a) Latency by physical SM (normalized index); both are non-uniform, and the 5090 sits higher. (b) Per-SM latency spectra in nanoseconds; the two devices occupy disjoint bands.
The fingerprint is therefore architecture-specific. The L40-trained oracle of Section 4, applied to 5090 fingerprints, drops to 0.6% accuracy—exactly chance for 170 classes—while a 5090-native oracle trained by the same procedure reaches 86.2%. The method transfers across architectures; the learned model does not. This confirms that the latency map is a genuine device signature and that the central finding is not an L40 artifact.
6
Device Fingerprinting and Physical-Location Inference
Because the per-SM map is specific to each physical device, it acts as a fingerprint. We separate two questions: which device a kernel runs on, and where on it (which SM).
6.1
Same-model device fingerprinting
The cross-architecture case is easy: the L40 and RTX 5090 latency bands are disjoint (Section 5), so a fingerprint names the architecture trivially. The substantive question is whether two identical-model devices are distinguishable. We measured a second physical L40—distinct UUID, different driver—with the same probe. The two L40s have nearly identical global mean latency (277.5 versus 277.8 cycles, a 0.28-cycle offset), so they cannot be told apart by average latency or clock. Yet their per-SM maps differ markedly: the maps correlate at only r = 0.63 (Figure 5), and after removing the global offset the per-SM difference has a standard deviation of 12.4 cycles (up to 37.7), far above the 0.006-cycle noise floor. A classifier separates the two devices at 100%, and stays at 100% after de-meaning each fingerprint, so the signal is a genuine per-die pattern rather than a global offset. The placement oracle is itself die-specific: the first-L40 oracle scores 0% on the second (below the 0.7% chance rate), while a second-L40-native oracle reaches 98.6%. Each physical L40 thus carries a stable, distinct per-SM L2 signature—a hardware identity—most plausibly from the combination of process
8
variation and the device’s particular disabled-SM pattern (each L40 enables 142 of the 144 AD102 SMs).
L40 #2 per-SM mean latency (cycles)
Per-SM map: two same-model L40s (r = 0.63) 300 290 280 270 260 250 250
260 270 280 290 300 L40 #1 per-SM mean latency (cycles)
Figure 5: Two physically distinct, identical-model L40s: their per-SM latency maps correlate at only r = 0.63 despite near-identical means, so each die has its own pattern. The two devices are separated at 100% from these fingerprints and the per-SM oracle does not transfer between dies (figures in text).
6.2
Physical-location inference
Pooling the first L40 and the RTX 5090 into a 312-location label space (142 + 170 SMs), one classifier on the shared 32-probe feature recovers the full physical location on 92.1% of held-out fingerprints against a 0.3% chance rate; a single probe already reaches 64.6%. Here the device split is the easy, cross-architecture case—the two devices occupy disjoint regions of fingerprint space (Figure 6)—and the substantive part is the within-device SM, identified at 92.1%. No privileges, performance counters, or topology APIs are used—the only input is the timing of ordinary global loads.
Fingerprint space (PCA) L40 (Ada) RTX 5090
60
PC 2 (cycles)
40 20 0 20 40 60
400
300
200
100 0 100 PC 1 (cycles)
200
300
400
Figure 6: Physical-location inference: the fingerprint space (first two principal components); the two devices occupy disjoint regions, each with an internal per-SM gradient. The 312-way location accuracy (64.6% from one probe, 92.1% from 32) is given in the text.
9
Scope. What is demonstrated is self -localization and device identification: a kernel reads its own physical placement and the identity of the device it runs on. This is a localization and fingerprinting primitive—a stepping stone toward de-anonymization or co-residency detection [1], not a complete exploit. We do not infer the placement of a separate victim kernel: that would require a contention or interference channel that we do not measure, and we extract no secret data. The same primitive is the security face of the property that drives the performance result of Section 7.
7
A Consequence: NUCA-Aware Work Placement
If L2-hit latency depends on SM placement, then distributing latency-bound work uniformly wastes time on the slow SMs—the GPU analogue of locality-aware data placement in distributed CPU caches [2]. We test this directly. A fixed amount of dependent-load work is spread across the SMs under three policies, and the wall-clock makespan is measured with CUDA events: oblivious (equal work per SM, no topology knowledge), aware (work per SM proportional to 1/a(sm) from the measured map), and dynamic (a global atomic work queue that balances itself at runtime without a model). The total work is identical across policies; only its distribution differs.
Makespan reduction vs oblivious (%)
+10.9%
10
NUCA-aware work placement aware (model, static) dynamic (work-stealing) +8.9%
8 6 4 2 0
+0.9%
L2-resident (1 warp/SM)
L2-resident (8 warps/SM)
DRAM-bound (27 GiB)
Figure 7: Makespan reduction over the oblivious baseline for a fixed latency-bound workload. The model-based static schedule (aware) matches or beats runtime work-stealing in the L2-resident regime and gives essentially nothing once the workload is DRAM-bandwidth bound.
Figure 7 reports the result. When the working set is L2-resident and the workload is latency-bound, the aware schedule reduces makespan by 10.9% with one warp per SM and 8.9% with eight, matching the maxi ai /HM(a) ratio predicted by the per-SM map. The model-based static schedule equals or slightly beats the dynamic work-stealing scheduler (7.3–8.7%) while incurring no runtime coordination. The gain is bounded to the regime where it should exist: when the per-thread footprint is enlarged to 27 GiB so the workload is DRAM-bandwidth bound, the aware schedule gives only 0.9%, because L2 latency no longer governs throughput. The non-uniform topology thus has a regime-specific scheduling consequence: it matters for latency-bound L2 reuse and not for streaming.
8
Stability Under Sustained Load
A fingerprint or a scheduling decision is only useful if the topology is stable. We held each device at full utilization for one hour—a 38 GiB buffer on the L40 and a 26 GiB buffer on the 10
RTX 5090, hammered by a compute-and-memory load kernel—and took a 32-probe per-SM snapshot every 60 s, 60 snapshots per device.
Per-SM latency under 1h load (cycles)
L40 map: idle vs sustained load (r = 1.000) 300 290 280 270 260 250 250
260
270 280 290 Idle per-SM latency (cycles)
300
Figure 8: The L40 per-SM map after one hour of full GPU and VRAM load matches the independently measured idle map at r = 1.000. Across the run the snapshot-to-snapshot correlation stays at 1.000 on both devices (in text).
The per-SM map is essentially invariant. On the L40 the snapshot-to-snapshot correlation has median 1.000 and the per-SM mean drifts by at most 0.08 cycles over the hour at 100% utilization and 63 ◦ C; the loaded map matches the independently measured idle map at r = 1.000 (Figure 8). The RTX 5090 behaves the same (correlation 1.000, drift 0.35 cycles, 100% utilization, 64 ◦ C, 402 W, 2857 MHz). The topology is therefore a property of the silicon, not of a transient idle state. One nuance bounds the fingerprint’s portability: an oracle trained on idle fingerprints transfers poorly to under-load fingerprints (8.5% on the L40), because the fine per-probe detail shifts with the operating condition even though the per-SM mean does not; a load-calibrated oracle recovers 91.4%. The signature persists under load, but a deployed locator should be calibrated to the operating point.
9
Single-Thread Residency Controls
The remaining experiments are single-thread residency measurements. Each one implicitly treats the L2-hit latency as one constant near 279 cycles; the topology result shows that this constant is the average of a 222.5–339.2-cycle distribution and, for a one-block probe, the value of whichever SM the scheduler happened to use. The residency conclusions survive this reinterpretation because they depend on the transition between the L2-hit and DRAM regimes, not on the absolute hit latency, but the framing matters: a “single hit latency” is a per-SM quantity. Figure 9 shows a capacity sweep with a 128 B stride. The median latency is flat through 32 MiB at 279.3 cycles, crosses a narrow transition near the nominal 96 MiB boundary (301.3 cycles at 92 MiB, 369.1 at 96 MiB, 537.2 at 100 MiB), and plateaus near 633 cycles in the DRAM-backed regime. This establishes the two regimes; it does not by itself identify the replacement variable, because a dense chain changes address span and unique tag count together.
11
Dependent-load latency (cycles)
650
default, stride 128 B
Nominal L2
600 550 500 450 400 350 300 0
25
50
75 100 125 Address-span footprint (MiB)
150
175
200
Figure 9: Dependent-load capacity sweep (128 B stride). The median rises from 301.3 cycles at 92 MiB to 631.9 cycles at 104 MiB, bracketing the nominal 96 MiB L2 boundary.
Varying the stride separates the two. Table 3 lists the interpolated transition midpoints. Raw address-span midpoints cover a 7.6× range (98.16 to 745.48 MiB), but normalizing to 128 B tag-equivalent footprint collapses them to 98.56 ± 3.46 MiB (coefficient of variation 3.5%), near the physical 96 MiB capacity. Replacement is governed by unique 128 B line tags, not by virtual address span—a consistency check against the line-tag model rather than a new result. Table 3: Capacity-transition midpoints for different strides.
Stride (B)
Raw midpoint (MiB)
Tag-equivalent midpoint (MiB)
Midpoint cycles
32 64 128 256 512 1024
102.46 103.05 98.16 197.26 383.49 745.48
102.46 103.05 98.16 98.63 95.87 93.18
457.99 454.98 459.18 460.97 460.08 443.80
PTX accepts load modifiers that request L2 prefetch sizes [11]. Compiling the same dependent-chain kernel with default ld.global.cg and with L2::64B, L2::128B, and L2::256B does not move the replacement boundary: Table 4 shows all four forms crossing at 98.1 MiB for 128 B stride and 206.0 MiB for 256 B stride, differing only by a few cycles in the high plateau. The governing capacity variable remains the unique line-tag count.
12
Table 4: Midpoint crossings for prefetch modifiers.
Load form
Stride
Midpoint (MiB)
High plateau (cycles)
default L2::64B L2::128B L2::256B default L2::64B L2::128B L2::256B
128 B 128 B 128 B 128 B 256 B 256 B 256 B 256 B
98.09 98.14 98.14 98.12 206.02 206.02 206.03 206.04
634.95 637.23 641.66 641.69 635.07 637.21 641.67 641.66
Finally, the CUDA persisting-L2 set-aside has a measurable boundary. CUDA reports a 66 MiB maximum set-aside on this L40 [9]. With a 256 MiB cold stream between warm-up and probe, Table 5 shows that persisting policy keeps 16–64 MiB hot sets in the L2-hit regime, partially protects 72 MiB, and gives no benefit for 80–88 MiB hot sets, matching the reported set-aside. This is a measured sufficient condition for a serialized hot/cold pattern, not a universal replacement guarantee. Table 5: Persisting-L2 hot-set boundary after a 256 MiB cold stream (128 B stride, 106 probe loads, seven repetitions).
10
Hot set (MiB)
Normal, cold 256 (cycles)
Persist, cold 256 (cycles)
Persist benefit (cycles)
Cold 0 baseline (cycles)
16 32 48 64 72 80 88
324.73 370.00 418.69 465.74 488.95 512.14 536.07
279.37 279.37 282.71 284.36 333.81 512.08 535.92
45.36 90.64 135.98 181.38 155.14 0.07 0.15
279.39 279.38 282.70 284.36 284.92 285.37 287.67
Sustained High-Utilization Controls
Latency microbenchmarks underutilize the GPU. To record the operating envelope of the same instance, a separate 40 GiB read-modify-write kernel reads and writes one float4 per element while varying fused multiply-add iterations from 1 to 1024; the working set occupies 41,395 of the 46,068 MiB reported by nvidia-smi. Figure 10 traces the transition from memory-bound to power-bound execution: low arithmetic intensity sustains a median 597.4 GB/s read+write rate at 100% reported utilization, and high intensity reaches 48.38 TFLOP/s as board power meets the 300 W limit. The instance reaches both its memory-bandwidth and board-power envelopes in the same session as the latency study.
13
1
4
(a) Bandwidth regime 16
(b) Compute regime
50
64
1024 512
256
Throughput (TFLOP/s)
Effective read+write bandwidth (GB/s)
600 550 500 450 400 350 300 250 200
512
40 256
30 20 10
1024
101 100 102 Arithmetic intensity (flop/byte)
0
64 1
4
16
101 100 102 Arithmetic intensity (flop/byte)
Figure 10: Sustained 40 GiB operating-regime sweep. Low arithmetic intensity is bandwidth-limited near 597 GB/s; high intensity becomes power-limited near 300 W and reaches 48.38 TFLOP/s.
11
Related Work and Discussion
GPU memory hierarchies have been reverse-engineered by microbenchmarking for several generations, establishing cache sizes, line sizes, latencies, and associativities for Volta, Ampere, Hopper, and earlier parts [7, 3, 6, 13, 12]. These studies report a representative L2-hit latency per device. Our result is that on the L40 a single representative number discards a 52% effect: the hit latency is a function of physical SM placement. The closest work is the real-GPU NoC characterization of Jin et al. [4], which measured core-to-partition latency on V100, A100, and H100 and found non-uniformity driven by physical placement, with a discrete step across the A100/H100 dual-partition boundary and an approximately uniform per-core average. The L40 differs on both counts. It has no discrete partition step—the per-SM averages form a continuous gradient—yet its per-SM average is non-uniform by 23%, which the multi-partition parts did not show. The two-fold symmetry at period 72 indicates a symmetric two-cluster GPC arrangement on the single AD102 die, and the 512 B slice-interleave period and 57/40 cycle SM/slice terms give a compact characterization of the fabric without access to performance counters. Together these extend real-GPU NoC measurement to the Ada generation and to a single-partition data-center GPU that the prior study did not include. The practical consequences are twofold. For performance, a latency-bound, reuse-heavy kernel pinned to a near SM sees its L2 hits roughly 57 cycles faster than the same kernel on a far SM; Section 7 turns this into a measured 11% makespan reduction for NUCA-aware work placement, and performance models that assume one L2 latency mispredict by up to 52% on latency-bound phases. The effect is not confined to the L40: Section 5 finds the same structure on a Blackwell RTX 5090, with a slower absolute L2 and a weaker two-fold symmetry. For security, the placement signal is a self-localization and fingerprinting primitive: Section 4 shows it leaks 6–7 bits about a kernel’s own physical scheduling (exact SM at 99.2%), and Section 6.1 shows two identical-model L40s are separable at 100% from their per-SM signature. This extends practical GPU side-channel concerns [8, 1] to placement and device inference on Ada, while remaining a stepping stone—not a complete exploit—as it reads the measuring kernel’s own placement and extracts no victim data.
12
Limitations
The study covers two physical L40s and one RTX 5090, with their own driver/runtime stacks and host platforms. Two L40s are enough to demonstrate same-model fingerprinting but not to characterize its distribution across a population of dies, and the cross-architecture comparison 14
rests on a single 5090. The security capability is scoped to self-localization and device fingerprinting: a kernel reads its own placement and device identity. Inferring the placement of a separate victim kernel, or extracting its data, would require a contention or interference channel we do not measure; the primitive is a stepping stone for de-anonymization or co-residency detection, not a complete exploit. The container does not permit privileged performance counters, so we do not report the measured slice-to-controller mapping, associativity, or internal routes; the topology is inferred from timing, its two-fold symmetry, the slice-interleave period, and cross-pattern agreement. The mapping of %smid to the physical die floorplan is consistent with the AD102 GPC count but is not independently confirmed by a die photograph. The latency probe uses one in-flight dependent load to isolate the fabric path; real kernels issue independent and coalesced requests, use L1 and shared memory, and overlap memory with computation, so the absolute cycle counts are not application throughput. The bandwidth uniformity reported for other GPUs [4] is not measured here; our throughput data are aggregate controls, not per-SM.
13
Reproducibility and Artifact
The artifact is organized for ancillary-file submission. The TeX source, BibTeX file, and figures remain in the submission root. Code, configurations, raw CSV results, manifests, and regeneration scripts are placed under anc/. Raw observations are CSV or JSON and compress well below the 50 MB ancillary budget. Every randomized chain has an explicit 64-bit seed; each run records command lines, source SHA-256 hashes, host information, CUDA/NVIDIA versions, and before/after telemetry. The per-SM map regenerates from topology_slice_sweep.csv and topology_chain_uniform.csv via analyze_topology.py. The placement oracle is built from the labeled fingerprint shots (topology_fp_a*.csv.zst) by train_oracle.py; the trained model, its probe-bank metadata, and a held-out fingerprint sample are included so predict_sm.py runs the oracle offline without a GPU. The cross-architecture comparison (analyze_cross_device.py), the same-model fingerprinting (analyze_samemodel.py, from the second L40’s runs), the sustained-load stability runs (run_stability.py, analyze_stability.py, with the topology_stability kernel), and the placement campaign (run_placement.py, analyze_placement.py, with the placement_sched kernel) all regenerate from their archived runs. Reproducing the GPU measurements requires an NVIDIA L40 or a compatible device (the cross-device run additionally uses a Blackwell RTX 5090), and the scripts record changed device properties in a new manifest rather than overwriting archived results.
14
Conclusion
On the tested NVIDIA L40 the L2-hit latency is not a single number but a function of physical SM placement, spanning 222.5–339.2 cycles (52%) across the 142 SMs. An additive model with a 57.2-cycle SM term and a 39.5-cycle slice term explains 87% of the variation, the SM term is two-fold symmetric at period 72 in line with the AD102 GPC layout, and the effect is confirmed physical by independent access patterns (r = 1.000) and a sub-0.01-cycle noise floor. The L40 lacks the discrete inter-partition step of A100/H100 but, unlike them, has a non-uniform per-SM average. The map is a stable per-die signature: it is unchanged after an hour at full load, two identical-model L40s are separable at 100% from it despite near-identical mean latency, and it lets a kernel self-localize to its own SM at 99.2%—a fingerprinting primitive, not a victim attack. The same structure appears on a Blackwell RTX 5090, so the finding is not L40-specific, and distributing latency-bound work by the per-SM map cuts makespan by up to 11%—a concrete use of the topology, bounded to the latency-bound regime. The single-thread residency results—an effective-capacity boundary near 96 MiB, line-tag scaling, a 15
prefetch-modifier null result, and a 66 MiB persisting boundary—hold as controls, but each is an average over the non-uniform latency this paper maps.
References [1] Sankha Baran Dutta, Hoda Naghibijouybari, Arjun Gupta, Nael Abu-Ghazaleh, Andres Marquez, and Kevin Barker. Spy in the GPU-box: Covert and side channel attacks on multi-GPU systems. arXiv preprint arXiv:2203.15981, 2022. [2] Nikos Hardavellas, Michael Ferdman, Babak Falsafi, and Anastasia Ailamaki. Reactive NUCA: Near-optimal block placement and replication in distributed caches. In Proceedings of the 36th Annual International Symposium on Computer Architecture (ISCA), pages 184–195, 2009. [3] Zhe Jia, Marco Maggioni, Benjamin Staiger, and Daniele P. Scarpazza. Dissecting the nvidia volta gpu architecture via microbenchmarking. arXiv preprint arXiv:1804.06826, 2018. [4] Zhixian Jin, Christopher Rocca, Jiho Kim, Hans Kasan, Minsoo Rhu, Ali Bakhoda, Tor M. Aamodt, and John Kim. Uncovering real gpu noc characteristics: Implications on interconnect architecture. In Proceedings of the 57th Annual IEEE/ACM International Symposium on Microarchitecture, pages 885–898, 2024. [5] Changkyu Kim, Doug Burger, and Stephen W. Keckler. An adaptive, non-uniform cache structure for wire-delay dominated on-chip caches. In Proceedings of the 10th International Conference on Architectural Support for Programming Languages and Operating Systems (ASPLOS X), pages 211–222, 2002. [6] Chao Luo, Rengan Fan, Zeyu Li, Dayou Du, and Qiang Chen. Dissecting the nvidia hopper architecture through microbenchmarking and multiple level analysis. arXiv preprint arXiv:2501.12084, 2025. [7] Xinxin Mei and Xiaowen Chu. Dissecting gpu memory hierarchy through microbenchmarking. IEEE Transactions on Parallel and Distributed Systems, 28(1):72–86, 2017. [8] Hoda Naghibijouybari, Ajaya Neupane, Zhiyun Qian, and Nael Abu-Ghazaleh. Rendered insecure: GPU side channel attacks are practical. In Proceedings of the 2018 ACM SIGSAC Conference on Computer and Communications Security, pages 2139–2153, 2018. [9] NVIDIA Corporation. CUDA Programming Guide: L2 Cache Control. https://docs.n vidia.com/cuda/cuda-programming-guide/04-special-topics/l2-cache-control.h tml, 2026. Accessed 2026-06-21. [10] NVIDIA Corporation. Nsight Compute Profiling Guide. https://docs.nvidia.com/ns ight-compute/ProfilingGuide/, 2026. Accessed 2026-06-21. [11] NVIDIA Corporation. Parallel Thread Execution ISA, Version 9.3. https://docs.nvidi a.com/cuda/parallel-thread-execution/, 2026. Accessed 2026-06-21. [12] Rogier van Stigt, Simon N. Swatman, and Ana Lucia Varbanescu. Isolating gpu architectural features using parallelism-aware microbenchmarks. In Proceedings of the 2022 ACM/SPEC International Conference on Performance Engineering, pages 77–88, 2022.
16
[13] Xianwei Zhang and Evgeny Shcherbakov. DELTA: Validate gpu memory profiling with microbenchmarks. In Proceedings of the International Symposium on Memory Systems, pages 97–104, 2020.
17