Skip to content
KitploitKITPLOIT
FerramentasBlog
Enviar
FerramentasBlog
Enviar

Ferramentas de Hacking, PenTest e Cibersegurança para o seu Arsenal de Segurança!

Kitploit é um diretório de ferramentas de hacking, cibersegurança e pentesting. Descubra as últimas atualizações de projetos para encontrar vulnerabilidades, analisar sistemas, automatizar testes e fortalecer sua segurança.

··Feeds·Contato·Privacidade·© 2026 Kitploit

Diretório de Ferramentas

Categorias

Ver todas as categorias
Loading categories
sass-king — 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. | Kitploit
Ferramentas/GitHubGitHub/florianmattana/sass-king
Segurança de Sistemas EmbarcadosAnálise EstáticaAnálise de CódigoEngenharia ReversaSegurança de HardwareSegurança de Hardware e IoTAnálise de BináriosPapers e PesquisaAprendizado e EducaçãoRecursos CuradosAnálise de Firmware
31715há 3 mesesRevisado pelo Kitploit

Mais Populares

Ver todos →

Descubra as ferramentas mais usadas pela nossa comunidade.

Explore todas as ferramentas

Navegue pela nossa coleção de ferramentas

Ver todas as ferramentas →
Compartilhar
GitHub
florianmattana/sass-king

sass-king

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.

Ver RepositórioSite

Logotipo SASS King

SASS King

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

Arquitetura Status Licença

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.

Navegação Rápida

O repositório está organizado como um pipeline de evidências:

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

Por que existe

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.

Estado Atual

Entregável da Fase 3

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:

  • as estruturas repetidas encontradas nos capítulos 01-25 foram promovidas para 29 páginas nomeadas de padrões;
  • cada padrão possui uma explicação em inglês claro, assinatura SASS, variantes, antipadrões, lacunas abertas e nível de confiança;
  • as tags de afirmação permanecem limitadas às evidências de origem em knowledge/FINDINGS.md;
  • a navegação voltada para auditoria agora começa em patterns/README.md;
  • itens não resolvidos são explicitamente levados adiante como lacunas, em vez de serem ocultados dentro do texto do padrão.

Cada página de padrão inclui:

  • significado em inglês claro;
  • assinatura SASS;
  • variantes observadas;
  • limites de interpretação;
  • antipadrões;
  • lacunas abertas;
  • nível de confiança.

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.

Comece por aqui

  • Novo no projeto: leia Comece por aqui.
  • Quer o mapa geral do projeto: leia o índice da base de conhecimento.
  • Quer o mapa atual de instruções: leia Instruções SASS no SM120 / SM120a.
  • Quer notas de codificação: leia notas de codificação.
  • Quer a fonte bruta da verdade: leia descobertas.
  • Quer assinaturas de auditoria reutilizáveis: leia a biblioteca de padrões.
  • Quer evidências de tensor-core: comece com capítulos de tensor-core.
  • Quer contribuir com dumps ou correções: leia contribuindo.
  • Quer os limites da v0.1: leia notas de lançamento.

Artigos públicos:

  • Parte 1 - Lendo SASS da NVIDIA a partir dos Primeiros Princípios
  • Parte 2 - Lendo a Mente do Compilador

Metodologia

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.

O que está coberto

A primeira passagem foca no pipeline de tensor-core e memória SM120:

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

Roteiro

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 - Kernels de Ensino

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.

Fase 2 - Cobertura de Tensor-Core e SM120

Kernels 13-25 cobrem o caminho atual de tensor-core SM120:

Fase 2.5 - Validação Cruzada denvdis

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.

Fase 3 - Biblioteca de Padrões

Estruturas recorrentes formalizadas em assinaturas reutilizáveis:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • encadeamento de HMMA / QMMA / OMMA
  • STSM -> BAR -> LDS -> STG
  • reduções warp e coletivos entre lanes
  • assinaturas de spill de registradores
  • padrões de fluxo de controle escalar e uniforme

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

Fase 4 - Auditorias de Produção

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:

  • segmenta um kernel real em regiões SASS;
  • cita páginas correspondentes de PATTERN-NN;
  • atribui níveis de confiança a cada conclusão;
  • registra regiões não explicadas como novas lacunas;
  • evita construir uma ferramenta de auditoria até que pelo menos um relatório manual esteja estável.

Fase 5 - Ferramenta de Auditoria

Construir um pipeline que recebe um cubin, detecta padrões conhecidos e emite um relatório orientado a otimização.

Fase 6 - Entre Arquiteturas

Reproduzir a metodologia em alvos adicionais:

Mapa do Repositório

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

Ferramentas

  • cuobjdump --dump-sass para desmontagem bruta.
  • gpuasm.com para scoreboards, stalls, pressão e setas de dependência.
  • Nsight Compute para perfilamento e atribuição de stalls.
  • %clock microbenchmarks para sondas de latência de instrução.
  • nvcc -Xptxas -v para metadados de registradores e spills.

Trabalhos Relacionados

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" pela disciplina empírica de microbenchmarking por trás da validação de latência, rendimento e dependência.
  • kuterdinel.com/nv_isa pelo trabalho de codificação de ISA da NVIDIA via fuzzing, especialmente a ideia de derivar regras de codificação legíveis por máquina a partir do comportamento do desmontador.
  • redplait/denvdis por tabelas de opcode, desmontagem em nível de bit, inspeção de campos de codificação, análise de escalonamento, rastreamento de registradores e manipulação de cubin. As tabelas data12 SM120 extraídas são usadas como verificação cruzada de baixo nível enquanto dumps locais permanecem como evidência primária.
  • Ferramentas e notas do Redplait: ced cubin editor, SASS disassembly Perl bindings, SASS latency analysis e notas sobre libcuda/nvasm_internal.
  • Huerta et al. 2025 por engenharia reversa de escalonamento guiado por compilador, códigos de controle, contadores de dependência, flags de reutilização e comportamento de yield.
  • Yan et al. 2026 pela análise de lançamento em camada de driver e pushbuffer abaixo do SASS.
  • MaxAS e TuringAS como esforços públicos anteriores de montador SASS para arquiteturas mais antigas da NVIDIA.
  • Documentação da NVIDIA CUDA Binary Utilities para ferramentas oficiais de cubin, fatbin e desmontagem.

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.

