
Reverse engineering del dizionario delle istruzioni NVIDIA SASS, audit dei kernel e riconoscimento di pattern attraverso le architetture GPU.
Reverse engineering di NVIDIA SASS da kernel controllati a audit di produzione.
Articolo 1 · Articolo 2 · Base di conoscenza · Libreria di pattern · Glossario istruzioni SM120 · Note di codifica · Inizia qui · Struttura del progetto · Capitoli tensor core · Contribuire
SASS King è un progetto sistematico di reverse engineering per NVIDIA SASS, il set di istruzioni nativo della GPU emesso nei binari CUDA compilati. Il progetto parte dall'hardware consumer Blackwell SM120 / SM120a e si espande nel tempo verso una libreria completa di ISA e pattern multi-architettura.
L'obiettivo è pratico: aiutare un ingegnere di kernel ad aprire un dump SASS, riconoscere pattern del compilatore, identificare strutture rilevanti per le prestazioni e collegare il binario alle decisioni di ottimizzazione a livello di sorgente.
Il progetto ha completato la libreria di pattern iniziale della Fase 3: 29 firme SASS riutilizzabili sono ora formalizzate in patterns/, con knowledge/FINDINGS.md mantenuto come traccia completa delle evidenze. Il prossimo passo importante è la Fase 4: applicare questi pattern a kernel di produzione reali.
Il repository è organizzato come una pipeline di evidenze:
corpus/ kernel controllati e prove SASS grezze
knowledge/ risultati a livello di progetto, note sulle istruzioni e note di codifica
patterns/ firme di audit riutilizzabili della Fase 3
production/ audit su kernel reali della Fase 4
L'ultimo lavoro pubblico di reverse engineering SASS ampio e paragonabile per spirito è stato Jia et al. su Volta e Turing nel 2018. Ampere, Hopper e Blackwell hanno cambiato sostanzialmente il mix di istruzioni: percorsi di copia asincrona, famiglie di tensor core, istruzioni di load/store matriciale, forme MMA sparse e scalate, e nuovi flussi di registri uniformi.
SASS King colma questa lacuna combinando micro-kernel controllati, lettura SASS grezza, sonde runtime e audit di kernel di produzione.
La libreria formale di pattern è il principale output della Fase 3. Trasforma le evidenze locali dei capitoli in firme di audit riutilizzabili, in modo che un audit possa citare un pattern con nome invece di riscrivere ogni volta l'intera traccia di ricerca.
La Fase 3 è considerata completa perché:
knowledge/FINDINGS.md;patterns/README.md;Ogni pagina di pattern include:
Usa patterns/README.md come indice orientato all'audit. Usa knowledge/FINDINGS.md quando hai bisogno del contesto di ricerca più lungo dietro un pattern.
La Fase 3 non sostiene che ogni comportamento di NVIDIA SASS sia decodificato. Stabilisce un livello di pattern SM120 / SM120a riutilizzabile, sufficientemente buono per iniziare audit manuali di produzione. La decodifica del layout runtime, il posizionamento completo dei bit dei codici di controllo, la reportistica cubin automatizzata e la riproduzione cross-architettura rimangono lavoro futuro.
Articoli pubblici:
Variazione controllata. Due kernel differiscono per esattamente una variabile: dtype, ordine degli operandi, fattore di unrolling, layout di memoria o target di compilazione. Il diff SASS isola la decisione del compilatore.
Tag di affermazione stretti. Ogni affermazione tecnica usa un tag:
Top-down e bottom-up insieme. I micro-kernel isolano singole istruzioni e decisioni del compilatore. I kernel simili a produzione mostrano quali pattern contano nel codice reale.
Audit basati su pattern. Un audit di produzione dovrebbe citare una pagina formale PATTERN-NN solo dopo aver abbinato la firma SASS visibile e averne portato i limiti di confidenza, gli anti-pattern e le lacune aperte.
Il primo passaggio si concentra sul tensor core SM120 e sulla pipeline di memoria:
HMMA, QMMA, OMMALDSM, STSMLDGSTS, LDGDEPBAR, DEPBARLDG, STG, LDS, STS, REDGBRA, EXIT, BSSY, , Il progetto non pretende che l'ISA sia ancora completa. Il glossario pubblico tiene traccia di ciò che è osservato e spiegato; pagine più approfondite sotto knowledge/encoding/ tengono traccia delle famiglie con sufficienti evidenze per una documentazione in stile matcher.
SASS King non compete con i disassemblatori SASS a livello di bit. Il progetto usa dump locali come evidenza primaria e può usare redplait/denvdis come verifica incrociata per campi delle istruzioni, tabelle di scheduling, predicati e tracciamento dei registri. denvdis può convalidare interpretazioni di codifica a basso livello; SASS King possiede le evidenze a variazione controllata, il livello di pattern semantico e l'interpretazione dell'audit di produzione.
flowchart LR
P1["Phase 1<br/>Teaching kernels<br/>01-12"] --> P2["Phase 2<br/>SM120 tensor-core corpus<br/>13-25"]
P2 --> P25["Phase 2.5<br/>denvdis cross-validation<br/>bit-level backend"]
P25 --> P3["Phase 3<br/>Pattern library<br/>compiler signatures"]
P3 --> P4["Phase 4<br/>Production audits<br/>real kernels"]
P4 --> P5["Phase 5<br/>Audit tool<br/>cubin reports"]
P5 --> P6["Phase 6<br/>Cross-architecture replay<br/>SM80/86/89/90a/100a/120"]
classDef done fill:#0b6d55,color:#fff,stroke:#0b6d55;
classDef active fill:#f4c95d,color:#111,stroke:#b89422;
classDef planned fill:#1f2937,color:#fff,stroke:#6b7280;
class P1,P2,P25,P3 done;
class P4 active;
class P5,P6 planned;
I kernel 01-12 stabiliscono concetti SASS di base: fusione FMA, comportamento scoreboard, abbassamento dei loop, memoria condivisa, memoria globale, primitive warp, matematica a percorso lento e spill di memoria locale.
I kernel 13-25 coprono il percorso tensor core corrente di SM120:
Convalidare redplait/denvdis come backend di verifica incrociata a livello di bit per SM120 / SM120a prima che gli audit di produzione dipendano dalla libreria di pattern. Il passaggio esegue nvd -O, nvd -S, nvd -p e, dove utile, nvd -T su cubin o dump locali rappresentativi che coprono HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR e marcatori di divergenza.
L'output è knowledge/DENVDIS_INTEGRATION.md: una tabella di compatibilità fattuale dalla famiglia allo stato di riconoscimento denvdis, copertura dei modificatori, campi dei codici di controllo esposti e azione di SASS King. L'output di denvdis è evidenza di supporto, non un sostituto delle osservazioni dei dump locali.
Strutture ricorrenti formalizzate in firme riutilizzabili:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMA concatenatiSTSM -> BAR -> LDS -> STGLa libreria iniziale della Fase 3 contiene 29 pagine di pattern sotto patterns/. knowledge/FINDINGS.md rimane il registro di ricerca e la fonte di verità; patterns/ è il punto di ingresso orientato all'audit.
La Fase 3 è completa al livello della libreria iniziale. Gli elementi rimanenti come la decodifica del layout runtime, il posizionamento completo dei bit dei codici di controllo e la riproduzione cross-architettura sono tracciati come lacune o fasi future, non come blocchi per iniziare gli audit di produzione della Fase 4.
Applicare la libreria di pattern a kernel reali da librerie come FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad e progetti correlati. L'obiettivo è una copertura rappresentativa per pattern algoritmico, non un file markdown per kernel.
Il primo risultato della Fase 4 dovrebbe essere un rapporto di audit manuale che:
PATTERN-NN corrispondenti;Costruire una pipeline che prende un cubin, rileva pattern noti ed emette un report orientato all'ottimizzazione.
Riprodurre la metodologia su target aggiuntivi:
.
├── corpus/ # Kernel controllati, dump e articoli dei capitoli
│ ├── basics/ # Kernel 01-08: basi scalari/vettoriali e memoria
│ ├── warp_collectives/ # Kernel 09-10: shuffle, vote, riduzione
│ ├── math_and_spills/ # Kernel 11-12: percorsi lenti e spill
│ └── tensor_cores/ # Kernel 13-25: studi sui tensor core
├── knowledge/ # Risultati, glossario, note di codifica
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # Libreria formale di pattern Fase 3
├── production/ # Audit di kernel di produzione Fase 4
├── docs/ # Onboarding, struttura e note per il rilascio
└── guide/ # Sottomodulo guida esterna alla lettura SASS
Ogni cartella del capitolo contiene kernel sorgente, artefatti compilati quando pertinenti, dump SASS quando fanno parte del set di evidenze convalidate e un articolo conclusion<N>.md.
Per una spiegazione più completa di cosa appartiene a ogni directory, leggi Struttura del Progetto.
cuobjdump --dump-sass per il disassemblaggio grezzo.gpuasm.com per scoreboard, stall, pressione e frecce di dipendenza.%clock per sonde di latenza delle istruzioni.nvcc -Xptxas -v per metadati di registri e spill.SASS King opera al livello dei pattern algoritmici: riconoscere come sono strutturati i kernel compilati e collegare quelle strutture alle decisioni di ottimizzazione a livello di sorgente.
I contributi sono benvenuti, specialmente:
Vedi CONTRIBUTING.md per i metadati attesi e lo standard di scrittura.
Florian Mattana. florianmattana.com
| Se vuoi... | Inizia qui | Poi leggi |
|---|
| Capire il progetto in 10 minuti | docs/README.md | docs/START_HERE.md, poi docs/PROJECT_STRUCTURE.md |
| Riprodurre le evidenze | corpus/README.md | un capitolo conclusion*.md, poi il suo dump .sass |
| Trovare la fonte di verità | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md |
| Riconoscere un pattern in un nuovo dump | patterns/README.md | la pagina patterns/NN-*.md corrispondente |
| Iniziare un audit di produzione | production/README.md | pagine PATTERN-NN corrispondenti e prove sorgente |
| Contribuire con una correzione o un dump | CONTRIBUTING.md | docs/START_HERE.md |
| Area | Stato | Dove |
|---|
| Kernel didattici SM120 | Completati attraverso i kernel 01-12 | corpus/basics/01_vector_add/ a corpus/math_and_spills/12_register_spill/ |
| Studi sui tensor core | Completati attraverso il Kernel 25 | corpus/tensor_cores/ |
| Risultati globali | Fonte di verità attiva | knowledge/FINDINGS.md |
| Glossario istruzioni SM120 | Attivo, basato su evidenze | knowledge/SASS_INSTRUCTIONS_SM120.md |
| Piloti di codifica | Avviati con LDSM, STSM, QMMA | knowledge/encoding/ |
| Convalida incrociata denvdis | Primo passaggio completato; rimangono lacune nei codici di controllo più profondi | knowledge/DENVDIS_INTEGRATION.md |
| Libreria di pattern | Libreria iniziale Fase 3 completata | patterns/ |
| Audit di produzione | Prossima fase | production/ |
| Famiglia di pattern | Esempi | Dove |
|---|
| Calcolo tensor core | Catene di accumulatori HMMA, QMMA, OMMA; metadati sparsi; frammenti stretti | patterns/02-* a patterns/04-*, patterns/10-*, patterns/21-* |
| Memoria matriciale ed epiloghi | LDSM, STSM, pipeline di copia asincrona, epiloghi di riduzione REDG | patterns/05-*, patterns/06-*, patterns/07-*, patterns/28-* |
| Flusso di controllo | divergenza/riconvergenza, back-edge di loop, uscite predicate, cold trap, CALL locali | patterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-* |
| Memoria e registri | memoria globale vettorizzata, spill, staging in memoria condivisa, descrittori, flusso di registri uniformi | patterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-* |
| Aritmetica e scheduling | fusione FFMA, costanti, percorsi lenti MUFU, scoreboard, riciclo lifetime | patterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-* |
| Collettivi warp | riduzioni warp, shuffle/vote/match/sync primitive | patterns/01-*, patterns/25-* |
| Tag | Significato |
|---|
[OBS] | Osservato direttamente in un dump, log, output runtime o profilo. |
[INF] | Dedotto da evidenze osservate. |
[HYP] | Plausibile ma non confermato. |
[RES] | Un'ipotesi precedente risolta da evidenze successive. |
[GAP] | Domanda aperta documentata esplicitamente. |
BSYNCWARPSYNCSHFL, VOTE, REDUXS2UR, R2UR, UMOV, ULEA, LDCU| Fase | Stato | Output | Perché è importante |
|---|
| 1. Kernel didattici | Completato | corpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/ | Stabilisce il vocabolario di lettura da esperimenti CUDA-a-SASS controllati. |
| 2. Corpus tensor core SM120 | Completato | corpus/tensor_cores/13_hmma_fp16/ a 25_stsm_epilogue/ | Cattura il primo set di evidenze per tensor core SM120 / SM120a, memoria matriciale, flusso di controllo ed epiloghi. |
| 2.5. Convalida incrociata denvdis | Primo passaggio completato | knowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.md | Usa denvdis come verifica incrociata a livello di bit senza sostituire le evidenze dei dump locali. Il posizionamento completo dei bit di stall/yield rimane aperto. |
| 3. Libreria di pattern | Libreria iniziale completata | patterns/ | Trasforma strutture SASS/compilatore ripetute in firme riutilizzabili. |
| 4. Audit di produzione | Prossimo | production/ | Verifica se i pattern del corpus spiegano kernel reali da librerie di produzione. |
| 5. Strumento di audit | Pianificato | pipeline da cubin a report | Rende il livello di pattern scriptabile e ripetibile. |
| 6. Riproduzione cross-architettura | Pianificato | Confronti SM80, SM86, SM89, SM90a, SM100a, SM120 | Separa i fatti specifici dell'architettura dal comportamento generale di NVIDIA SASS. |
| Kernel | Argomento |
|---|
| 13 | Baseline HMMA, allocazione registri, concatenamento accumulatori |
| 14 | Baseline QMMA FP8 / FP6 / FP4 |
| 15 | Varianti MMA strette |
| 16 | Picco FP4 e OMMA/QMMA con blocco scalato |
| 17 | LDSM e comportamento del load matriciale |
| 18 | Tile MMA pipelined e staging di copia asincrona |
| 19 | Metadati MMA sparsi |
| 20 | Flusso di controllo e rilevamento back-edge |
| 21 | Divergenza e riconvergenza |
| 22 | Comportamento dello store matriciale STSM |
| 23 | Sonde di layout frammento FP4 / FP6 |
| 24 | Audit mini-GEMM di produzione |
| 25 | Layout epilogo STSM e semantica storeback |
| Arch | GPU rappresentativa | Perché |
|---|
| SM80 | A100 | Baseline datacenter Ampere |
| SM86 | RTX 3090 | Corpus Ampere consumer |
| SM89 | RTX 4090 | Scheda comune per inferenza consumer |
| SM90a | H100 | TMA, WGMMA, specializzazione warp, cluster |
| SM100a | B200 | tcgen05.mma, TMEM |
| SM120 | RTX 5070 Ti / 5090 | Punto di partenza Blackwell consumer |