Skip to content
KitploitKITPLOIT
ToolsBlog
Einreichen
ToolsBlog
Einreichen

Hacking-, PenTest- und Cybersicherheits-Tools für Ihr Sicherheitsarsenal!

Kitploit ist ein Verzeichnis von Hacking-, Cybersicherheits- und Pentesting-Tools. Entdecken Sie die neuesten Projekt-Updates, um Schwachstellen zu finden, Systeme zu analysieren, Tests zu automatisieren und Ihre Sicherheit zu stärken.

··Feeds·Kontakt·Datenschutz·© 2026 Kitploit

Tool-Verzeichnis

Kategorien

Alle Kategorien anzeigen
Loading categories
sass-king — Reverse Engineering des NVIDIA SASS-Befehlsverzeichnisses, Kernel-Audits und Mustererkennung auf verschiedenen GPU-Architekturen. | Kitploit
Tools/GitHubGitHub/florianmattana/sass-king
Embedded-System-SicherheitStatische AnalyseCode-AnalyseReverse EngineeringHardware-SicherheitHardware- & IoT-SicherheitBinäranalysePapers & ForschungLernen & BildungKuratierte RessourcenFirmware-Analyse
31715vor 3 MonatenVon Kitploit geprüft

Beliebteste

Alle anzeigen →

Entdecken Sie die meistgenutzten Tools unserer Community.

Alle Tools erkunden

Durchsuchen Sie unsere Tool-Sammlung

Alle Tools anzeigen →
Teilen
GitHub
florianmattana/sass-king

sass-king

Reverse Engineering des NVIDIA SASS-Befehlsverzeichnisses, Kernel-Audits und Mustererkennung auf verschiedenen GPU-Architekturen.

Repository anzeigenWebseite

SASS King Logo

SASS King

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

Architecture Status License

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.

Schnelle Navigation

Das Repository ist als Beweispipeline organisiert:

root@kitploit:~
corpus/      kontrollierte Kernel und rohe SASS-Beweise
knowledge/   projektweite Erkenntnisse, Befehlsnotizen und Codierungsnotizen
patterns/    wiederverwendbare Phase-3-Audit-Signaturen
production/  Phase-4-Realkernel-Audits

Warum es das gibt

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.

Aktueller Stand

Phasen-3-Ergebnis

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:

  • die in den Kapiteln 01-25 gefundenen wiederholten Strukturen in 29 benannte Musterseiten überführt wurden;
  • jedes Muster eine verständliche Erklärung, SASS-Signatur, Varianten, Anti-Patterns, offene Lücken und ein Vertrauensniveau besitzt;
  • Anspruchs-Tags an die Quellnachweise in knowledge/FINDINGS.md gebunden bleiben;
  • die audit-orientierte Navigation nun von patterns/README.md ausgeht;
  • ungelöste Punkte explizit als Lücken weitergeführt werden, anstatt im Mustertext versteckt zu sein.

Jede Musterseite enthält:

  • verständliche Bedeutung;
  • SASS-Signatur;
  • beobachtete Varianten;
  • Interpretationsgrenzen;
  • Anti-Patterns;
  • offene Lücken;
  • Vertrauensniveau.

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.

Hier starten

  • Neu im Projekt: Lesen Sie Hier starten.
  • Möchten Sie die projektweite Karte: Lesen Sie das Wissensbasisverzeichnis.
  • Möchten Sie die aktuelle Befehlskarte: Lesen Sie SASS instructions on SM120 / SM120a.
  • Möchten Sie Codierungsnotizen: Lesen Sie encoding notes.
  • Möchten Sie die rohe Quelle der Wahrheit: Lesen Sie findings.
  • Möchten Sie wiederverwendbare Audit-Signaturen: Lesen Sie die Musterbibliothek.
  • Möchten Sie Tensor-Core-Beweise: Beginnen Sie mit tensor-core chapters.
  • Möchten Sie Dumps oder Korrekturen beitragen: Lesen Sie contributing.
  • Möchten Sie die v0.1-Grenze: Lesen Sie release notes.

Öffentliche Artikel:

  • Part 1 - Reading NVIDIA SASS from First Principles
  • Part 2 - Reading the Compiler's Mind

Methodik

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.

Was abgedeckt ist

Der erste Durchgang konzentriert sich auf die SM120-Tensor-Core- und Speicher-Pipeline:

  • HMMA, QMMA, OMMA
  • LDSM, STSM
  • LDGSTS, LDGDEPBAR, DEPBAR
  • LDG, STG, LDS, STS, REDG
  • BRA, 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.

Roadmap

root@kitploit:~
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;

Phase 1 - Lehrkernel

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.

Phase 2 - Tensor-Core- und SM120-Abdeckung

Kernel 13-25 decken den aktuellen SM120-Tensor-Core-Pfad ab:

