Skip to content
KitploitKITPLOIT
OutilsBlog
Soumettre
OutilsBlog
Soumettre

Outils de Hacking, PenTest et Cybersécurité pour votre Arsenal de Sécurité !

Kitploit est un répertoire d'outils de hacking, de cybersécurité et de pentesting. Découvrez les dernières mises à jour des projets pour trouver des vulnérabilités, analyser des systèmes, automatiser les tests et renforcer votre sécurité.

··Flux·Contact·Confidentialité·© 2026 Kitploit

Répertoire d'outils

Catégories

Voir toutes les catégories
Loading categories
sass-king — Ingénierie inverse du dictionnaire d'instructions NVIDIA SASS, audits de noyaux et reconnaissance de motifs à travers les architectures GPU. | Kitploit
Outils/GitHubGitHub/florianmattana/sass-king
Sécurité des Systèmes EmbarquésAnalyse StatiqueAnalyse de CodeRétro-ingénierieSécurité MatérielleSécurité Matériel et IoTAnalyse de BinairesArticles et RechercheApprentissage et ÉducationRessources OrganiséesAnalyse de Micrologiciel
31715il y a 3 moisVérifié par Kitploit

Populaires

Voir tout →

Découvrez les outils les plus utilisés par notre communauté.

Explorer tous les outils

Parcourez notre collection d'outils

Voir tous les outils →
GitHub
florianmattana/sass-king

sass-king

Ingénierie inverse du dictionnaire d'instructions NVIDIA SASS, audits de noyaux et reconnaissance de motifs à travers les architectures GPU.

Voir le dépôtSite web
Partager

Logo SASS King

SASS King

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

Architecture Statut Licence

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.

Navigation rapide

Le dépôt est organisé comme un pipeline de preuves :

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

Pourquoi cela existe

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.

État actuel

Livrable de la Phase 3

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 :

  • les structures répétées trouvées dans les chapitres 01-25 ont été promues dans 29 pages de motifs nommés ;
  • chaque motif a une explication en anglais simple, une signature SASS, des variantes, des anti-motifs, des lacunes ouvertes et un niveau de confiance ;
  • les étiquettes de revendication restent liées aux preuves sources dans knowledge/FINDINGS.md ;
  • la navigation orientée audit commence maintenant depuis patterns/README.md ;
  • les éléments non résolus sont explicitement reportés comme des lacunes plutôt que d'être cachés dans le texte du motif.

Chaque page de motif comprend :

  • signification en anglais simple ;
  • signature SASS ;
  • variantes observées ;
  • limites d'interprétation ;
  • anti-motifs ;
  • lacunes ouvertes ;
  • niveau de confiance.

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.

Commencez ici

  • Nouveau sur le projet : lisez Commencez ici.
  • Vous voulez la carte globale du projet : lisez l'index de la base de connaissances.
  • Vous voulez la carte actuelle des instructions : lisez Instructions SASS sur SM120 / SM120a.
  • Vous voulez des notes d'encodage : lisez notes d'encodage.
  • Vous voulez la source brute de vérité : lisez découvertes.
  • Vous voulez des signatures d'audit réutilisables : lisez la bibliothèque de motifs.
  • Vous voulez des preuves sur les tensor-cores : commencez par les chapitres sur les tensor-cores.
  • Vous voulez contribuer des dumps ou des corrections : lisez contribuer.
  • Vous voulez la frontière de la v0.1 : lisez notes de version.

Articles publics :

  • Partie 1 - Lire le SASS NVIDIA à partir des premiers principes
  • Partie 2 - Lire dans les pensées du compilateur

Méthodologie

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.

Ce qui est couvert

Le premier passage se concentre sur le pipeline des tensor-cores et de la mémoire SM120 :

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

Feuille de route

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

Phase 1 - Kernels pédagogiques

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.

Phase 2 - Couverture des tensor-cores et du SM120

Les kernels 13-25 couvrent le chemin actuel des tensor-cores SM120 :

Phase 2.5 - Validation croisée denvdis

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.

