💻 computer science

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

Dit artikel onthult dat de L2-cache latentie op NVIDIA L40 en Blackwell GPU's niet-uniform is en sterk afhankelijk is van de specifieke fysieke Streaming Multiprocessor (SM) die de load initieert, wat een stabiel, gebruikersniveau primitief mogelijk maakt voor kernel zelf-lokalisatie, device fingerprinting en geoptimaliseerde werkverdeling die de maaktijd met tot wel 11% vermindert.

Oorspronkelijke auteurs: Faruk Alpay, Baris Basaran

Gepubliceerd 2026-06-23

📖 5 min leestijd🧠 Diepgaand

CC BY 4.0

Oorspronkelijke auteurs: Faruk Alpay, Baris Basaran

Oorspronkelijk artikel gelicentieerd onder CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ Dit is een AI-gegenereerde uitleg van het onderstaande artikel. Het is niet geschreven of goedgekeurd door de auteurs. Raadpleeg het oorspronkelijke artikel voor technische nauwkeurigheid. Lees de volledige disclaimer

Stel je een enorme, razendsnelle bibliotheek voor (de GPU) waar duizenden bibliothecarissen (genaamd SM's, of Streaming Multiprocessors) samenwerken om boeken (data) te halen uit een gigantische, gedeelde opslagruimte (de L2-cache).

Jarenlang geloofden computerwetenschappers dat deze bibliotheek werkte als een perfect uniform systeem: ongeacht welke bibliothecaris je ook vroeg, de tijd die het kostte om een boek uit de opslagruimte te pakken was exact hetzelfde. Ze dachten dat de opslagruimte een enkele, vlakke poel van informatie was waar afstand niet toe deed.

Dit artikel onthult echter dat dit geloof onjuist is. De auteurs ontdekten dat op de NVIDIA L40 GPU waar een bibliothecaris zich fysiek bevindt, bepaalt hoe snel hij een boek kan pakken.

Hier is de uitsplitsing van hun bevindingen met behulp van eenvoudige analogieën:

1. De ontdekking: "Afstand doet ertoe"

Denk aan de opslagruimte niet als een vlakke vloer, maar als een groot magazijn met gangpaden.

Het oude beeld: Iedereen dacht dat elke bibliothecaris op precies dezelfde afstand van de planken stond. Als je Bibliothecaris A of Z vroeg, was de wacjesduur identiek (ongeveer 279 "ticks" van de klok).
De nieuwe realiteit: De auteurs maten de wachttijd voor elke van de 142 bibliothecarissen. Ze ontdekten dat sommige bibliothecarissen direct naast de planken staan en boeken krijgen in 222 ticks. Anderen staan aan het verre uiteinde van het magazijn en doen er 339 ticks over.
Het resultaat: Dat is een verschil van 52% in snelheid. Het is alsof één persoon naar de voordeur rent terwijl een ander naar de achteruitgang moet rennen om hetzelfde pakketje te krijgen.

2. Het "Spiegelbeeld"-patroon

Toen de auteurs in kaart brachten wie snel en wie traag was, zagen ze geen willekeurige ruis. Ze zagen een perfect patroon.

De 142 bibliothecarissen zijn verdeeld in twee identieke helften (zoals twee vleugels van een gebouw).
Bibliothecaris #1 in de eerste vleugel heeft exact hetzelfde snelheidsprofiel als Bibliothecaris #1 in de tweede vleugel.
Deze "spiegelymmetrie" komt overeen met de werkelijke fysieke blauwdruk van de computerchip, wat bewijst dat dit geen fout in de software is — het is een feit van de fysieke constructie van de hardware.

3. Het "Vingerafdruk"-effect

Omdat elke bibliothecaris een unieke snelheid heeft op basis van zijn fysieke plek, heeft de bibliotheek een geheime identiteit.

Zelflokalisatie: Als jij een bibliothecaris (een computerprogramma) bent en je vraagt: "Hoe lang duurt het mij om een boek te pakken?", dan kun je direct uitvinden op welke specifieke plek in het magazijn je staat. De auteurs bouwden een hulpmiddel dat jouw specifieke plek met 99% nauwkeurigheid kan identificeren.
Device Fingerprinting: Zelfs als je twee gloednieuwe, identieke L40 GPU's hebt (twee identieke bibliotheken), zijn ze licht van elkaar verschillend. De ene bibliotheek kan bijvoorbeeld hebben dat Bibliothecaris #5 iets dichter bij de planken staat dan Bibliothecaris #5 van de andere bibliotheek, door kleine productieverschillen. De auteurs konden de twee identieke machines 100% van de tijd van elkaar onderscheiden door alleen deze minuscule snelheidverschillen te meten. Het is alsof je twee identieke tweelingen van elkaar kunt onderscheiden door de manier waarop ze lopen.

4. Is dit een beveiligingsrisico?

De auteurs verduidelijken zorgvuldig wat dit betekent voor de beveiliging.

Wat het WEL is: Een manier voor een programma om te weten waar het zich bevindt binnen de computer. Het is als een persoon die een kamer binnenloopt en beseft: "Oh, ik sta in de hoek bij het raam."
Wat het NIET is: Het is geen manier om geheimen van andere programma's te stelen. De auteurs benadruken dat dit alleen de eigen snelheid van het programma meet. Het kan niet meekijken met wat andere programma's doen of hun data stelen. Het is een "zelflokalisatie"-hulpmiddel, geen "spion"-hulpmiddel.

5. Het praktische voordeel: Slimmere planning

Waarom is dit belangrijk voor de prestaties?

Stel je voor dat je een lijst hebt met 100 taken die uitgevoerd moeten worden.
De oude manier: Je wijst taken willekeurig toe. Sommige gaan naar de bibliothecarissen die ver weg staan (traag), en andere naar degenen die dichtbij staan (snel). De hele klus wacht tot de langzaamste mensen klaar zijn.
De nieuwe manier: Nu we de kaart hebben, kunnen we het zware werk toewijzen aan de bibliothecarissen die het dichtst bij de planken staan.
Het resultaat: Door dit te doen, lieten de auteurs zien dat ze het werk 11% sneller konden voltooien voor taken die sterk afhankelijk zijn van de cache. Echter, als de taak veel data moet ophalen uit het hoofdgeheugen (een andere, tragere opslag), helpt deze truc niet.

6. Het is niet beperkt tot één chip

De auteurs testten dit ook op een nieuwere, andere chip (de RTX 5090). Ze vonden dat dezelfde "afstand doet ertoe"-regel daar ook geldt, hoewel het patroon iets anders is. Dit bewijst dat het oude idee van een "uniforme" cache waarschijnlijk onjuist is voor veel moderne high-end computers.

Samenvatting

Het artikel verbrijzelt de illusie dat alle delen van een GPU-cache gelijk zijn. Het onthult dat fysieke locatie de snelheid bepaalt. Door deze verborgen snelheidverschillen in kaart te brengen, kunnen we:

Identificeren waar een programma exact wordt uitgevoerd.
Twee identieke machines van elkaar onderscheiden.
Specifieke taken versnellen door ze toe te wijzen aan de snelste fysieke plekken.

Het blijkt dat de computerchip geen plat, uniform veld is; het is een landschap met heuvels en dalen, en het bezitten van de kaart maakt je sneller.

Technische Samenvatting: Niet-uniforme L2-cachelatentie op NVIDIA L40

Probleemstelling
De NVIDIA L40 GPU, gebaseerd op de volledige AD102-implementatie, exposeert een 96 MiB L2-cache die conventioneel in prestatie-instrumenten en academische literatuur wordt gemodelleerd als een uniforme pool met een enkele, apparaatbrede hit-latentie. Deze abstractie gaat ervan uit dat elke Streaming Multiprocessor (SM) die de L2-cache aanroept, dezelfde kosten met zich meebrengt. De fysieke realiteit van de L2-cache omvat echter meerdere slices die verbonden zijn met geheugencontrollers, die worden aangeroepen via een on-chip netwerk (NoC). Hoewel eerder werk op V100-, A100- en H100-architecturen vaststelde dat NoC-latentie niet-uniform is, benadrukten die studies dat de gemiddelde latentie per kern ongeveer uniform was, waarbij niet-uniformiteit voornamelijk verscheen als discrete stappen over partitiegrenzen heen. Het bleef een openstaande vraag of een single-die, single-partition Ada Lovelace datacenter-GPU (zoals de L40) een vergelijkbare structuur vertoont, en zo ja, of de gemiddelde latentie per SM uniform of plaatsingsafhankelijk is.

Methodologie
De auteurs ontwikkelden een microbenchmarking-methodologie op gebruikersniveau om de L2-hit-latentie vanaf elke fysieke SM te meten zonder verhoogde privileges of hardwareprestatie-tellers te vereisen.

Probe-ontwerp: De meting maakt gebruik van een single-thread, single-warp afhankelijke load-keten (ld.global.cg) om een enkele round-trip door de cache-fabric te garanderen. Twee bufferformaten werden toegepast: een "single-line" vorm om specifieke 128 B slices te beproeven en een "chain" vorm (2 MiB willekeurige permutatie) om de gemiddelde afstand tot alle slices te meten.
SM-isolatie: Om on-chip congestie te elimineren, worden blocks geserialiseerd met behulp van een globale turn-counter. Elk block claimt een beurt via atomicAdd, wacht, voert een warm-up uit en meet zijn afhankelijke loads in isolatie.
Dekking: De probe dekt alle 142 fysieke SM's op de L40 en 256 slice-probes. Een vergelijkbare probe werd uitgevoerd op een NVIDIA RTX 5090 (Blackwell GB202) om generalisatie te testen.
Modellering: De latentiegegevens worden geanalyseerd met een additief Non-Uniform Cache Access (NUCA) model: $L = \mu + a(sm) + b(slice)$ , waarbij $a(sm)$ de SM-plaatsingsterm vertegenwoordigt en $b(slice)$ de slice-plaatsingsterm vertegenwoordigt.

Belangrijkste Resultaten

Significante Latentievariatie: In tegenstelling tot het uniforme model is de L2-hit-latentie op de L40 niet constant. Deze varieert van 222,5 tot 339,2 cycli (een bereik van 52%), met een gemiddelde variatie van 57,2 cycli over de SM's. De ruisvloer is verwaarloosbaar (<0,01 cycli), wat bevestigt dat de variatie fysiek is.
Tweevoudige Symmetrie: De SM-plaatsingsterm vertoont een sterke tweevoudige symmetrie. Het splitsen van de 142 SM's bij identifier 72 onthult twee helften (72 en 70 SM's) met bijna identieke latentieprofielen (correlatie $r=0,999$ ). Dit komt overeen met de AD102-vloerplan van twee symmetrische zes-GPC clusters.
Continue Gradiënt versus Discrete Stappen: In tegen tegenstelling tot de A100/H100, die een discrete latentiestap vertonen bij de partitiegrens, vertoont de L40 een continue gradiënt van latentie over de SM's zonder bimodale kloof.
Passendheid van het Additieve Model: Het additieve model verklaart 87% van de latentievariatie ( $R^2=0,87$ ). Het toevoegen van een enkele rank-1 interactieterm verhoogt dit naar 0,98, wat aangeeft dat de latentiekaart een laagdimensionale geometrie volgt die de GPC-lay-out volgt.
Generalisatie en Specificiteit:
- Cross-Architectuur: De RTX 5090 vertoont ook niet-uniformiteit (46% spreiding), maar met een zwakkere symmetrie ( $r=0,80$ ) en een tragere absolute latentie.
- Device Fingerprinting: Twee fysiek verschillende, identieke L40-modellen zijn met 100% nauwkeurigheid te onderscheiden op basis van hun per-SM kaarten, ondanks dat ze bijna identieke gemiddelde latenties hebben (verschil van 0,28 cyclus). De per-SM kaart fungeert als een stabiele hardware-identiteit.
Stabiliteit: De per-SM latentiekaart blijft ongewijzigd na één uur van 100% benutting op beide apparaten, wat bevestigt dat de topologie een eigenschap van de silicium is en geen transiënte staat.

