Skip to content
KitploitKITPLOIT
StrumentiBlog
Invia
StrumentiBlog
Invia

Strumenti di Hacking, PenTest e Cybersecurity per il tuo Arsenale di Sicurezza!

Kitploit è una directory di strumenti di hacking, cybersecurity e pentesting. Scopri gli ultimi aggiornamenti dei progetti per trovare vulnerabilità, analizzare sistemi, automatizzare i test e rafforzare la tua sicurezza.

··Feed·Contatto·Privacy·© 2026 Kitploit

Directory degli strumenti

Categorie

Vedi tutte le categorie
Loading categories
sass-king — Reverse engineering del dizionario delle istruzioni NVIDIA SASS, audit dei kernel e riconoscimento di pattern attraverso le architetture GPU. | Kitploit
Strumenti/GitHubGitHub/florianmattana/sass-king
Sicurezza Sistemi EmbeddedAnalisi StaticaAnalisi del CodiceReverse EngineeringSicurezza HardwareSicurezza Hardware e IoTAnalisi di BinariPaper e RicercaApprendimento e FormazioneRisorse CurateAnalisi del Firmware
317153 mesi faRevisionato da Kitploit

Più Popolari

Vedi tutti →

Scopri gli strumenti più utilizzati dalla nostra community.

Esplora tutti gli strumenti

Sfoglia la nostra collezione di strumenti

Vedi tutti gli strumenti →
Condividi
GitHub
florianmattana/sass-king

sass-king

Reverse engineering del dizionario delle istruzioni NVIDIA SASS, audit dei kernel e riconoscimento di pattern attraverso le architetture GPU.

Vedi RepositorySito web

Logo di SASS King

SASS King

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

Architecture Status License

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.

Navigazione rapida

Il repository è organizzato come una pipeline di evidenze:

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

Perché esiste

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.

Stato attuale

Risultato della Fase 3

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é:

  • le strutture ripetute trovate nei capitoli 01-25 sono state promosse in 29 pagine di pattern con nome;
  • ogni pattern ha una spiegazione in inglese semplice, firma SASS, varianti, anti-pattern, lacune aperte e livello di confidenza;
  • i tag delle affermazioni rimangono ancorati alle evidenze sorgente in knowledge/FINDINGS.md;
  • la navigazione orientata all'audit ora parte da patterns/README.md;
  • gli elementi irrisolti vengono esplicitamente portati avanti come lacune anziché essere nascosti all'interno del testo del pattern.

Ogni pagina di pattern include:

  • significato in inglese semplice;
  • firma SASS;
  • varianti osservate;
  • limiti di interpretazione;
  • anti-pattern;
  • lacune aperte;
  • livello di confidenza.

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.

Inizia qui

  • Nuovo al progetto: leggi Inizia qui.
  • Vuoi la mappa a livello di progetto: leggi l'indice della base di conoscenza.
  • Vuoi la mappa corrente delle istruzioni: leggi Istruzioni SASS su SM120 / SM120a.
  • Vuoi note di codifica: leggi note di codifica.
  • Vuoi la fonte di verità grezza: leggi risultati.
  • Vuoi firme di audit riutilizzabili: leggi la libreria di pattern.
  • Vuoi evidenze sui tensor core: inizia con capitoli tensor core.
  • Vuoi contribuire con dump o correzioni: leggi contribuire.
  • Vuoi il limite v0.1: leggi note di rilascio.

Articoli pubblici:

  • Parte 1 - Leggere NVIDIA SASS dai Primi Principi
  • Parte 2 - Leggere la Mente del Compilatore

Metodologia

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.

Cosa è coperto

Il primo passaggio si concentra sul tensor core SM120 e sulla pipeline di memoria:

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

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;

Fase 1 - Kernel Didattici

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.

Fase 2 - Copertura Tensor Core e SM120

I kernel 13-25 coprono il percorso tensor core corrente di SM120:

Fase 2.5 - Convalida Incrociata denvdis

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.

