
Ingénierie inverse du dictionnaire d'instructions NVIDIA SASS, audits de noyaux et reconnaissance de motifs à travers les architectures GPU.
Rétro-ingénierie du SASS NVIDIA, des kernels contrôlés aux audits de production.
Article 1 · Article 2 · Base de connaissances · Bibliothèque de motifs · Glossaire des instructions SM120 · Notes sur l'encodage · Commencez ici · Structure du projet · Chapitres sur les tensor-cores · Contribuer
SASS King est un projet systématique de rétro-ingénierie du SASS NVIDIA, le jeu d'instructions natif du GPU émis dans les binaires CUDA compilés. Le projet commence avec le matériel Blackwell grand public SM120 / SM120a et s'étend progressivement vers une bibliothèque complète d'ISA et de motifs multi-architectures.
L'objectif est pratique : aider un ingénieur de kernel à ouvrir un dump SASS, reconnaître les motifs du compilateur, identifier les structures pertinentes pour la performance, et relier le binaire aux décisions d'optimisation au niveau du source.
Le projet a terminé sa phase initiale de bibliothèque de motifs Phase 3 : 29 signatures SASS réutilisables sont maintenant formalisées sous patterns/, avec knowledge/FINDINGS.md conservé comme la trace complète des preuves. La prochaine étape majeure est la Phase 4 : appliquer ces motifs à des kernels de production réels.
Le dépôt est organisé comme un pipeline de preuves :
corpus/ kernels contrôlés et preuves SASS brutes
knowledge/ découvertes à l'échelle du projet, notes d'instructions et notes d'encodage
patterns/ signatures d'audit réutilisables de la Phase 3
production/ audits de kernels réels de la Phase 4
Le dernier grand travail public de rétro-ingénierie SASS comparable en esprit a été celui de Jia et al. sur Volta et Turing en 2018. Ampere, Hopper et Blackwell ont considérablement modifié le mélange d'instructions : chemins de copie asynchrones, familles de tensor-cores, instructions de chargement/stockage de matrices, formes MMA éparses et à échelle, et nouveaux flux de registres uniformes.
SASS King comble cette lacune en combinant des micro-kernels contrôlés, la lecture brute du SASS, des sondes d'exécution et des audits de kernels de production.
La bibliothèque formelle de motifs est le principal résultat de la Phase 3. Elle transforme les preuves locales des chapitres en signatures d'audit réutilisables, afin qu'un audit puisse citer un motif nommé au lieu de réécrire la piste de recherche complète à chaque fois.
La Phase 3 est considérée comme terminée parce que :
knowledge/FINDINGS.md ;patterns/README.md ;Chaque page de motif comprend :
Utilisez patterns/README.md comme index orienté audit. Utilisez knowledge/FINDINGS.md lorsque vous avez besoin du contexte de recherche plus long derrière un motif.
La Phase 3 ne prétend pas que chaque comportement SASS de NVIDIA est décodé. Elle établit une couche de motifs SM120 / SM120a réutilisable suffisamment bonne pour commencer des audits manuels de production. Le décodage de la disposition d'exécution, le placement complet des bits de code de contrôle, les rapports automatisés de cubin et la relecture multi-architecture restent des travaux futurs.
Articles publics :
Variation contrôlée. Deux kernels diffèrent par exactement une variable : dtype, ordre des opérandes, facteur de déroulage, disposition mémoire ou cible de compilation. Le diff SASS isole la décision du compilateur.
Étiquettes de revendication strictes. Chaque revendication technique utilise une étiquette :
Top-down et bottom-up ensemble. Les micro-kernels isolent les instructions individuelles et les décisions du compilateur. Les kernels de type production montrent quels motifs sont importants dans le code réel.
Audits orientés motifs. Un audit de production doit citer une page formelle PATTERN-NN seulement après avoir fait correspondre la signature SASS visible et avoir reporté ses limites de confiance, ses anti-motifs et ses lacunes ouvertes.
Le premier passage se concentre sur le pipeline des tensor-cores et de la mémoire SM120 :
HMMA, QMMA, OMMALDSM, STSMLDGSTS, LDGDEPBAR, DEPBARLDG, STG, LDS, STS, REDGBRA, EXIT, BSSY, , Le projet ne prétend pas que l'ISA est encore complète. Le glossaire public suit ce qui est observé et expliqué ; des pages plus approfondies sous knowledge/encoding/ suivent les familles avec suffisamment de preuves pour une documentation de type "matcher".
SASS King ne concurrence pas les désassembleurs SASS au niveau des bits. Le projet utilise les dumps locaux comme preuve principale et peut utiliser redplait/denvdis comme une vérification croisée pour les champs d'instructions, les tables d'ordonnancement, les prédicats et le suivi des registres. denvdis peut valider les interprétations d'encodage de bas niveau ; SASS King possède les preuves de variation contrôlée, la couche de motifs sémantiques et l'interprétation des audits de production.
flowchart LR
P1["Phase 1<br/>Kernels pédagogiques<br/>01-12"] --> P2["Phase 2<br/>Corpus tensor-core SM120<br/>13-25"]
P2 --> P25["Phase 2.5<br/>Validation croisée denvdis<br/>Backend niveau bit"]
P25 --> P3["Phase 3<br/>Bibliothèque de motifs<br/>Signatures du compilateur"]
P3 --> P4["Phase 4<br/>Audits de production<br/>Kernels réels"]
P4 --> P5["Phase 5<br/>Outil d'audit<br/>Rapports cubin"]
P5 --> P6["Phase 6<br/>Relecture multi-architecture<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;
Les kernels 01-12 établissent les concepts de base du SASS : fusion FMA, comportement du scoreboard, abaissement de boucle, mémoire partagée, mémoire globale, primitives warp, mathématiques à chemin lent et débordements de mémoire locale.
Les kernels 13-25 couvrent le chemin actuel des tensor-cores SM120 :
Valider redplait/denvdis comme backend de vérification croisée au niveau bit pour SM120 / SM120a avant que les audits de production ne dépendent de la bibliothèque de motifs. Le passage exécute nvd -O, nvd -S, nvd -p, et si utile nvd -T sur des cubins ou dumps locaux représentatifs couvrant HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR et les marqueurs de divergence.
Le résultat est knowledge/DENVDIS_INTEGRATION.md : un tableau de compatibilité factuel allant de la famille au statut de reconnaissance denvdis, à la couverture des modificateurs, aux champs de code de contrôle exposés et à l'action SASS King. La sortie denvdis est une preuve complémentaire, pas un remplacement des observations de dumps locaux.
Formalisation des structures récurrentes en signatures réutilisables :
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMA chaînésSTSM -> BAR -> LDS -> STGLa bibliothèque initiale de la Phase 3 contient 29 pages de motifs sous patterns/. knowledge/FINDINGS.md reste le journal de recherche et la source de vérité ; patterns/ est le point d'entrée orienté audit.
La Phase 3 est terminée au niveau de la bibliothèque initiale. Les éléments restants tels que le décodage de la disposition d'exécution, le placement complet des bits de code de contrôle et la relecture multi-architecture sont suivis comme des lacunes ou des phases futures, et non comme des blocages pour commencer les audits de production de la Phase 4.
Appliquer la bibliothèque de motifs à des kernels réels provenant de bibliothèques telles que FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad et projets connexes. L'objectif est une couverture représentative par motif algorithmique, et non un fichier markdown par kernel.
Le premier livrable de la Phase 4 devrait être un rapport d'audit manuel qui :
PATTERN-NN correspondantes ;Construire un pipeline qui prend un cubin, détecte les motifs connus et émet un rapport orienté optimisation.
Rejouer la méthodologie sur des cibles supplémentaires :
.
├── corpus/ # Kernels contrôlés, dumps et comptes rendus de chapitres
│ ├── basics/ # Kernels 01-08 : bases scalaires/vectorielles et mémoire
│ ├── warp_collectives/ # Kernels 09-10 : shuffle, vote, réduction
│ ├── math_and_spills/ # Kernels 11-12 : chemins lents et débordements
│ └── tensor_cores/ # Kernels 13-25 : études des tensor-cores
├── knowledge/ # Découvertes, glossaire, notes d'encodage
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # Bibliothèque formelle de motifs Phase 3
├── production/ # Audits de kernels de production Phase 4
├── docs/ # Notes d'intégration, structure et version
└── guide/ # Sous-module du guide de lecture SASS externe
Chaque dossier de chapitre contient les kernels sources, les artefacts compilés quand ils sont pertinents, les dumps SASS lorsqu'ils font partie de l'ensemble de preuves validé, et un compte rendu conclusion<N>.md.
Pour une explication plus complète de ce que chaque répertoire contient, lisez Structure du projet.
cuobjdump --dump-sass pour le désassemblage brut.gpuasm.com pour les scoreboards, stalls, pression et flèches de dépendance.%clock pour les sondes de latence des instructions.nvcc -Xptxas -v pour les métadonnées de registres et de débordements.SASS King opère au niveau de la couche de motifs algorithmiques : reconnaître comment les kernels compilés sont structurés et relier ces structures aux décisions d'optimisation au niveau source.
Les contributions sont les bienvenues, en particulier :
Voir CONTRIBUTING.md pour les métadonnées attendues et le standard d'écriture.
Florian Mattana. florianmattana.com
| Si vous voulez... | Commencez ici | Puis lisez |
|---|
| Comprendre le projet en 10 minutes | docs/README.md | docs/START_HERE.md, puis docs/PROJECT_STRUCTURE.md |
| Reproduire les preuves | corpus/README.md | un chapitre conclusion*.md, puis son dump .sass |
| Trouver la source de vérité | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md |
| Reconnaître un motif dans un nouveau dump | patterns/README.md | la page patterns/NN-*.md correspondante |
| Démarrer un audit de production | production/README.md | les pages PATTERN-NN correspondantes et les preuves sources |
| Contribuer une correction ou un dump | CONTRIBUTING.md | docs/START_HERE.md |
| Domaine | Statut | Où |
|---|
| Kernels pédagogiques SM120 | Terminés, kernels 01-12 | corpus/basics/01_vector_add/ à corpus/math_and_spills/12_register_spill/ |
| Études sur les tensor-cores | Terminées, jusqu'au Kernel 25 | corpus/tensor_cores/ |
| Découvertes globales | Source de vérité active | knowledge/FINDINGS.md |
| Glossaire des instructions SM120 | Actif, basé sur des preuves | knowledge/SASS_INSTRUCTIONS_SM120.md |
| Pilotes d'encodage | Commencé avec LDSM, STSM, QMMA | knowledge/encoding/ |
| Validation croisée denvdis | Premier passage terminé ; des lacunes dans les codes de contrôle subsistent | knowledge/DENVDIS_INTEGRATION.md |
| Bibliothèque de motifs | Bibliothèque initiale de la Phase 3 terminée | patterns/ |
| Audits de production | Phase suivante | production/ |
| Famille de motifs | Exemples | Où |
|---|
| Calcul des tensor-cores | Chaînes d'accumulateurs HMMA, QMMA, OMMA ; métadonnées éparses ; fragments étroits | patterns/02-* à patterns/04-*, patterns/10-*, patterns/21-* |
| Mémoire matricielle et épilogues | LDSM, STSM, pipelines de copie asynchrone, épilogues de réduction REDG | patterns/05-*, patterns/06-*, patterns/07-*, patterns/28-* |
| Flux de contrôle | divergence/reconvergence, back-edges de boucle, sorties prédicatées, pièges froids, CALLs locaux | patterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-* |
| Mémoire et registres | mémoire globale vectorisée, débordements, mise en mémoire tampon de la mémoire partagée, descripteurs, flux de registres uniformes | patterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-* |
| Arithmétique et ordonnancement | Fusion FFMA, constantes, chemins lents MUFU, scoreboards, recyclage de durée de vie | patterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-* |
| Collectifs warp | réductions warp, shuffle/vote/match/sync primitives | patterns/01-*, patterns/25-* |
| Étiquette | Signification |
|---|
[OBS] | Observé directement dans un dump, un journal, une sortie d'exécution ou un profil. |
[INF] | Inféré à partir de preuves observées. |
[HYP] | Plausible mais non confirmé. |
[RES] | Une hypothèse antérieure résolue par des preuves ultérieures. |
[GAP] | Question ouverte documentée explicitement. |
BSYNCWARPSYNCSHFL, VOTE, REDUXS2UR, R2UR, UMOV, ULEA, LDCU| Phase | Statut | Résultat | Pourquoi c'est important |
|---|
| 1. Kernels pédagogiques | Terminée | corpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/ | Établit le vocabulaire de lecture à partir d'expériences contrôlées CUDA vers SASS. |
| 2. Corpus tensor-core SM120 | Terminée | corpus/tensor_cores/13_hmma_fp16/ à 25_stsm_epilogue/ | Capture le premier ensemble de preuves SM120 / SM120a sur les tensor-cores, la mémoire matricielle, le flux de contrôle et les épilogues. |
| 2.5. Validation croisée denvdis | Premier passage terminé | knowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.md | Utilise denvdis comme vérification croisée au niveau bit sans remplacer les preuves de dumps locaux. Le placement complet des bits de stall/yield reste ouvert. |
| 3. Bibliothèque de motifs | Bibliothèque initiale terminée | patterns/ | Transforme les structures répétées du compilateur/SASS en signatures réutilisables. |
| 4. Audits de production | Prochaine | production/ | Teste si les motifs du corpus expliquent les kernels réels des bibliothèques de production. |
| 5. Outil d'audit | Planifié | pipeline cubin vers rapport | Rend la couche de motifs scriptable et reproductible. |
| 6. Relecture multi-architecture | Planifié | Comparaisons SM80, SM86, SM89, SM90a, SM100a, SM120 | Sépare les faits spécifiques à l'architecture du comportement général du SASS NVIDIA. |
| Kernel | Sujet |
|---|
| 13 | Base HMMA, allocation de registres, chaînage d'accumulateurs |
| 14 | Base QMMA FP8 / FP6 / FP4 |
| 15 | Variantes MMA étroites |
| 16 | Pic FP4 et OMMA/QMMA à échelle de bloc |
| 17 | Comportement LDSM et chargement de matrice |
| 18 | Tuile MMA pipeline et mise en tampon de copie asynchrone |
| 19 | Métadonnées MMA éparses |
| 20 | Flux de contrôle et détection de back-edge |
| 21 | Divergence et reconvergence |
| 22 | Comportement STSM de stockage de matrice |
| 23 | Sondes de disposition de fragments FP4 / FP6 |
| 24 | Audit de mini-GEMM de production |
| 25 | Disposition d'épilogue STSM et sémantiques de storeback |
| Arch | GPU représentatif | Pourquoi |
|---|
| SM80 | A100 | Base Ampere datacenter |
| SM86 | RTX 3090 | Corpus Ampere grand public |
| SM89 | RTX 4090 | Carte d'inférence grand public courante |
| SM90a | H100 | TMA, WGMMA, spécialisation warp, clusters |
| SM100a | B200 | tcgen05.mma, TMEM |
| SM120 | RTX 5070 Ti / 5090 | Point de départ Blackwell grand public |