
Reverse Engineering des NVIDIA SASS-Befehlsverzeichnisses, Kernel-Audits und Mustererkennung auf verschiedenen GPU-Architekturen.
Reverse Engineering von NVIDIA SASS – von kontrollierten Kernels bis zu Produktions-Audits.
Artikel 1 · Artikel 2 · Wissensbasis · Musterbibliothek · SM120-Befehlsglossar · Codierungsnotizen · Hier starten · Projektstruktur · Tensor-Core-Kapitel · Mitwirken
SASS King ist ein systematisches Reverse-Engineering-Projekt für NVIDIA SASS, den nativen GPU-Befehlssatz, der in kompilierten CUDA-Binärdateien enthalten ist. Das Projekt startet mit SM120 / SM120a Consumer-Blackwell-Hardware und erweitert sich im Laufe der Zeit zu einer vollständigen architekturübergreifenden ISA- und Musterbibliothek.
Das Ziel ist praktisch: einem Kernel-Ingenieur helfen, einen SASS-Dump zu öffnen, Compiler-Muster zu erkennen, leistungsrelevante Strukturen zu identifizieren und die Binärdatei mit Optimierungsentscheidungen auf Quellcode-Ebene zu verknüpfen.
Das Projekt hat seine erste Phase-3-Musterbibliothek abgeschlossen: 29 wiederverwendbare SASS-Signaturen sind nun unter patterns/ formalisiert, mit knowledge/FINDINGS.md als vollständiger Beweiskette. Der nächste große Schritt ist Phase 4: die Anwendung dieser Muster auf echte Produktions-Kernel.
Das Repository ist als Beweispipeline organisiert:
corpus/ kontrollierte Kernel und rohe SASS-Beweise
knowledge/ projektweite Erkenntnisse, Befehlsnotizen und Codierungsnotizen
patterns/ wiederverwendbare Phase-3-Audit-Signaturen
production/ Phase-4-Realkernel-Audits
Die letzte breite öffentliche SASS-Reverse-Engineering-Arbeit mit vergleichbarem Geist war Jia et al. zu Volta und Turing im Jahr 2018. Ampere, Hopper und Blackwell haben die Befehlsmischung erheblich verändert: asynchrone Copy-Pfade, Tensor-Core-Familien, Matrix-Load/Store-Befehle, sparse und skalierte MMA-Formen sowie neue Uniform-Register-Flows.
SASS King schließt diese Lücke, indem es kontrollierte Mikro-Kernel, reines SASS-Lesen, Runtime-Probes und Produktionskernel-Audits kombiniert.
Die formale Musterbibliothek ist das Hauptprodukt von Phase 3. Sie wandelt die kapitelspezifischen Beweise in wiederverwendbare Audit-Signaturen um, sodass ein Audit ein benanntes Muster zitieren kann, anstatt jedes Mal die vollständige Recherchespur neu schreiben zu müssen.
Phase 3 gilt als abgeschlossen, weil:
knowledge/FINDINGS.md gebunden bleiben;patterns/README.md ausgeht;Jede Musterseite enthält:
Verwenden Sie patterns/README.md als audit-orientierten Index. Verwenden Sie knowledge/FINDINGS.md, wenn Sie den längeren Forschungskontext hinter einem Muster benötigen.
Phase 3 behauptet nicht, dass jedes NVIDIA-SASS-Verhalten dekodiert ist. Sie erstellt eine wiederverwendbare SM120 / SM120a-Musterschicht, die gut genug ist, um manuelle Produktions-Audits zu beginnen. Runtime-Layout-Dekodierung, vollständige Control-Code-Bit-Platzierung, automatisiertes cubin-Reporting und architekturübergreifende Wiederholung bleiben zukünftige Arbeiten.
Öffentliche Artikel:
Kontrollierte Variation. Zwei Kernel unterscheiden sich um genau eine Variable: dtype, Operandenreihenfolge, Entrollfaktor, Speicherlayout oder Kompilierungsziel. Der SASS-Diff isoliert die Compiler-Entscheidung.
Strenge Anspruchs-Tags. Jeder technische Anspruch verwendet ein Tag:
Top-down und Bottom-up zusammen. Mikro-Kernel isolieren einzelne Befehle und Compiler-Entscheidungen. Produktionsähnliche Kernel zeigen, welche Muster in echtem Code relevant sind.
Muster-zuerst-Audits. Ein Produktions-Audit sollte eine formale PATTERN-NN-Seite nur dann zitieren, wenn die sichtbare SASS-Signatur übereinstimmt und ihre Vertrauensgrenzen, Anti-Patterns und offenen Lücken übernommen werden.
Der erste Durchgang konzentriert sich auf die SM120-Tensor-Core- und Speicher-Pipeline:
HMMA, QMMA, OMMALDSM, STSMLDGSTS, LDGDEPBAR, DEPBARLDG, STG, LDS, STS, REDGBRA, EXIT, BSSY, , Das Projekt gibt nicht vor, dass die ISA bereits vollständig ist. Das öffentliche Glossar verfolgt, was beobachtet und erklärt ist; tiefere Seiten unter knowledge/encoding/ verfolgen Familien mit genügend Beweisen für eine matcher-artige Dokumentation.
SASS King konkurriert nicht mit Bit-Level-SASS-Disassemblierern. Das Projekt verwendet lokale Dumps als primäre Beweise und kann redplait/denvdis als Kreuzprüfung für Befehlsfelder, Scheduling-Tabellen, Prädikate und Registerverfolgung einsetzen. denvdis kann Low-Level-Codierungsinterpretationen validieren; SASS King besitzt die kontrollierte Variationsbeweise, die semantische Musterschicht und die Produktions-Audit-Interpretation.
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;
Kernel 01-12 etablieren grundlegende SASS-Konzepte: FMA-Fusion, Scoreboard-Verhalten, Schleifensenkung, Shared Memory, Global Memory, Warp-Primitive, Slow-Path-Mathematik und Local-Memory-Spills.
Kernel 13-25 decken den aktuellen SM120-Tensor-Core-Pfad ab:
Validieren Sie redplait/denvdis als Bit-Level-Kreuzprüfungs-Backend für SM120 / SM120a, bevor Produktions-Audits von der Musterbibliothek abhängen. Der Durchlauf führt nvd -O, nvd -S, nvd -p und wo nützlich nvd -T auf repräsentativen lokalen Cubins oder Dumps durch, die HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR und Divergenzmarkierungen abdecken.
Das Ergebnis ist knowledge/DENVDIS_INTEGRATION.md: eine faktenbasierte Kompatibilitätstabelle von Familie zu denvdis-Erkennungsstatus, Modifikatorabdeckung, exponierten Control-Code-Feldern und der SASS-King-Aktion. denvdis-Ausgabe ist unterstützender Beweis, kein Ersatz für lokale Dump-Beobachtungen.
Formalisierte wiederkehrende Strukturen in wiederverwendbare Signaturen:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMASTSM -> BAR -> LDS -> STGDie erste Phase-3-Bibliothek enthält 29 Musterseiten unter patterns/. knowledge/FINDINGS.md bleibt das Forschungstagebuch und die Quelle der Wahrheit; patterns/ ist der audit-orientierte Einstiegspunkt.
Phase 3 ist auf der Ebene der ersten Bibliothek abgeschlossen. Verbleibende Punkte wie Runtime-Layout-Dekodierung, vollständige Control-Code-Bit-Platzierung und architekturübergreifende Wiederholung werden als Lücken oder zukünftige Phasen verfolgt, nicht als Blocker für den Start von Phase-4-Produktions-Audits.
Wenden Sie die Musterbibliothek auf reale Kernel aus Bibliotheken wie FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad und verwandten Projekten an. Ziel ist eine repräsentative Abdeckung nach algorithmischem Muster, nicht eine Markdown-Datei pro Kernel.
Das erste Phase-4-Ergebnis sollte ein manueller Audit-Bericht sein, der:
PATTERN-NN-Seiten zitiert;Bauen Sie eine Pipeline, die ein cubin nimmt, bekannte Muster erkennt und einen optimierungsorientierten Bericht ausgibt.
Wiederholen Sie die Methodik auf zusätzlichen Zielen:
.
├── corpus/ # Kontrollierte Kernel, Dumps und Kapitel-Ausarbeitungen
│ ├── basics/ # Kernel 01-08: Skalar/Vector und Speichergrundlagen
│ ├── warp_collectives/ # Kernel 09-10: Shuffle, Vote, Reduction
│ ├── math_and_spills/ # Kernel 11-12: Slow-Paths und Spills
│ └── tensor_cores/ # Kernel 13-25: Tensor-Core-Studien
├── knowledge/ # Erkenntnisse, Glossar, Codierungsnotizen
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # Formale Phase-3-Musterbibliothek
├── production/ # Phase-4-Produktionskernel-Audits
├── docs/ # Onboarding, Struktur und versionierungsbezogene Notizen
└── guide/ # Externer SASS-Leseleitfaden als Submodul
Jeder Kapitelordner enthält Quellkernel, kompilierte Artefakte (falls relevant), SASS-Dumps (wenn sie Teil des validierten Beweissatzes sind) und eine conclusion<N>.md-Ausarbeitung.
Für eine ausführlichere Erklärung, was in jedes Verzeichnis gehört, lesen Sie Projektstruktur.
cuobjdump --dump-sass für rohe Disassemblierung.gpuasm.com für Scoreboards, Stalls, Druck und Abhängigkeitspfeile.%clock-Mikrobenchmarks für Befehls-Latenzsonden.nvcc -Xptxas -v für Register- und Spill-Metadaten.SASS King arbeitet auf der algorithmischen Musterschicht: Erkennen, wie kompilierte Kernel strukturiert sind, und Verbinden dieser Strukturen mit Optimierungsentscheidungen auf Quellcode-Ebene.
Beiträge sind willkommen, insbesondere:
Siehe CONTRIBUTING.md für die erwarteten Metadaten und den Schreibstandard.
Florian Mattana. florianmattana.com
| Wenn Sie... | Starten Sie hier | Dann lesen Sie |
|---|
| Das Projekt in 10 Minuten verstehen möchten | docs/README.md | docs/START_HERE.md, dann docs/PROJECT_STRUCTURE.md |
| Die Beweise reproduzieren möchten | corpus/README.md | ein Kapitel conclusion*.md, dann dessen .sass-Dump |
| Die Quelle der Wahrheit finden möchten | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md |
| Ein Muster in einem neuen Dump erkennen möchten | patterns/README.md | die passende patterns/NN-*.md-Seite |
| Ein Produktions-Audit beginnen möchten | production/README.md | passende PATTERN-NN-Seiten und Quellnachweise |
| Eine Korrektur oder einen Dump beitragen möchten | CONTRIBUTING.md | docs/START_HERE.md |
| Bereich | Status | Wo |
|---|
| SM120-Lehrkernel | Abgeschlossen durch Kernel 01-12 | corpus/basics/01_vector_add/ bis corpus/math_and_spills/12_register_spill/ |
| Tensor-Core-Studien | Abgeschlossen bis Kernel 25 | corpus/tensor_cores/ |
| Globale Erkenntnisse | Aktive Quelle der Wahrheit | knowledge/FINDINGS.md |
| SM120-Befehlsglossar | Aktiv, evidenzgestützt | knowledge/SASS_INSTRUCTIONS_SM120.md |
| Codierungspiloten | Gestartet mit LDSM, STSM, QMMA | knowledge/encoding/ |
| denvdis-Kreuzvalidierung | Erste Durchführung abgeschlossen; tiefere Control-Code-Lücken bestehen noch | knowledge/DENVDIS_INTEGRATION.md |
| Musterbibliothek | Erste Phase-3-Bibliothek abgeschlossen | patterns/ |
| Produktions-Audits | Nächste Phase | production/ |
| Musterfamilie | Beispiele | Wo |
|---|
| Tensor-Core-Berechnung | HMMA-, QMMA-, OMMA-Akkumulator-Ketten; sparse Metadaten; schmale Fragmente | patterns/02-* bis patterns/04-*, patterns/10-*, patterns/21-* |
| Matrixspeicher und Epiloge | LDSM, STSM, asynchrone Copy-Pipelines, REDG-Reduktions-Epiloge | patterns/05-*, patterns/06-*, patterns/07-*, patterns/28-* |
| Kontrollfluss | Divergenz/Reconvergenz, Schleifen-Rückkanten, prädizierte Ausgänge, Cold Traps, lokale CALLs | patterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-* |
| Speicher und Register | vektorisierter globaler Speicher, Spills, Shared-Memory-Staging, Deskriptoren, Uniform-Register-Flow | patterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-* |
| Arithmetik und Scheduling | FFMA-Fusion, Konstanten, MUFU-Slowpaths, Scoreboards, Lifetime-Recycling | patterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-* |
| Warp-Kollektive | Warp-Reduktionen, Shuffle/Vote/Match/Sync-Primitive | patterns/01-*, patterns/25-* |
| Tag | Bedeutung |
|---|
[OBS] | Direkt in einem Dump, Log, Laufzeitoutput oder Profil beobachtet. |
[INF] | Aus beobachteten Beweisen abgeleitet. |
[HYP] | Plausibel, aber nicht bestätigt. |
[RES] | Eine frühere Hypothese, die durch spätere Beweise aufgelöst wurde. |
[GAP] | Offene Frage explizit dokumentiert. |
BSYNCWARPSYNCSHFL, VOTE, REDUXS2UR, R2UR, UMOV, ULEA, LDCU| Phase | Status | Output | Warum es wichtig ist |
|---|
| 1. Lehrkernel | Abgeschlossen | corpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/ | Etabliert das Lesevokabular aus kontrollierten CUDA-zu-SASS-Experimenten. |
| 2. SM120-Tensor-Core-Korpus | Abgeschlossen | corpus/tensor_cores/13_hmma_fp16/ bis 25_stsm_epilogue/ | Erfasst den ersten SM120 / SM120a Tensor-Core-, Matrix-Speicher-, Kontrollfluss- und Epilog-Beweissatz. |
| 2.5. denvdis-Kreuzvalidierung | Erste Durchführung abgeschlossen | knowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.md | Verwendet denvdis als Bit-Level-Kreuzprüfung, ohne lokale Dump-Beweise zu ersetzen. Vollständige Stall/Yield-Bit-Platzierung bleibt offen. |
| 3. Musterbibliothek | Erste Bibliothek abgeschlossen | patterns/ | Wandelt wiederholte Compiler/SASS-Strukturen in wiederverwendbare Signaturen um. |
| 4. Produktions-Audits | Als Nächstes | production/ | Testet, ob Korpusmuster reale Kernel aus Produktionsbibliotheken erklären. |
| 5. Audit-Tool | Geplant | cubin-to-report-Pipeline | Macht die Musterschicht skriptbar und wiederholbar. |
| 6. Architekturübergreifende Wiederholung | Geplant | SM80-, SM86-, SM89-, SM90a-, SM100a-, SM120-Vergleiche | Trennt architekturspezifische Fakten von allgemeinem NVIDIA-SASS-Verhalten. |
| Kernel | Thema |
|---|
| 13 | HMMA-Baseline, Registerallokation, Akkumulator-Verkettung |
| 14 | QMMA FP8 / FP6 / FP4-Baseline |
| 15 | Schmale MMA-Varianten |
| 16 | FP4-Spitze und block-skalierte OMMA/QMMA |
| 17 | LDSM und Matrix-Load-Verhalten |
| 18 | Pipelined MMA-Tile und asynchrones Copy-Staging |
| 19 | Sparse-MMA-Metadaten |
| 20 | Kontrollfluss und Rückkantenerkennung |
| 21 | Divergenz und Reconvergenz |
| 22 | STSM-Matrix-Store-Verhalten |
| 23 | FP4 / FP6-Fragmentlayout-Sonden |
| 24 | Produktions-Mini-GEMM-Audit |
| 25 | STSM-Epilog-Layout und Storeback-Semantik |
| Architektur | Repräsentative GPU | Warum |
|---|
| SM80 | A100 | Datacenter-Ampere-Baseline |
| SM86 | RTX 3090 | Consumer-Ampere-Korpus |
| SM89 | RTX 4090 | Häufige Consumer-Inference-Karte |
| SM90a | H100 | TMA, WGMMA, Warp-Spezialisierung, Cluster |
| SM100a | B200 | tcgen05.mma, TMEM |
| SM120 | RTX 5070 Ti / 5090 | Consumer-Blackwell-Startpunkt |