Phase 3 - Bibliothèque de motifs

Formalisation des structures récurrentes en signatures réutilisables :

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • HMMA / QMMA / OMMA chaînés
  • STSM -> BAR -> LDS -> STG
  • réductions warp et collectifs inter-lanes
  • signatures de débordement de registres
  • motifs de flux de contrôle scalaires et uniformes

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

Phase 4 - Audits de production

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 :

  • segmente un kernel réel en régions SASS ;
  • cite les pages PATTERN-NN correspondantes ;
  • attribue des niveaux de confiance à chaque conclusion ;
  • enregistre les régions inexpliquées comme de nouvelles lacunes ;
  • évite de construire un outil d'audit avant qu'au moins un rapport manuel ne soit stable.

Phase 5 - Outil d'audit

Construire un pipeline qui prend un cubin, détecte les motifs connus et émet un rapport orienté optimisation.

Phase 6 - Multi-architecture

Rejouer la méthodologie sur des cibles supplémentaires :

Carte du dépôt

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

Outillage

  • cuobjdump --dump-sass pour le désassemblage brut.
  • gpuasm.com pour les scoreboards, stalls, pression et flèches de dépendance.
  • Nsight Compute pour le profilage et l'attribution des stalls.
  • Microbenchmarks %clock pour les sondes de latence des instructions.
  • nvcc -Xptxas -v pour les métadonnées de registres et de débordements.

Travaux connexes

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" pour la discipline de microbenchmarking empirique derrière la validation de latence, débit et dépendance.
  • kuterdinel.com/nv_isa pour le travail sur l'encodage de l'ISA NVIDIA par fuzzing, en particulier l'idée de dériver des règles d'encodage lisibles par machine à partir du comportement du désassembleur.
  • redplait/denvdis pour les tables d'opcodes, le désassemblage au niveau bit, l'inspection des champs d'encodage, l'analyse d'ordonnancement, le suivi des registres et la manipulation des cubins. Les tables SM120 data12 extraites sont utilisées comme vérification croisée de bas niveau tandis que les dumps locaux restent la preuve principale.
  • Outillage et notes de Redplait : ced - éditeur cubin, liaisons Perl pour désassemblage SASS, analyse de latence SASS, et notes libcuda/nvasm_internal.
  • Huerta et al. 2025 pour la rétro-ingénierie de l'ordonnancement guidé par le compilateur, les codes de contrôle, les compteurs de dépendance, les indicateurs de réutilisation et le comportement yield.
  • Yan et al. 2026 pour l'analyse de la couche pilote de lancement et de pushbuffer en dessous du SASS.
  • MaxAS et TuringAS comme précédents efforts publics d'assembleur SASS pour des architectures NVIDIA plus anciennes.
  • Documentation NVIDIA CUDA Binary Utilities pour l'outillage officiel des cubins, fatbins et désassemblage.

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.

Contribuer

Les contributions sont les bienvenues, en particulier :

  • les dumps SASS bruts provenant de matériel non directement disponible ici ;
  • les études de kernels contrôlés qui isolent une décision du compilateur ;
  • les corrections des observations existantes ;
  • les nouvelles propositions de motifs pour des kernels de production ;
  • les comparaisons multi-architecture.

Voir CONTRIBUTING.md pour les métadonnées attendues et le standard d'écriture.

Auteur

Florian Mattana. florianmattana.com