Fase 3 - Libreria di Pattern

Strutture ricorrenti formalizzate in firme riutilizzabili:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • HMMA / QMMA / OMMA concatenati
  • STSM -> BAR -> LDS -> STG
  • riduzioni warp e collettivi cross-lane
  • firme di spill di registro
  • pattern di flusso di controllo scalari e uniformi

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

Fase 4 - Audit di Produzione

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:

  • segmenta un kernel reale in regioni SASS;
  • cita pagine PATTERN-NN corrispondenti;
  • assegna livelli di confidenza a ogni conclusione;
  • registra le regioni non spiegate come nuove lacune;
  • evita di costruire uno strumento di audit finché almeno un rapporto manuale non è stabile.

Fase 5 - Strumento di Audit

Costruire una pipeline che prende un cubin, rileva pattern noti ed emette un report orientato all'ottimizzazione.

Fase 6 - Cross-Architettura

Riprodurre la metodologia su target aggiuntivi:

Mappa del repository

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

Strumenti

  • cuobjdump --dump-sass per il disassemblaggio grezzo.
  • gpuasm.com per scoreboard, stall, pressione e frecce di dipendenza.
  • Nsight Compute per profilazione e attribuzione degli stall.
  • Microbenchmark %clock per sonde di latenza delle istruzioni.
  • nvcc -Xptxas -v per metadati di registri e spill.

Lavori Correlati

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" per la disciplina empirica di microbenchmarking dietro la validazione di latenza, throughput e dipendenza.
  • kuterdinel.com/nv_isa per il lavoro di codifica ISA NVIDIA fuzzato, specialmente l'idea di derivare regole di codifica leggibili da macchina dal comportamento del disassemblatore.
  • redplait/denvdis per tabelle opcode, disassemblaggio a livello di bit, ispezione dei campi di codifica, analisi dello scheduling, tracciamento dei registri e manipolazione dei cubin. Le tabelle data12 SM120 estratte sono usate come verifica incrociata a basso livello mentre i dump locali rimangono evidenza primaria.
  • Strumenti e note Redplait: ced cubin editor, binding Perl disassemblaggio SASS, analisi latenza SASS e note libcuda/nvasm_internal.
  • Huerta et al. 2025 per reverse engineering dello scheduling guidato dal compilatore, codici di controllo, contatori di dipendenza, flag di riutilizzo e comportamento yield.
  • Yan et al. 2026 per l'analisi del livello driver di lancio e pushbuffer sotto SASS.
  • MaxAS e TuringAS come precedenti sforzi pubblici di assemblatore SASS per architetture NVIDIA più vecchie.
  • Documentazione NVIDIA CUDA Binary Utilities per strumenti ufficiali cubin, fatbin e disassemblaggio.

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.

Contribuire

I contributi sono benvenuti, specialmente:

  • dump SASS grezzi da hardware non direttamente disponibile qui;
  • studi su kernel controllati che isolano una decisione del compilatore;
  • correzioni a osservazioni esistenti;
  • proposte di nuovi pattern per kernel di produzione;
  • confronti cross-architettura.

Vedi CONTRIBUTING.md per i metadati attesi e lo standard di scrittura.

Autore

Florian Mattana. florianmattana.com