Contribuindo

Contribuições são bem-vindas, especialmente:

  • dumps SASS brutos de hardware não disponível diretamente aqui;
  • estudos controlados de kernel que isolam uma decisão do compilador;
  • correções para observações existentes;
  • novas propostas de padrões de kernel de produção;
  • comparações entre arquiteturas.

Veja CONTRIBUTING.md para os metadados esperados e padrão de escrita.

Autor

Florian Mattana. florianmattana.com

Baixar ferramenta
Se você quiser...Comece por aquiDepois leia
Entender o projeto em 10 minutosdocs/README.mddocs/START_HERE.md, depois docs/PROJECT_STRUCTURE.md
Reproduzir as evidênciascorpus/README.mdum capítulo conclusion*.md, depois seu dump .sass
Encontrar a fonte da verdadeknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md
Reconhecer um padrão em um novo dumppatterns/README.mda página correspondente patterns/NN-*.md
Iniciar uma auditoria de produçãoproduction/README.mdpáginas PATTERN-NN correspondentes e evidências de origem
Contribuir com uma correção ou dumpCONTRIBUTING.mddocs/START_HERE.md
ÁreaStatusOnde
Kernels de ensino SM120Completos até kernels 01-12corpus/basics/01_vector_add/ a corpus/math_and_spills/12_register_spill/
Estudos de tensor-coreCompletos até Kernel 25corpus/tensor_cores/
Descobertas globaisFonte ativa da verdadeknowledge/FINDINGS.md
Glossário de instruções SM120Ativo, baseado em evidênciasknowledge/SASS_INSTRUCTIONS_SM120.md
Pilotamentos de codificaçãoIniciado com LDSM, STSM, QMMAknowledge/encoding/
Validação cruzada denvdisPrimeira passada completa; lacunas mais profundas de código de controle permanecemknowledge/DENVDIS_INTEGRATION.md
Biblioteca de padrõesBiblioteca inicial da Fase 3 completapatterns/
Auditorias de produçãoPróxima faseproduction/
Família de padrõesExemplosOnde
Computação de tensor-coreCadeias de acumuladores HMMA, QMMA, OMMA; metadados esparsos; fragmentos estreitospatterns/02-* a patterns/04-*, patterns/10-*, patterns/21-*
Memória de matriz e epílogosLDSM, STSM, pipelines de cópia assíncrona, epílogos de redução REDGpatterns/05-*, patterns/06-*, patterns/07-*, patterns/28-*
Fluxo de controledivergência/reconvergência, arestas de retorno de loop, saídas predicadas, armadilhas frias, CALLs locaispatterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-*
Memória e registradoresmemória global vetorizada, spills, staging em memória compartilhada, descritores, fluxo de registradores uniformespatterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-*
Aritmética e escalonamentoFusão FFMA, constantes, caminhos lentos MUFU, scoreboards, reciclagem de tempo de vidapatterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-*
Coletivos de warpreduções warp, shuffle/vote/match/sync primitivaspatterns/01-*, patterns/25-*
TagSignificado
[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.
BSYNC
WARPSYNC
  • SHFL, VOTE, REDUX
  • fluxo de registradores uniformes: S2UR, R2UR, UMOV, ULEA, LDCU
  • FaseStatusSaídaPor que é importante
    1. Kernels de ensinoConcluídocorpus/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 SM120Concluídocorpus/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 denvdisPrimeira passada completaknowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.mdUsa 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õesBiblioteca inicial completapatterns/Transforma estruturas repetidas do compilador/SASS em assinaturas reutilizáveis.
    4. Auditorias de produçãoPróximaproduction/Testa se os padrões do corpus explicam kernels reais de bibliotecas de produção.
    5. Ferramenta de auditoriaPlanejadopipeline cubin-para-relatórioTorna a camada de padrão scriptável e repetível.
    6. Reprodução entre arquiteturasPlanejadoComparações SM80, SM86, SM89, SM90a, SM100a, SM120Separa fatos específicos de arquitetura do comportamento geral do SASS da NVIDIA.
    KernelTópico
    13Linha de base HMMA, alocação de registradores, encadeamento de acumuladores
    14Linha de base QMMA FP8 / FP6 / FP4
    15Variantes estreitas de MMA
    16Pico FP4 e OMMA/QMMA com escala de bloco
    17LDSM e comportamento de carga de matriz
    18Tile MMA pipeline e staging de cópia assíncrona
    19Metadados esparsos MMA
    20Fluxo de controle e detecção de arestas de retorno
    21Divergência e reconvergência
    22Comportamento de armazenamento de matriz STSM
    23Sondas de layout de fragmento FP4 / FP6
    24Auditoria de mini-GEMM de produção
    25Layout de epílogo STSM e semânticas de armazenamento de volta
    ArquiteturaGPU representativaPor que
    SM80A100Linha de base Ampere para datacenter
    SM86RTX 3090Corpus Ampere consumidor
    SM89RTX 4090Placa de inferência consumidora comum
    SM90aH100TMA, WGMMA, especialização warp, clusters
    SM100aB200tcgen05.mma, TMEM
    SM120RTX 5070 Ti / 5090Ponto de partida Blackwell consumidor