Bijdragen en Betekenis
Het artikel claimt de volgende bijdragen:

Topologie Mapping: Het biedt de eerste per-SM L2-latentiekaart voor de L40, die een 52% spreiding onthult die de uniforme cache-aanname voor latentie-gebonden workloads invalideert.
Self-Localization Primitive: De auteurs demonstreren dat een enkele L2-latentie probe een kernel in staat stelt zijn eigen fysieke SM-plaatsing te identificeren met 99,2% nauwkeurigheid (of 92% over een gemengde device-set) en de specifieke device-die te identificeren met 100% nauwkeurigheid. Dit wordt geframed als een "self-localization" primitive, en niet als een aanval op een victim-kernel, aangezien het geen geheime data extraheert en vertrouwt op de eigen plaatsing van de kernel.
Prestatie-implicaties: Door latentie-gebonden werk te distribueren volgens de per-SM kaart (werk te plaatsen op snellere SM's), bereikten de auteurs een reductie van 10,9% in de maaktijd voor L2-resident workloads. Dit voordeel verdwijnt in DRAM-gebonden regimes, wat de regime-specifieke aard van de impact van de topologie benadrukt.
Herinterpretatie van Residency: Het artikel herinterpreteert bestaande single-thread residency controles (capaciteitsgrenzen, line-tag scaling en persisting-L2 resultaten) als gemiddelden over deze niet-uniforme distributie, in plaats van constanten.

Omvang en Beperkingen
De auteurs beperken expliciet de omvang van hun beveiligingsclaims. De techniek maakt self-localization mogelijk (een kernel die zijn eigen plaatsing leest) en device fingerprinting. Het kan niet de plaatsing van een andere, co-resident "victim" kernel afleiden, noch extraheert het geheime data; dergelijke capaciteiten zouden een congestie- of interferentiekanaal vereisen dat niet in deze studie is gemeten. De studie is beperkt tot twee L40's en één RTX 5090, en de mapping van %smid naar de fysieke vloerplan wordt afgeleid van symmetrie en GPC-aantallen in plaats van onafhankelijk bevestigd door die-fotografie. De absolute cyclusaantallen zijn specifief voor het microbenchmark-ontwerp en vertalen zich niet direct naar applicatie-doorvoer waar onafhankelijke verzoeken en L1-gebruik voorkomen.