Scarica lo strumento
Se vuoi...Inizia quiPoi leggi
Capire il progetto in 10 minutidocs/README.mddocs/START_HERE.md, poi docs/PROJECT_STRUCTURE.md
Riprodurre le evidenzecorpus/README.mdun capitolo conclusion*.md, poi il suo dump .sass
Trovare la fonte di veritàknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md
Riconoscere un pattern in un nuovo dumppatterns/README.mdla pagina patterns/NN-*.md corrispondente
Iniziare un audit di produzioneproduction/README.mdpagine PATTERN-NN corrispondenti e prove sorgente
Contribuire con una correzione o un dumpCONTRIBUTING.mddocs/START_HERE.md
AreaStatoDove
Kernel didattici SM120Completati attraverso i kernel 01-12corpus/basics/01_vector_add/ a corpus/math_and_spills/12_register_spill/
Studi sui tensor coreCompletati attraverso il Kernel 25corpus/tensor_cores/
Risultati globaliFonte di verità attivaknowledge/FINDINGS.md
Glossario istruzioni SM120Attivo, basato su evidenzeknowledge/SASS_INSTRUCTIONS_SM120.md
Piloti di codificaAvviati con LDSM, STSM, QMMAknowledge/encoding/
Convalida incrociata denvdisPrimo passaggio completato; rimangono lacune nei codici di controllo più profondiknowledge/DENVDIS_INTEGRATION.md
Libreria di patternLibreria iniziale Fase 3 completatapatterns/
Audit di produzioneProssima faseproduction/
Famiglia di patternEsempiDove
Calcolo tensor coreCatene di accumulatori HMMA, QMMA, OMMA; metadati sparsi; frammenti strettipatterns/02-* a patterns/04-*, patterns/10-*, patterns/21-*
Memoria matriciale ed epiloghiLDSM, STSM, pipeline di copia asincrona, epiloghi di riduzione REDGpatterns/05-*, patterns/06-*, patterns/07-*, patterns/28-*
Flusso di controllodivergenza/riconvergenza, back-edge di loop, uscite predicate, cold trap, CALL localipatterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-*
Memoria e registrimemoria globale vettorizzata, spill, staging in memoria condivisa, descrittori, flusso di registri uniformipatterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-*
Aritmetica e schedulingfusione FFMA, costanti, percorsi lenti MUFU, scoreboard, riciclo lifetimepatterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-*
Collettivi warpriduzioni warp, shuffle/vote/match/sync primitivepatterns/01-*, patterns/25-*
TagSignificato
[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.
BSYNC
WARPSYNC
  • SHFL, VOTE, REDUX
  • flusso di registri uniformi: S2UR, R2UR, UMOV, ULEA, LDCU
  • FaseStatoOutputPerché è importante
    1. Kernel didatticiCompletatocorpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/Stabilisce il vocabolario di lettura da esperimenti CUDA-a-SASS controllati.
    2. Corpus tensor core SM120Completatocorpus/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 denvdisPrimo passaggio completatoknowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.mdUsa 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 patternLibreria iniziale completatapatterns/Trasforma strutture SASS/compilatore ripetute in firme riutilizzabili.
    4. Audit di produzioneProssimoproduction/Verifica se i pattern del corpus spiegano kernel reali da librerie di produzione.
    5. Strumento di auditPianificatopipeline da cubin a reportRende il livello di pattern scriptabile e ripetibile.
    6. Riproduzione cross-architetturaPianificatoConfronti SM80, SM86, SM89, SM90a, SM100a, SM120Separa i fatti specifici dell'architettura dal comportamento generale di NVIDIA SASS.
    KernelArgomento
    13Baseline HMMA, allocazione registri, concatenamento accumulatori
    14Baseline QMMA FP8 / FP6 / FP4
    15Varianti MMA strette
    16Picco FP4 e OMMA/QMMA con blocco scalato
    17LDSM e comportamento del load matriciale
    18Tile MMA pipelined e staging di copia asincrona
    19Metadati MMA sparsi
    20Flusso di controllo e rilevamento back-edge
    21Divergenza e riconvergenza
    22Comportamento dello store matriciale STSM
    23Sonde di layout frammento FP4 / FP6
    24Audit mini-GEMM di produzione
    25Layout epilogo STSM e semantica storeback
    ArchGPU rappresentativaPerché
    SM80A100Baseline datacenter Ampere
    SM86RTX 3090Corpus Ampere consumer
    SM89RTX 4090Scheda comune per inferenza consumer
    SM90aH100TMA, WGMMA, specializzazione warp, cluster
    SM100aB200tcgen05.mma, TMEM
    SM120RTX 5070 Ti / 5090Punto di partenza Blackwell consumer