Télécharger l’outil
Si vous voulez...Commencez iciPuis lisez
Comprendre le projet en 10 minutesdocs/README.mddocs/START_HERE.md, puis docs/PROJECT_STRUCTURE.md
Reproduire les preuvescorpus/README.mdun chapitre conclusion*.md, puis son dump .sass
Trouver la source de véritéknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md
Reconnaître un motif dans un nouveau dumppatterns/README.mdla page patterns/NN-*.md correspondante
Démarrer un audit de productionproduction/README.mdles pages PATTERN-NN correspondantes et les preuves sources
Contribuer une correction ou un dumpCONTRIBUTING.mddocs/START_HERE.md
DomaineStatutOù
Kernels pédagogiques SM120Terminés, kernels 01-12corpus/basics/01_vector_add/ à corpus/math_and_spills/12_register_spill/
Études sur les tensor-coresTerminées, jusqu'au Kernel 25corpus/tensor_cores/
Découvertes globalesSource de vérité activeknowledge/FINDINGS.md
Glossaire des instructions SM120Actif, basé sur des preuvesknowledge/SASS_INSTRUCTIONS_SM120.md
Pilotes d'encodageCommencé avec LDSM, STSM, QMMAknowledge/encoding/
Validation croisée denvdisPremier passage terminé ; des lacunes dans les codes de contrôle subsistentknowledge/DENVDIS_INTEGRATION.md
Bibliothèque de motifsBibliothèque initiale de la Phase 3 terminéepatterns/
Audits de productionPhase suivanteproduction/
Famille de motifsExemplesOù
Calcul des tensor-coresChaînes d'accumulateurs HMMA, QMMA, OMMA ; métadonnées éparses ; fragments étroitspatterns/02-* à patterns/04-*, patterns/10-*, patterns/21-*
Mémoire matricielle et épiloguesLDSM, STSM, pipelines de copie asynchrone, épilogues de réduction REDGpatterns/05-*, patterns/06-*, patterns/07-*, patterns/28-*
Flux de contrôledivergence/reconvergence, back-edges de boucle, sorties prédicatées, pièges froids, CALLs locauxpatterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-*
Mémoire et registresmémoire globale vectorisée, débordements, mise en mémoire tampon de la mémoire partagée, descripteurs, flux de registres uniformespatterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-*
Arithmétique et ordonnancementFusion FFMA, constantes, chemins lents MUFU, scoreboards, recyclage de durée de viepatterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-*
Collectifs warpréductions warp, shuffle/vote/match/sync primitivespatterns/01-*, patterns/25-*
ÉtiquetteSignification
[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.
BSYNC
WARPSYNC
  • SHFL, VOTE, REDUX
  • flux de registres uniformes : S2UR, R2UR, UMOV, ULEA, LDCU
  • PhaseStatutRésultatPourquoi c'est important
    1. Kernels pédagogiquesTerminéecorpus/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 SM120Terminéecorpus/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 denvdisPremier passage terminéknowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.mdUtilise 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 motifsBibliothèque initiale terminéepatterns/Transforme les structures répétées du compilateur/SASS en signatures réutilisables.
    4. Audits de productionProchaineproduction/Teste si les motifs du corpus expliquent les kernels réels des bibliothèques de production.
    5. Outil d'auditPlanifiépipeline cubin vers rapportRend la couche de motifs scriptable et reproductible.
    6. Relecture multi-architecturePlanifiéComparaisons SM80, SM86, SM89, SM90a, SM100a, SM120Sépare les faits spécifiques à l'architecture du comportement général du SASS NVIDIA.
    KernelSujet
    13Base HMMA, allocation de registres, chaînage d'accumulateurs
    14Base QMMA FP8 / FP6 / FP4
    15Variantes MMA étroites
    16Pic FP4 et OMMA/QMMA à échelle de bloc
    17Comportement LDSM et chargement de matrice
    18Tuile MMA pipeline et mise en tampon de copie asynchrone
    19Métadonnées MMA éparses
    20Flux de contrôle et détection de back-edge
    21Divergence et reconvergence
    22Comportement STSM de stockage de matrice
    23Sondes de disposition de fragments FP4 / FP6
    24Audit de mini-GEMM de production
    25Disposition d'épilogue STSM et sémantiques de storeback
    ArchGPU représentatifPourquoi
    SM80A100Base Ampere datacenter
    SM86RTX 3090Corpus Ampere grand public
    SM89RTX 4090Carte d'inférence grand public courante
    SM90aH100TMA, WGMMA, spécialisation warp, clusters
    SM100aB200tcgen05.mma, TMEM
    SM120RTX 5070 Ti / 5090Point de départ Blackwell grand public