Phase 2.5 - denvdis-Kreuzvalidierung

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.

Phase 3 - Musterbibliothek

Formalisierte wiederkehrende Strukturen in wiederverwendbare Signaturen:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • verkettetes HMMA / QMMA / OMMA
  • STSM -> BAR -> LDS -> STG
  • Warp-Reduktionen und fahrbahnübergreifende Kollektive
  • Register-Spill-Signaturen
  • skalare und uniforme Kontrollflussmuster

Die 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.

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:

  • einen realen Kernel in SASS-Regionen segmentiert;
  • passende PATTERN-NN-Seiten zitiert;
  • Vertrauensniveaus für jede Schlussfolgerung zuweist;
  • unerklärte Regionen als neue Lücken aufzeichnet;
  • den Bau eines Audit-Tools vermeidet, bis mindestens ein manueller Bericht stabil ist.

Phase 5 - Audit-Tool

Bauen Sie eine Pipeline, die ein cubin nimmt, bekannte Muster erkennt und einen optimierungsorientierten Bericht ausgibt.

Phase 6 - Architekturübergreifend

Wiederholen Sie die Methodik auf zusätzlichen Zielen:

Repository-Karte

root@kitploit:~
.
├── 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.

Werkzeuge

  • cuobjdump --dump-sass für rohe Disassemblierung.
  • gpuasm.com für Scoreboards, Stalls, Druck und Abhängigkeitspfeile.
  • Nsight Compute für Profiling und Stall-Zuordnung.
  • %clock-Mikrobenchmarks für Befehls-Latenzsonden.
  • nvcc -Xptxas -v für Register- und Spill-Metadaten.

Verwandte Arbeiten

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" für die empirische Mikrobenchmarking-Disziplin hinter Latenz, Durchsatz und Abhängigkeitsvalidierung.
  • kuterdinel.com/nv_isa für gefuzete NVIDIA-ISA-Codierungsarbeit, insbesondere die Idee, maschinenlesbare Codierungsregeln aus Disassembler-Verhalten abzuleiten.
  • redplait/denvdis für Opcode-Tabellen, Bit-Level-Disassemblierung, Codierungsfeld-Inspektion, Scheduling-Analyse, Registerverfolgung und Cubin-Manipulation. Die extrahierten SM120 data12-Tabellen werden als Low-Level-Kreuzprüfung verwendet, während lokale Dumps primäre Beweise bleiben.
  • Redplait-Werkzeuge und Notizen: ced cubin editor, SASS disassembly Perl bindings, SASS latency analysis und libcuda/nvasm_internal notes.
  • Huerta et al. 2025 für Reverse-Engineering des compilergeführten Schedulings, Control-Codes, Abhängigkeitszähler, Wiederverwendungs-Flags und Yield-Verhalten.
  • Yan et al. 2026 für die Treiberschicht-Start- und Pushbuffer-Analyse unterhalb von SASS.
  • MaxAS und TuringAS als frühere öffentliche SASS-Assembler-Bemühungen für ältere NVIDIA-Architekturen.
  • NVIDIA CUDA Binary Utilities Dokumentation für offizielle cubin-, fatbin- und Disassemblierungswerkzeuge.

SASS King arbeitet auf der algorithmischen Musterschicht: Erkennen, wie kompilierte Kernel strukturiert sind, und Verbinden dieser Strukturen mit Optimierungsentscheidungen auf Quellcode-Ebene.

Mitwirken

Beiträge sind willkommen, insbesondere:

  • rohe SASS-Dumps von Hardware, die hier nicht direkt verfügbar ist;
  • kontrollierte Kernelstudien, die eine Compiler-Entscheidung isolieren;
  • Korrekturen vorhandener Beobachtungen;
  • neue Produktionskernel-Mustervorschläge;
  • architekturübergreifende Vergleiche.

Siehe CONTRIBUTING.md für die erwarteten Metadaten und den Schreibstandard.

Autor

Florian Mattana. florianmattana.com

