
Engenharia reversa do dicionário de instruções SASS da NVIDIA, auditorias de kernel e reconhecimento de padrões em todas as arquiteturas de GPU.
Engenharia reversa do SASS da NVIDIA desde kernels controlados até auditorias de produção.
Artigo 1 · Artigo 2 · Base de conhecimento · Biblioteca de padrões · Glossário de instruções SM120 · Notas de codificação · Comece por aqui · Estrutura do projeto · Capítulos de tensor-core · Contribuindo
SASS King é um projeto sistemático de engenharia reversa do SASS da NVIDIA, o conjunto de instruções nativas da GPU emitido dentro de binários CUDA compilados. O projeto começa com hardware Blackwell de consumo SM120 / SM120a e expande em direção a uma biblioteca completa de ISA e padrões entre arquiteturas ao longo do tempo.
O objetivo é prático: ajudar um engenheiro de kernel a abrir um dump SASS, reconhecer padrões do compilador, identificar estruturas relevantes para desempenho e conectar o binário de volta às decisões de otimização no nível do código-fonte.
O projeto concluiu sua biblioteca de padrões inicial da Fase 3: 29 assinaturas SASS reutilizáveis agora estão formalizadas em patterns/, com knowledge/FINDINGS.md mantida como trilha de evidências completa. O próximo grande passo é a Fase 4: aplicar esses padrões a kernels reais de produção.
O repositório está organizado como um pipeline de evidências:
corpus/ kernels controlados e evidências SASS brutas
knowledge/ descobertas do projeto, notas de instruções e notas de codificação
patterns/ assinaturas de auditoria reutilizáveis da Fase 3
production/ auditorias de kernels reais da Fase 4
O último trabalho público amplo de engenharia reversa do SASS comparável em espírito foi Jia et al. sobre Volta e Turing em 2018. Ampere, Hopper e Blackwell mudaram substancialmente o mix de instruções: caminhos de cópia assíncrona, famílias de tensor-core, instruções de load/store de matriz, formas MMA esparsas e escalonadas e novos fluxos de registradores uniformes.
SASS King preenche essa lacuna combinando micro-kernels controlados, leitura bruta de SASS, sondas de tempo de execução e auditorias de kernels de produção.
A biblioteca formal de padrões é o principal resultado da Fase 3. Ela transforma as evidências locais dos capítulos em assinaturas de auditoria reutilizáveis, para que uma auditoria possa citar um padrão nomeado em vez de reescrever toda a trilha de pesquisa a cada vez.
A Fase 3 é considerada concluída porque:
knowledge/FINDINGS.md;patterns/README.md;Cada página de padrão inclui:
Use patterns/README.md como o índice voltado para auditoria. Use knowledge/FINDINGS.md quando precisar do contexto de pesquisa mais longo por trás de um padrão.
A Fase 3 não afirma que todo comportamento do SASS da NVIDIA está decodificado. Ela estabelece uma camada de padrão SM120 / SM120a reutilizável suficientemente boa para iniciar auditorias manuais de produção. O decodificação de layout em tempo de execução, o posicionamento completo de bits de código de controle, a geração automatizada de relatórios cubin e a reprodução entre arquiteturas permanecem como trabalho futuro.
Artigos públicos:
Variação controlada. Dois kernels diferem exatamente por uma variável: dtype, ordem dos operandos, fator de desenrolamento, layout de memória ou alvo de compilação. O diff SASS isola a decisão do compilador.
Tags de afirmação rigorosas. Toda afirmação técnica usa uma tag:
Top-down e bottom-up juntos. Micro-kernels isolam instruções individuais e decisões do compilador. Kernels com aparência de produção mostram quais padrões importam em código real.
Auditorias centradas em padrões. Uma auditoria de produção deve citar uma página formal PATTERN-NN apenas após combinar a assinatura SASS visível e carregar seus limites de confiança, antipadrões e lacunas abertas.
A primeira passagem foca no pipeline de tensor-core e memória SM120:
HMMA, QMMA, OMMALDSM, STSMLDGSTS, LDGDEPBAR, DEPBARLDG, STG, LDS, STS, REDGBRA, EXIT, BSSY, , O projeto não alega que a ISA está completa ainda. O glossário público rastreia o que é observado e explicado; páginas mais profundas em knowledge/encoding/ rastreiam famílias com evidência suficiente para documentação no estilo matcher.
O SASS King não compete com desmontadores SASS de nível de bit. O projeto usa dumps locais como evidência primária e pode usar redplait/denvdis como verificação cruzada para campos de instrução, tabelas de escalonamento, predicados e rastreamento de registradores. denvdis pode validar interpretações de codificação de baixo nível; SASS King detém a evidência de variação controlada, a camada de padrão semântico e a interpretação de auditoria de produção.
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;
Kernels 01-12 estabelecem conceitos básicos de SASS: fusão FMA, comportamento de scoreboard, redução de loop, memória compartilhada, memória global, primitivas warp, matemática de caminho lento e spills de memória local.
Kernels 13-25 cobrem o caminho atual de tensor-core SM120:
Validar redplait/denvdis como o backend de verificação cruzada de nível de bit para SM120 / SM120a antes que as auditorias de produção dependam da biblioteca de padrões. A passagem executa nvd -O, nvd -S, nvd -p e, onde útil, nvd -T em cubins ou dumps locais representativos cobrindo HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR e marcadores de divergência.
O resultado é knowledge/DENVDIS_INTEGRATION.md: uma tabela factual de compatibilidade da família até o status de reconhecimento do denvdis, cobertura de modificadores, campos de código de controle expostos e a ação do SASS King. A saída do denvdis é evidência de suporte, não uma substituição para observações de dump local.
Estruturas recorrentes formalizadas em assinaturas reutilizáveis:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMASTSM -> BAR -> LDS -> STGA biblioteca inicial da Fase 3 contém 29 páginas de padrões em patterns/. knowledge/FINDINGS.md permanece como o registro de pesquisa e fonte da verdade; patterns/ é o ponto de entrada voltado para auditoria.
A Fase 3 está completa no nível da biblioteca inicial. Itens restantes como decodificação de layout em tempo de execução, posicionamento completo de bits de código de controle e reprodução entre arquiteturas são rastreados como lacunas ou fases futuras, não como bloqueadores para iniciar as auditorias de produção da Fase 4.
Aplicar a biblioteca de padrões a kernels reais de bibliotecas como FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad e projetos relacionados. O objetivo é cobertura representativa por padrão algorítmico, não um arquivo markdown por kernel.
A primeira entrega da Fase 4 deve ser um relatório de auditoria manual que:
PATTERN-NN;Construir um pipeline que recebe um cubin, detecta padrões conhecidos e emite um relatório orientado a otimização.
Reproduzir a metodologia em alvos adicionais:
.
├── corpus/ # Kernels controlados, dumps e artigos de capítulo
│ ├── basics/ # Kernels 01-08: conceitos básicos escalar/vetorial e memória
│ ├── warp_collectives/ # Kernels 09-10: shuffle, voto, redução
│ ├── math_and_spills/ # Kernels 11-12: caminhos lentos e spills
│ └── tensor_cores/ # Kernels 13-25: estudos de tensor-core
├── knowledge/ # Descobertas, glossário, notas de codificação
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # Biblioteca formal de padrões da Fase 3
├── production/ # Auditorias de kernels de produção da Fase 4
├── docs/ # Documentação de integração, estrutura e notas de lançamento
└── guide/ # Submódulo do guia externo de leitura SASS
Cada pasta de capítulo contém kernels de origem, artefatos compilados quando relevante, dumps SASS quando fazem parte do conjunto de evidências validado e um artigo conclusion<N>.md.
Para uma explicação mais completa do que pertence a cada diretório, leia Estrutura do Projeto.
cuobjdump --dump-sass para desmontagem bruta.gpuasm.com para scoreboards, stalls, pressão e setas de dependência.%clock microbenchmarks para sondas de latência de instrução.nvcc -Xptxas -v para metadados de registradores e spills.O SASS King opera na camada de padrão algorítmico: reconhecendo como kernels compilados são estruturados e conectando essas estruturas a decisões de otimização no nível do código-fonte.
Contribuições são bem-vindas, especialmente:
Veja CONTRIBUTING.md para os metadados esperados e padrão de escrita.
Florian Mattana. florianmattana.com
| Se você quiser... | Comece por aqui | Depois leia |
|---|
| Entender o projeto em 10 minutos | docs/README.md | docs/START_HERE.md, depois docs/PROJECT_STRUCTURE.md |
| Reproduzir as evidências | corpus/README.md | um capítulo conclusion*.md, depois seu dump .sass |
| Encontrar a fonte da verdade | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md |
| Reconhecer um padrão em um novo dump | patterns/README.md | a página correspondente patterns/NN-*.md |
| Iniciar uma auditoria de produção | production/README.md | páginas PATTERN-NN correspondentes e evidências de origem |
| Contribuir com uma correção ou dump | CONTRIBUTING.md | docs/START_HERE.md |
| Área | Status | Onde |
|---|
| Kernels de ensino SM120 | Completos até kernels 01-12 | corpus/basics/01_vector_add/ a corpus/math_and_spills/12_register_spill/ |
| Estudos de tensor-core | Completos até Kernel 25 | corpus/tensor_cores/ |
| Descobertas globais | Fonte ativa da verdade | knowledge/FINDINGS.md |
| Glossário de instruções SM120 | Ativo, baseado em evidências | knowledge/SASS_INSTRUCTIONS_SM120.md |
| Pilotamentos de codificação | Iniciado com LDSM, STSM, QMMA | knowledge/encoding/ |
| Validação cruzada denvdis | Primeira passada completa; lacunas mais profundas de código de controle permanecem | knowledge/DENVDIS_INTEGRATION.md |
| Biblioteca de padrões | Biblioteca inicial da Fase 3 completa | patterns/ |
| Auditorias de produção | Próxima fase | production/ |
| Família de padrões | Exemplos | Onde |
|---|
| Computação de tensor-core | Cadeias de acumuladores HMMA, QMMA, OMMA; metadados esparsos; fragmentos estreitos | patterns/02-* a patterns/04-*, patterns/10-*, patterns/21-* |
| Memória de matriz e epílogos | LDSM, STSM, pipelines de cópia assíncrona, epílogos de redução REDG | patterns/05-*, patterns/06-*, patterns/07-*, patterns/28-* |
| Fluxo de controle | divergência/reconvergência, arestas de retorno de loop, saídas predicadas, armadilhas frias, CALLs locais | patterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-* |
| Memória e registradores | memória global vetorizada, spills, staging em memória compartilhada, descritores, fluxo de registradores uniformes | patterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-* |
| Aritmética e escalonamento | Fusão FFMA, constantes, caminhos lentos MUFU, scoreboards, reciclagem de tempo de vida | patterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-* |
| Coletivos de warp | reduções warp, shuffle/vote/match/sync primitivas | patterns/01-*, patterns/25-* |
| Tag | Significado |
|---|
[OBS] | Diretamente observado em um dump, log, saída em tempo de execução ou perfil. |
[INF] | Inferido a partir de evidências observadas. |
[HYP] | Plausível mas não confirmado. |
[RES] | Uma hipótese anterior resolvida por evidências posteriores. |
[GAP] | Pergunta em aberto documentada explicitamente. |
BSYNCWARPSYNCSHFL, VOTE, REDUXS2UR, R2UR, UMOV, ULEA, LDCU| Fase | Status | Saída | Por que é importante |
|---|
| 1. Kernels de ensino | Concluído | corpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/ | Estabelece o vocabulário de leitura a partir de experimentos controlados CUDA-para-SASS. |
| 2. Corpus tensor-core SM120 | Concluído | corpus/tensor_cores/13_hmma_fp16/ a 25_stsm_epilogue/ | Captura o primeiro conjunto de evidências SM120 / SM120a de tensor-core, memória de matriz, fluxo de controle e epílogo. |
| 2.5. Validação cruzada denvdis | Primeira passada completa | knowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.md | Usa denvdis como verificação cruzada de nível de bit sem substituir a evidência de dump local. O posicionamento completo de bits de stall/yield permanece em aberto. |
| 3. Biblioteca de padrões | Biblioteca inicial completa | patterns/ | Transforma estruturas repetidas do compilador/SASS em assinaturas reutilizáveis. |
| 4. Auditorias de produção | Próxima | production/ | Testa se os padrões do corpus explicam kernels reais de bibliotecas de produção. |
| 5. Ferramenta de auditoria | Planejado | pipeline cubin-para-relatório | Torna a camada de padrão scriptável e repetível. |
| 6. Reprodução entre arquiteturas | Planejado | Comparações SM80, SM86, SM89, SM90a, SM100a, SM120 | Separa fatos específicos de arquitetura do comportamento geral do SASS da NVIDIA. |
| Kernel | Tópico |
|---|
| 13 | Linha de base HMMA, alocação de registradores, encadeamento de acumuladores |
| 14 | Linha de base QMMA FP8 / FP6 / FP4 |
| 15 | Variantes estreitas de MMA |
| 16 | Pico FP4 e OMMA/QMMA com escala de bloco |
| 17 | LDSM e comportamento de carga de matriz |
| 18 | Tile MMA pipeline e staging de cópia assíncrona |
| 19 | Metadados esparsos MMA |
| 20 | Fluxo de controle e detecção de arestas de retorno |
| 21 | Divergência e reconvergência |
| 22 | Comportamento de armazenamento de matriz STSM |
| 23 | Sondas de layout de fragmento FP4 / FP6 |
| 24 | Auditoria de mini-GEMM de produção |
| 25 | Layout de epílogo STSM e semânticas de armazenamento de volta |
| Arquitetura | GPU representativa | Por que |
|---|
| SM80 | A100 | Linha de base Ampere para datacenter |
| SM86 | RTX 3090 | Corpus Ampere consumidor |
| SM89 | RTX 4090 | Placa de inferência consumidora comum |
| SM90a | H100 | TMA, WGMMA, especialização warp, clusters |
| SM100a | B200 | tcgen05.mma, TMEM |
| SM120 | RTX 5070 Ti / 5090 | Ponto de partida Blackwell consumidor |