💻 computer science

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

Questo articolo rivela che la latenza di hit della cache L2 sulle GPU NVIDIA L40 e Blackwell è non uniforme e altamente dipendente dal particolare Streaming Multiprocessor (SM) fisico che emette il carico, consentendo un primitivo stabile a livello utente per l'auto-localizzazione del kernel, il fingerprinting del dispositivo e la distribuzione ottimizzata del lavoro che riduce il makespan fino all'11%.

Autori originali: Faruk Alpay, Baris Basaran

Pubblicato 2026-06-23

📖 5 min di lettura🧠 Approfondimento

CC BY 4.0

Autori originali: Faruk Alpay, Baris Basaran

Articolo originale sotto licenza CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ Questa è una spiegazione generata dall'IA dell'articolo qui sotto. Non è stata scritta né approvata dagli autori. Per precisione tecnica, consulta l'articolo originale. Leggi il disclaimer completo

Immagina una biblioteca enorme e ad alta velocità (la GPU) dove migliaia di bibliotecari (chiamati SM, o Streaming Multiprocessors) lavorano insieme per recuperare libri (dati) da un enorme archivio condiviso (la cache L2).

Per anni, gli scienziati informatici hanno creduto che questa biblioteca funzionasse come un sistema perfettamente uniforme: non importava quale bibliotecario interpellassi, il tempo necessario per prendere un libro dall'archivio era esattamente lo stesso. Pensavano che l'archivio fosse un unico pool piatto di informazioni dove la distanza non contava nulla.

Questo articolo, tuttavia, rivela che questa convinzione è errata. Gli autori hanno scoperto che su NVIDIA L40 GPU, dove si trova fisicamente un bibliotecario determina la velocità con cui può recuperare un libro.

Ecco la suddivisione delle loro scoperte utilizzando analogie semplici:

1. La scoperta del "La distanza conta"

Pensa all'archivio non come a un pavimento piatto, ma come a un grande magazzino con corridoi.

