Accurate Models of NVIDIA Tensor Cores

Titolo: Modelli Accurati dei Tensor Core NVIDIA

Autori: Faizan A. Khattak e Mantas Mikaitis (Università di Leeds, Regno Unito)

1. Il Problema

La moltiplicazione di matrici è un'operazione fondamentale per l'addestramento delle reti neurali e l'inferenza. Per accelerare questo processo, le GPU moderne (in particolare NVIDIA) utilizzano unità hardware specializzate chiamate Tensor Core. Tuttavia, questi acceleratori presentano diverse criticità per il calcolo scientifico e la riproducibilità:

• Non conformità agli standard: I Tensor Core non aderiscono al comportamento dell'aritmetica in virgola mobile IEEE 754. Implementano logiche di arrotondamento, normalizzazione e allineamento specifiche per massimizzare il throughput, spesso a scapito della precisione standard.
• Mancanza di standardizzazione: Diversi vendor e diverse generazioni di architetture (Volta, Ampere, Hopper, Blackwell) implementano le moltiplicazioni miste (mixed-precision) in modi diversi. Questo porta a risultati non riproducibili tra diverse GPU.
• Difficoltà di analisi: L'assenza di documentazione pubblica dettagliata sulle caratteristiche numeriche interne (come la larghezza dell'accumulatore, i bit di guardia, i punti di normalizzazione) rende difficile per gli scienziati computazionali eseguire analisi di errore di arrotondamento accurate o sviluppare algoritmi robusti.
• Limiti dei test esistenti: I vettori di test precedenti spesso non riescono a distinguere tra modelli hardware diversi o non sono stati validati su un'ampia gamma di piattaforme a causa della scarsa disponibilità di hardware.

2. Metodologia

Gli autori hanno sviluppato un approccio iterativo per creare modelli software in MATLAB che emulano il comportamento bit-per-bit dei Tensor Core. Il processo si articola in due fasi principali:

1. Generalised Numerical Feature Testing (GNFT):

   • Vengono utilizzati vettori di test progettati per attivare casi specifici (es. numeri subnormali, overflow, allineamento di mantissa) per dedurre le caratteristiche numeriche di base (es. presenza di bit di allineamento extra, modalità di arrotondamento).
   • Questo metodo fornisce una prima approssimazione del modello.

2. Input Space Search Method (ISSM):

   • Per affinare il modello, viene eseguita una ricerca semi-esauriva nello spazio degli input. Vengono generati milioni di vettori di input casuali (distribuzioni normali e uniformi, inclusi casi limite come subnormali, Inf e NaN).
   • I risultati del modello software vengono confrontati con quelli ottenuti eseguendo lo stesso codice su hardware reale (GPU).
   • In caso di discrepanze (mismatch), gli autori analizzano i casi di fallimento per identificare caratteristiche numeriche mancanti o errate nel modello (es. logica di normalizzazione dei prodotti parziali) e aggiornano il modello.
   • Il ciclo si ripete fino a quando il modello produce risultati bit-accurati rispetto all'hardware.

3. Contributi Chiave

• Modellazione di 9 GPU e 31 configurazioni: Sono stati sviluppati modelli accurati per i Tensor Core di nove diverse GPU NVIDIA (V100, A2, A30, A100, H100, H200, L40S, RTX 1000 Ada, B200) coprendo 31 combinazioni di formati di input/output (fp8, fp16, bf16, tf19, fp32).
• Scoperta di nuove caratteristiche numeriche:
  • Bit di allineamento extra ($n_{eab}$): Determinazione precisa del numero di bit extra utilizzati durante l'allineamento delle mantisse (es. $n_{eab}=0$ su V100, $n_{eab}=1$ su A100, $n_{eab}=2$ su H100/H200/B200).
  • Limitazione degli esponenti: Rilevamento di un limite massimo sull'esponente utilizzato per l'allineamento quando si sommano prodotti subnormali (es. limite a -133 su A100/H100).
  • Pattern di intercalamento (Interleaving): Scoperta che su H100/H200/B200, per gli input fp8 tramite istruzioni mma.sync, i vettori di input vengono elaborati in modo intercalato su due invochi del core, un comportamento non documentato.
  • Gestione dei prodotti non normalizzati: Conferma che i prodotti parziali rimangono in forma non normalizzata durante l'accumulo in molte architetture, influenzando la precisione.
• Toolbox MATLAB Open Source: Pubblicazione della libreria MATLAB Tensor Core v0.4.1 su GitHub. Include modelli predefiniti per le GPU sopra citate e un modello generalizzato (Generic_BFMA_TC) che permette agli utenti di definire parametri arbitrari (dimensione FMA, bit di allineamento, modalità di arrotondamento) per simulare varianti custom.

4. Risultati

• Accuratezza Bit-per-Bit: I modelli software sono stati validati contro l'hardware reale utilizzando $10^7$ (e in alcuni casi $10^8$) vettori di input casuali. In tutte le configurazioni testate, i modelli hanno prodotto risultati identici a livello di bit rispetto alle GPU fisiche.
• Dettagli Architetturali:
  • V100: $N_{FMA}=4$, nessun bit di allineamento extra, allineamento a 25 bit.
  • A100/A2/A30: $N_{FMA}=8$ (fp16/bf16), 1 bit di allineamento extra, gestione specifica degli esponenti subnormali.
  • H100/H200: $N_{FMA}=16$ (fp16/bf16), 2 bit di allineamento extra, comportamento diverso per fp8 (conversione interna a fp16 con pattern intercalato).
  • B200: Comportamento numerico identico a H100/H200 per la maggior parte dei formati, ma con supporto nativo per istruzioni di nuova generazione (tcgen05) che offrono $N_{FMA}=32$ per fp8 con 23 bit frazionari.
• Validazione degli Algoritmi: I modelli sono stati utilizzati per replicare esperimenti di aritmetica multi-parola (multi-word arithmetic) per l'emulazione di alta precisione. I risultati hanno mostrato differenze significative negli errori di norma tra le diverse generazioni di GPU (es. V100 vs B200), dimostrando l'utilità del toolbox per prevedere il comportamento degli algoritmi prima dell'implementazione su hardware reale.

5. Significato e Impatto

• Riproducibilità Scientifica: Fornisce agli scienziati computazionali uno strumento affidabile per ottenere risultati riproducibili su diverse piattaforme, superando l'opacità dell'hardware proprietario.
• Sviluppo di Algoritmi: Consente lo sviluppo e il test di algoritmi di precisione mista (mixed-precision) in ambiente MATLAB senza necessità di accesso immediato a hardware costoso o specifico.
• Standardizzazione: I dati raccolti e i modelli sviluppati contribuiscono a colmare il vuoto informativo sulla numerica dei Tensor Core, fornendo dati concreti che potrebbero influenzare future iniziative di standardizzazione (come IEEE P3109).
• Accessibilità: La disponibilità del codice open source democratizza la ricerca sui Tensor Core, permettendo a ricercatori senza accesso diretto a GPU di ultima generazione di studiare il loro comportamento numerico.

In sintesi, il paper rappresenta un passo fondamentale verso la trasparenza e la modellazione accurata dell'hardware per l'IA, trasformando i Tensor Core da "scatole nere" a componenti comprensibili e simulabili con precisione.