Tool herunterladen
Wenn Sie...Starten Sie hierDann lesen Sie
Das Projekt in 10 Minuten verstehen möchtendocs/README.mddocs/START_HERE.md, dann docs/PROJECT_STRUCTURE.md
Die Beweise reproduzieren möchtencorpus/README.mdein Kapitel conclusion*.md, dann dessen .sass-Dump
Die Quelle der Wahrheit finden möchtenknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md
Ein Muster in einem neuen Dump erkennen möchtenpatterns/README.mddie passende patterns/NN-*.md-Seite
Ein Produktions-Audit beginnen möchtenproduction/README.mdpassende PATTERN-NN-Seiten und Quellnachweise
Eine Korrektur oder einen Dump beitragen möchtenCONTRIBUTING.mddocs/START_HERE.md
BereichStatusWo
SM120-LehrkernelAbgeschlossen durch Kernel 01-12corpus/basics/01_vector_add/ bis corpus/math_and_spills/12_register_spill/
Tensor-Core-StudienAbgeschlossen bis Kernel 25corpus/tensor_cores/
Globale ErkenntnisseAktive Quelle der Wahrheitknowledge/FINDINGS.md
SM120-BefehlsglossarAktiv, evidenzgestütztknowledge/SASS_INSTRUCTIONS_SM120.md
CodierungspilotenGestartet mit LDSM, STSM, QMMAknowledge/encoding/
denvdis-KreuzvalidierungErste Durchführung abgeschlossen; tiefere Control-Code-Lücken bestehen nochknowledge/DENVDIS_INTEGRATION.md
MusterbibliothekErste Phase-3-Bibliothek abgeschlossenpatterns/
Produktions-AuditsNächste Phaseproduction/
MusterfamilieBeispieleWo
Tensor-Core-BerechnungHMMA-, QMMA-, OMMA-Akkumulator-Ketten; sparse Metadaten; schmale Fragmentepatterns/02-* bis patterns/04-*, patterns/10-*, patterns/21-*
Matrixspeicher und EpilogeLDSM, STSM, asynchrone Copy-Pipelines, REDG-Reduktions-Epilogepatterns/05-*, patterns/06-*, patterns/07-*, patterns/28-*
KontrollflussDivergenz/Reconvergenz, Schleifen-Rückkanten, prädizierte Ausgänge, Cold Traps, lokale CALLspatterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-*
Speicher und Registervektorisierter globaler Speicher, Spills, Shared-Memory-Staging, Deskriptoren, Uniform-Register-Flowpatterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-*
Arithmetik und SchedulingFFMA-Fusion, Konstanten, MUFU-Slowpaths, Scoreboards, Lifetime-Recyclingpatterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-*
Warp-KollektiveWarp-Reduktionen, Shuffle/Vote/Match/Sync-Primitivepatterns/01-*, patterns/25-*
TagBedeutung
[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.
BSYNC
WARPSYNC
  • SHFL, VOTE, REDUX
  • Uniform-Register-Flow: S2UR, R2UR, UMOV, ULEA, LDCU
  • PhaseStatusOutputWarum es wichtig ist
    1. LehrkernelAbgeschlossencorpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/Etabliert das Lesevokabular aus kontrollierten CUDA-zu-SASS-Experimenten.
    2. SM120-Tensor-Core-KorpusAbgeschlossencorpus/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-KreuzvalidierungErste Durchführung abgeschlossenknowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.mdVerwendet denvdis als Bit-Level-Kreuzprüfung, ohne lokale Dump-Beweise zu ersetzen. Vollständige Stall/Yield-Bit-Platzierung bleibt offen.
    3. MusterbibliothekErste Bibliothek abgeschlossenpatterns/Wandelt wiederholte Compiler/SASS-Strukturen in wiederverwendbare Signaturen um.
    4. Produktions-AuditsAls Nächstesproduction/Testet, ob Korpusmuster reale Kernel aus Produktionsbibliotheken erklären.
    5. Audit-ToolGeplantcubin-to-report-PipelineMacht die Musterschicht skriptbar und wiederholbar.
    6. Architekturübergreifende WiederholungGeplantSM80-, SM86-, SM89-, SM90a-, SM100a-, SM120-VergleicheTrennt architekturspezifische Fakten von allgemeinem NVIDIA-SASS-Verhalten.
    KernelThema
    13HMMA-Baseline, Registerallokation, Akkumulator-Verkettung
    14QMMA FP8 / FP6 / FP4-Baseline
    15Schmale MMA-Varianten
    16FP4-Spitze und block-skalierte OMMA/QMMA
    17LDSM und Matrix-Load-Verhalten
    18Pipelined MMA-Tile und asynchrones Copy-Staging
    19Sparse-MMA-Metadaten
    20Kontrollfluss und Rückkantenerkennung
    21Divergenz und Reconvergenz
    22STSM-Matrix-Store-Verhalten
    23FP4 / FP6-Fragmentlayout-Sonden
    24Produktions-Mini-GEMM-Audit
    25STSM-Epilog-Layout und Storeback-Semantik
    ArchitekturRepräsentative GPUWarum
    SM80A100Datacenter-Ampere-Baseline
    SM86RTX 3090Consumer-Ampere-Korpus
    SM89RTX 4090Häufige Consumer-Inference-Karte
    SM90aH100TMA, WGMMA, Warp-Spezialisierung, Cluster
    SM100aB200tcgen05.mma, TMEM
    SM120RTX 5070 Ti / 5090Consumer-Blackwell-Startpunkt