La vecchia visione: Tutti pensavano che ogni bibliotecario si trovasse alla stessa identica distanza dagli scaffali. Se chiedevi al Bibliotecario A o al Bibliotecario Z, il tempo di attesa era identico (circa 279 "tick" dell'orologio).
La nuova realtà: Gli autori hanno misurato il tempo di attesa per ogni singolo uno dei 142 bibliotecari. Hanno scoperto che alcuni bibliotecari si trovano proprio accanto agli scaffali e recuperano i libri in 222 tick. Altri si trovano all'estremità opposta del magazzino e impiegano 339 tick.
Il risultato: C'è una differenza del 52% nella velocità. È come se una persona corresse verso la porta anteriore mentre un'altra deve correre verso l'uscita posteriore solo per ritirare lo stesso pacco.

2. Il modello a "Immagine Speculare"

Quando gli autori hanno mappato chi era veloce e chi era lento, non hanno visto rumore casuale. Hanno visto un modello perfetto.

I 142 bibliotecari sono divisi in due metà identiche (come due ali di un edificio).
Il Bibliotecario n. 1 nella prima ala ha esattamente lo stesso profilo di velocità del Bibliotecario n. 1 nella seconda ala.
Questa "simmetria speculare" corrisponde al vero progetto fisico del chip del computer, dimostrando che questo non è un glitch del software — è un fatto della costruzione fisica dell'hardware.

3. L'effetto "Impronta Digitale"

Poiché ogni bibliotecario ha una velocità unica basata sulla sua posizione fisica, la biblioteca ha un'identità segreta.

Auto-localizzazione: Se sei un bibliotecario (un programma per computer) e chiedi: "Quanto tempo mi serve per prendere un libro?", puoi capire istantaneamente in quale punto esatto del magazzino ti trovi. Gli autori hanno costruito uno strumento in grado di identificare il tuo punto specifico con un'accuratezza del 99%.
Fingerprinting del dispositivo: Anche se hai due GPU L40 nuove e identiche (due biblioteche identiche), sono leggermente diverse. In una biblioteca il Bibliotecario n. 5 potrebbe trovarsi leggermente più vicino agli scaffali rispetto al Bibliotecario n. 5 dell'altra biblioteca a causa di minuscole differenze di produzione. Gli autori sono riusciti a distinguere le due macchine identiche il 100% delle volte, misurando queste piccole differenze di velocità. È come essere in grado di distinguere due gemelli identici dal modo in cui camminano.

4. È un rischio per la sicurezza?

Gli autori chiariscono con cura cosa significhi questo per la sicurezza.

Ciò che È: Un modo per un programma di sapere dove si trova all'interno del computer. È come una persona che entra in una stanza e si rende conto: "Oh, sono in un angolo vicino alla finestra".
Ciò che NON È: Non è un modo per rubare segreti ad altri programmi. Gli autori sottolineano che questo misura solo la propria velocità del programma. Non può sbirciare ciò che fanno altri programmi o rubare i loro dati. È uno strumento di "auto-localizzazione", non uno strumento di "spionaggio".

5. Il beneficio pratico: Scheduling più intelligente

Perché questo è importante per le prestazioni?

Immagina di avere una lista di 100 compiti da svolgere.
Il vecchio modo: Assegni i compiti casualmente. Alcuni vanno ai bibliotecari che si trovano lontani (lenti), altri a quelli che si trovano vicini (veloci). L'intero lavoro aspetta che le persone più lente finiscano.
Il nuovo modo: Ora che abbiamo la mappa, possiamo assegnare i lavori pesanti ai bibliotecari che si trovano più vicini agli scaffali.
Il risultato: Facendo questo, gli autori hanno dimostrato di poter completare il lavoro il 11% più velocemente per i compiti che dipendono fortemente dalla cache. Tuttavia, se il compito richiede il recupero di dati dalla memoria principale (un archivio diverso e più lento), questo trucco non aiuta.

6. Non riguarda solo un chip

Gli autori hanno testato questo anche su un chip più recente e diverso (l'RTX 5090). Hanno scoperto che la stessa regola del "la distanza conta" si applica anche lì, anche se il modello è leggermente diverso. Questo dimostra che la vecchia idea di una cache "uniforme" è probabilmente errata per molti moderni computer ad alte prestazioni.

Riassunto

L'articolo distrugge l'illusione che tutte le parti della cache di una GPU siano uguali. Rivela che la posizione fisica determina la velocità. Mappando queste velocità nascoste, possiamo:

Identificare esattamente dove sta girando un programma.
Distinguere tra due macchine identiche.
Accelerare compiti specifici assegnandoli ai punti fisici più veloci.

Si scopre che il chip del computer non è un campo piatto e uniforme; è un paesaggio con colline e valli, e conoscere la mappa ti rende più veloce.

Sintesi Tecnica: Latenza della Cache L2 Non Uniforme su NVIDIA L40

Definizione del Problema
La GPU NVIDIA L40, basata sull'implementazione completa di AD102, espone una cache L2 da 96 MiB che viene convenzionalmente modellata negli strumenti di analisi delle prestazioni e nella letteratura accademica come un pool uniforme con una singola latenza di hit a livello di intero dispositivo. Questa astrazione assume che qualsiasi Streaming Multiprocessor (SM) che accede alla cache L2 incorra nello stesso costo. Tuttavia, la realtà fisica della cache L2 prevede molteplici slice collegate ai controller di memoria, accessibili tramite una rete on-chip (NoC). Sebbene studi precedenti sulle architetture V100, A100 e H100 abbiano stabilito che la latenza NoC è non uniforme, tali studi sottolineavano che la latenza media per core era approssimativamente uniforme, con la non-uniformità che appariva principalmente come scatti discreti attraverso i confini delle partizioni. Restava aperto il quesito se una GPU single-die, single-partition della serie Ada Lovelace per data center (come la L40) esibisse una struttura comparabile e, in tal caso, se la latenza media per SM fosse uniforme o dipendente dal posizionamento.

Metodologia
Gli autori hanno sviluppato una metodologia di micro-benchmarking a livello utente per misurare la latenza di hit L2 da ogni SM fisico senza richiedere privilegi elevati o contatori di prestazioni hardware.

Design della Sonda: La misurazione utilizza una catena di carichi dipendenti a singolo thread e singolo warp (ld.global.cg) per garantire un singolo round-trip attraverso il fabric della cache. Sono stati impiegati due shape di buffer: uno a "singola riga" (single-line) per sondare specifiche slice da 128 B e uno a "catena" (chain) (permutazione casuale di 2 MiB) per misurare la distanza media da tutte le slice.
Isolamento degli SM: Per eliminare la contesa on-chip, i blocchi sono serializzati tramite un contatore di turno globale. Ogni blocco reclama un turno tramite atomicAdd, attende, esegue il warm-up e temporizza i propri carichi dipendenti in isolamento.
Copertura: La sonda copre tutti i 142 SM fisici sulla L40 e 256 sonde di slice. Una sonda simile è stata eseguita su una NVIDIA RTX 5090 (Blackwell GB202) per testarne la generalizzazione.
Modellazione: I dati di latenza sono analizzati utilizzando un modello di accesso alla cache non uniforme (NUCA) additivo: $L = \mu + a(sm) + b(slice)$ , dove $a(sm)$ rappresenta il termine di posizionamento dell'SM e $b(slice)$ rappresenta il termine di posizionamento della slice.

Risultati Chiave

Variazione Significativa della Latenza: Contrariamente al modello uniforme, la latenza di hit L2 sulla L40 non è costante. Spazia tra 222,5 e 339,2 cicli (un intervallo del 52%), con una variazione media di 57,2 cicli tra gli SM. Il rumore di fondo è trascurabile (<0,01 cicli), confermando che la variazione è di natura fisica.
Simmetria a Due Facce: Il termine di posizionamento dell'SM esibisce una forte simmetria a due facce. Dividendo i 142 SM all'identificatore 72 si rivelano due metà (72 e 70 SM) con profili di latenza quasi identici (correlazione $r=0,999$ ). Ciò si allinea con il floorplan di AD102 composto da due cluster GPC simmetrici.
Gradiente Continuo vs Scatti Discreti: A differenza di A100/H100, che mostrano uno scatto di latenza discreto al confine della partizione, la L40 esibisce un gradiente continuo di latenza attraverso gli SM senza un gap bimodale.
Fit del Modello Additivo: Il modello additivo spiega l'87% della variazione di latenza ( $R^2=0,87$ ). L'aggiunta di un singolo termine di interazione rank-1 aumenta questo valore a 0,98, indicando che la mappa di latenza è una geometria a bassa dimensionalità che segue il layout GPC.
Generalizzazione e Specificità:
- Cross-Architettura: Anche la RTX 5090 mostra non-uniformità (spread del 46%), ma con una simmetria più debole ( $r=0,80$ ) e una latenza assoluta più lenta.
- Fingerprinting del Dispositivo: Due L40 fisicamente distinti e di modello identico sono distinguibili con un'accuratezza del 100% basandosi sulle loro mappe per-SM, nonostante abbiano latenze medie quasi identiche (differenza di 0,28 cicli). La mappa per-SM agisce come una stabile identità hardware.
Stabilità: La mappa di latenza per-SM rimane invariata dopo un'ora di utilizzo al 100% su entrambi i dispositivi, confermando che la topologia è una proprietà del silicio piuttosto che uno stato transitorio.

Contributi e Significato
Il documento rivendica i seguenti contributi:

Mappatura della Topologia: Fornisce la prima mappa di latenza L2 per-SM per la L40, rivelando uno spread del 52% che invalida l'assunto di cache uniforme per i carichi di lavoro limitati dalla latenza.
Primitiva di Auto-Localizzazione: Gli autori dimostrano che una singola sonda di latenza L2 consente a un kernel di identificare il proprio posizionamento fisico con un'accuratezza del 99,2% (o del 92% su un set di dispositivi misti) e di identificare il die specifico del dispositivo con il 100% di accuratezza. Questo è presentato come una primitiva di "auto-localizzazione", non come un attacco a un kernel vittima, poiché non estrae dati segreti e si basa sul proprio posizionamento del kernel.
Implicazioni sulle Prestazioni: Distribuendo il lavoro limitato dalla latenza secondo la mappa per-SM (posizionando il lavoro sugli SM più veloci), gli autori hanno ottenuto una riduzione del 10,9% del makespan per i carichi di lavoro residenti in L2. Questo beneficio svanisce nei regimi limitati dalla DRAM, evidenziando la natura specifica del regime dell'impatto della topologia.
Reinterpretazione della Residenza: Il documento reinterpreta i controlli esistenti sulla residenza a singolo thread (limiti di capacità, scaling del tag di linea e risultati di persistenza L2) come medie su questa distribuzione non uniforme, piuttosto che come costanti.

Ambito e Limitazioni
Gli autori limitano esplicitamente l'ambito delle loro rivendicazioni di sicurezza. La tecnica consente l'auto-localizzazione (un kernel che legge il proprio posizionamento) e il fingerprinting del dispositivo. Non permette di inferire il posizionamento di un separato kernel "vittima" co-residente, né estrae dati segreti; tali capacità richiederebbero un canale di contesa o interferenza non misurato in questo studio. Lo studio è limitato a due L40 e una RTX 5090, e la mappatura di %smid sul floorplan fisico è inferita dalla simmetria e dai conteggi GPC piuttosto che confermata indipendentamente tramite fotografia del die. Il conteggio dei cicli assoluti è specifico del design del micro-benchmark e non si traduce direttamente nel throughput dell'applicazione dove avvengono richieste indipendenti e uso della L1.