Skip to content
KitploitKITPLOIT
HerramientasBlog
Enviar
HerramientasBlog
Enviar

¡Herramientas de Hacking, PenTest y Ciberseguridad para tu Arsenal de Seguridad!

Kitploit es un directorio de herramientas de hacking, ciberseguridad y pentesting. Descubre las últimas actualizaciones de proyectos para encontrar vulnerabilidades, analizar sistemas, automatizar pruebas y fortalecer tu seguridad.

··Feeds·Contacto·Privacidad·© 2026 Kitploit

Directorio de Herramientas

Categorías

Ver todas las categorías
Loading categories
sass-king — Ingeniería inversa del diccionario de instrucciones NVIDIA SASS, auditorías de kernels y reconocimiento de patrones en todas las arquitecturas de GPU. | Kitploit
Herramientas/GitHubGitHub/florianmattana/sass-king
Seguridad de Sistemas EmbebidosAnálisis EstáticoAnálisis de CódigoIngeniería InversaSeguridad de HardwareSeguridad de Hardware e IoTAnálisis de BinariosPapers e InvestigaciónAprendizaje y EducaciónRecursos CuradosAnálisis de Firmware
31715hace 3 mesesRevisado por Kitploit

Más Populares

Ver todos →

Descubre las herramientas más usadas por nuestra comunidad.

Explora todas las herramientas

Explora nuestra colección de herramientas

Ver todas las herramientas →
Compartir
GitHub
florianmattana/sass-king

sass-king

Ingeniería inversa del diccionario de instrucciones NVIDIA SASS, auditorías de kernels y reconocimiento de patrones en todas las arquitecturas de GPU.

Ver RepositorioSitio web

Logo de SASS King

SASS King

Ingeniería inversa de SASS de NVIDIA desde kernels controlados hasta auditorías de producción.

Artículo 1 · Artículo 2 · Base de conocimiento · Biblioteca de patrones · Glosario de instrucciones SM120 · Notas de codificación · Empezar aquí · Estructura del proyecto · Capítulos de núcleos tensoriales · Contribuir

Arquitectura Estado Licencia

SASS King es un proyecto sistemático de ingeniería inversa para SASS de NVIDIA, el conjunto de instrucciones nativas de la GPU emitidas dentro de los binarios CUDA compilados. El proyecto comienza con hardware Blackwell de consumo SM120 / SM120a y se expande hacia una biblioteca ISA y de patrones completa entre arquitecturas con el tiempo.

El objetivo es práctico: ayudar a un ingeniero de kernels a abrir un volcado SASS, reconocer patrones del compilador, identificar estructuras relevantes para el rendimiento y conectar el binario con decisiones de optimización a nivel de código fuente.

El proyecto ha completado su biblioteca de patrones inicial de la Fase 3: 29 firmas SASS reutilizables ahora están formalizadas en patterns/, con knowledge/FINDINGS.md mantenido como la cadena de evidencia completa. El siguiente paso importante es la Fase 4: aplicar esos patrones a kernels de producción reales.

Navegación rápida

El repositorio está organizado como un pipeline de evidencia:

root@kitploit:~
corpus/      kernels controlados y evidencia SASS en bruto
knowledge/   hallazgos a nivel de proyecto, notas de instrucciones y notas de codificación
patterns/    firmas de auditoría reutilizables de la Fase 3
production/  auditorías de kernels reales de la Fase 4

Por qué existe

El último trabajo público amplio de ingeniería inversa de SASS comparable en espíritu fue Jia et al. sobre Volta y Turing en 2018. Ampere, Hopper y Blackwell han cambiado sustancialmente la mezcla de instrucciones: rutas de copia asíncrona, familias de núcleos tensoriales, instrucciones de carga/almacenamiento de matrices, formas MMA dispersas y escaladas, y nuevos flujos de registros uniformes.

SASS King llena ese vacío combinando micro-kernels controlados, lectura directa de SASS, sondas de ejecución y auditorías de kernels de producción.

Estado actual

Entregable de la Fase 3

La biblioteca formal de patrones es la principal salida de la Fase 3. Convierte la evidencia local de los capítulos en firmas de auditoría reutilizables, de modo que una auditoría pueda citar un patrón nombrado en lugar de reescribir toda la traza de investigación cada vez.

La Fase 3 se considera completa porque:

  • las estructuras repetidas encontradas en los capítulos 01-25 se han promovido a 29 páginas de patrones nombrados;
  • cada patrón tiene una explicación en lenguaje sencillo, firma SASS, variantes, antipatrones, brechas abiertas y nivel de confianza;
  • las etiquetas de afirmaciones permancen vinculadas a la evidencia fuente en knowledge/FINDINGS.md;
  • la navegación orientada a auditoría ahora comienza desde patterns/README.md;
  • los elementos no resueltos se trasladan explícitamente como brechas en lugar de ocultarse dentro del texto del patrón.

Cada página de patrón incluye:

  • significado en lenguaje sencillo;
  • firma SASS;
  • variantes observadas;
  • límites de interpretación;
  • antipatrones;
  • brechas abiertas;
  • nivel de confianza.

Usa patterns/README.md como el índice orientado a auditoría. Usa knowledge/FINDINGS.md cuando necesites el contexto de investigación más largo detrás de un patrón.

La Fase 3 no afirma que todo comportamiento SASS de NVIDIA esté decodificado. Establece una capa de patrón SM120 / SM120a reutilizable lo suficientemente buena para comenzar auditorías manuales de producción. La decodificación de diseño en tiempo de ejecución, la colocación completa de bits de código de control, el informe automatizado de cubin y la reproducción entre arquitecturas siguen siendo trabajos futuros.

Empezar aquí

  • Nuevo en el proyecto: lee Empezar aquí.
  • Quieres el mapa a nivel de proyecto: lee el índice de la base de conocimiento.
  • Quieres el mapa de instrucciones actual: lee Instrucciones SASS en SM120 / SM120a.
  • Quieres notas de codificación: lee notas de codificación.
  • Quieres la fuente de verdad en bruto: lee hallazgos.
  • Quieres firmas de auditoría reutilizables: lee la biblioteca de patrones.
  • Quieres evidencia de núcleos tensoriales: comienza con capítulos de núcleos tensoriales.
  • Quieres contribuir con volcados o correcciones: lee contribuir.
  • Quieres el límite de v0.1: lee notas de la versión.

Artículos públicos:

  • Parte 1 - Leyendo SASS de NVIDIA desde Primeros Principios
  • Parte 2 - Leyendo la Mente del Compilador

Metodología

Variación controlada. Dos kernels difieren exactamente en una variable: dtype, orden de operandos, factor de desenrollado, diseño de memoria o objetivo de compilación. El diff SASS aísla la decisión del compilador.

Etiquetas de afirmación estrictas. Cada afirmación técnica utiliza una etiqueta:

De arriba abajo y de abajo arriba juntos. Los micro-kernels aíslan instrucciones individuales y decisiones del compilador. Los kernels similares a producción muestran qué patrones importan en código real.

Auditorías primero con patrones. Una auditoría de producción debe citar una página PATTERN-NN formal solo después de hacer coincidir la firma SASS visible y arrastrar sus límites de confianza, antipatrones y brechas abiertas.

Qué está cubierto

El primer pase se centra en el pipeline de núcleos tensoriales y memoria SM120:

  • HMMA, QMMA, OMMA
  • LDSM, STSM
  • LDGSTS, LDGDEPBAR, DEPBAR
  • LDG, STG, LDS, STS, REDG
  • BRA, EXIT, BSSY, ,

El proyecto no pretende que la ISA esté completa todavía. El glosario público rastrea lo que se observa y explica; las páginas más profundas en knowledge/encoding/ rastrean familias con suficiente evidencia para documentación tipo matcher.

SASS King no compite con desensambladores SASS a nivel de bits. El proyecto utiliza volcados locales como evidencia principal y puede usar redplait/denvdis como una verificación cruzada para campos de instrucción, tablas de planificación, predicados y seguimiento de registros. denvdis puede validar interpretaciones de codificación de bajo nivel; SASS King posee la evidencia de variación controlada, la capa de patrón semántico y la interpretación de auditoría de producción.

Hoja de ruta

root@kitploit:~
flowchart LR
    P1["Fase 1<br/>Kernels didácticos<br/>01-12"] --> P2["Fase 2<br/>Corpus de núcleos tensoriales SM120<br/>13-25"]
    P2 --> P25["Fase 2.5<br/>Validación cruzada denvdis<br/>Backend a nivel de bits"]
    P25 --> P3["Fase 3<br/>Biblioteca de patrones<br/>Firmas del compilador"]
    P3 --> P4["Fase 4<br/>Auditorías de producción<br/>Kernels reales"]
    P4 --> P5["Fase 5<br/>Herramienta de auditoría<br/>Informes cubin"]
    P5 --> P6["Fase 6<br/>Reproducción entre arquitecturas<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 didácticos

Los kernels 01-12 establecen conceptos básicos de SASS: fusión FMA, comportamiento de scoreboard, reducción de bucles, memoria compartida, memoria global, primitivas warp, matemáticas de camino lento y spills de memoria local.

Fase 2 - Cobertura de núcleos tensoriales y SM120

Los kernels 13-25 cubren la ruta actual de núcleos tensoriales SM120:

Fase 2.5 - Validación cruzada denvdis

Valida redplait/denvdis como el backend de verificación cruzada a nivel de bits para SM120 / SM120a antes de que las auditorías de producción dependan de la biblioteca de patrones. El pase ejecuta nvd -O, nvd -S, nvd -p, y cuando sea útil nvd -T en cubins o volcados locales representativos que cubran HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR y marcadores de divergencia.

La salida es knowledge/DENVDIS_INTEGRATION.md: una tabla de compatibilidad factual desde la familia hasta el estado de reconocimiento de denvdis, cobertura de modificadores, campos de código de control expuestos y la acción de SASS King. La salida de denvdis es evidencia de apoyo, no un reemplazo de las observaciones de volcado local.

Fase 3 - Biblioteca de patrones

Estructuras recurrentes formalizadas en firmas reutilizables:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • HMMA / QMMA / OMMA encadenados
  • STSM -> BAR -> LDS -> STG
  • reducciones warp y colectivos entre lanes
  • firmas de spill de registros
  • patrones de flujo de control escalar y uniforme

La biblioteca inicial de la Fase 3 contiene 29 páginas de patrones en patterns/. knowledge/FINDINGS.md sigue siendo el registro de investigación y la fuente de verdad; patterns/ es el punto de entrada orientado a auditoría.

La Fase 3 está completa al nivel de la biblioteca inicial. Los elementos restantes como la decodificación de diseño en tiempo de ejecución, la colocación completa de bits de código de control y la reproducción entre arquitecturas se rastrean como brechas o fases futuras, no como bloqueos para iniciar las auditorías de producción de la Fase 4.

Fase 4 - Auditorías de producción

Aplica la biblioteca de patrones a kernels reales de bibliotecas como FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad y proyectos relacionados. El objetivo es una cobertura representativa por patrón algorítmico, no un archivo markdown por kernel.

El primer entregable de la Fase 4 debe ser un informe de auditoría manual que:

  • segmenta un kernel real en regiones SASS;
  • cita páginas PATTERN-NN coincidentes;
  • asigna niveles de confianza a cada conclusión;
  • registra regiones no explicadas como nuevas brechas;
  • evita construir una herramienta de auditoría hasta que al menos un informe manual sea estable.

Fase 5 - Herramienta de auditoría

Construye un pipeline que toma un cubin, detecta patrones conocidos y emite un informe orientado a la optimización.

Fase 6 - Entre arquitecturas

Reproduce la metodología en objetivos adicionales:

Mapa del repositorio

root@kitploit:~
.
├── corpus/                                # Kernels controlados, volcados y artículos de capítulos
│   ├── basics/                            # Kernels 01-08: conceptos básicos escalares/vectoriales y memoria
│   ├── warp_collectives/                  # Kernels 09-10: shuffle, voto, reducción
│   ├── math_and_spills/                   # Kernels 11-12: caminos lentos y spills
│   └── tensor_cores/                      # Kernels 13-25: estudios de núcleos tensoriales
├── knowledge/                              # Hallazgos, glosario, notas de codificación
│   ├── FINDINGS.md
│   ├── SASS_INSTRUCTIONS_SM120.md
│   └── encoding/
├── patterns/                               # Biblioteca formal de patrones de la Fase 3
├── production/                             # Auditorías de kernels de producción de la Fase 4
├── docs/                                   # Incorporación, estructura y notas orientadas a versiones
└── guide/                                  # Submódulo de guía de lectura externa de SASS

Cada carpeta de capítulo contiene kernels fuente, artefactos compilados cuando corresponda, volcados SASS cuando forman parte del conjunto de evidencia validado y un artículo conclusion<N>.md.

Para una explicación más completa de lo que pertenece a cada directorio, lee Estructura del proyecto.

Herramientas

  • cuobjdump --dump-sass para desensamblado en bruto.
  • gpuasm.com para scoreboards, stalls, presión y flechas de dependencia.
  • Nsight Compute para perfilado y atribución de stalls.
  • Microbenchmarks %clock para sondas de latencia de instrucciones.
  • nvcc -Xptxas -v para metadatos de registros y spills.

Trabajo relacionado

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" por la disciplina de microbenchmarking empírico detrás de la validación de latencia, rendimiento y dependencia.
  • kuterdinel.com/nv_isa por el trabajo de codificación ISA de NVIDIA basado en fuzzing, especialmente la idea de derivar reglas de codificación legibles por máquina a partir del comportamiento del desensamblador.
  • redplait/denvdis por tablas de opcodes, desensamblado a nivel de bits, inspección de campos de codificación, análisis de planificación, seguimiento de registros y manipulación de cubins. Las tablas SM120 data12 extraídas se utilizan como una verificación cruzada de bajo nivel mientras los volcados locales siguen siendo la evidencia principal.
  • Herramientas y notas de Redplait: ced cubin editor, enlaces Perl de desensamblado SASS, análisis de latencia SASS y notas internas de libcuda/nvasm_internal.
  • Huerta et al. 2025 por la ingeniería inversa de la planificación guiada por el compilador, códigos de control, contadores de dependencia, banderas de reutilización y comportamiento de yield.
  • Yan et al. 2026 por el análisis de la capa de lanzamiento de drivers y pushbuffer por debajo de SASS.
  • MaxAS y TuringAS como esfuerzos previos públicos de ensamblador SASS para arquitecturas NVIDIA más antiguas.
  • Documentación de NVIDIA CUDA Binary Utilities para herramientas oficiales de cubin, fatbin y desensamblado.

SASS King opera en la capa de patrón algorítmico: reconocer cómo se estructuran los kernels compilados y conectar esas estructuras a decisiones de optimización a nivel de código fuente.

Contribuir

Las contribuciones son bienvenidas, especialmente:

  • volcados SASS en bruto de hardware no disponible directamente aquí;
  • estudios controlados de kernels que aíslen una decisión del compilador;
  • correcciones a observaciones existentes;
  • nuevas propuestas de patrones de kernels de producción;
  • comparaciones entre arquitecturas.

Consulta CONTRIBUTING.md para conocer los metadatos esperados y el estándar de escritura.

Autor

Florian Mattana. florianmattana.com

Descargar herramienta
Si quieres...Empieza aquíLuego lee
Entender el proyecto en 10 minutosdocs/README.mddocs/START_HERE.md, luego docs/PROJECT_STRUCTURE.md
Reproducir la evidenciacorpus/README.mdun capítulo conclusion*.md, luego su volcado .sass
Encontrar la fuente de verdadknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md
Reconocer un patrón en un volcado nuevopatterns/README.mdla página patterns/NN-*.md correspondiente
Iniciar una auditoría de producciónproduction/README.mdpáginas PATTERN-NN correspondientes y evidencia fuente
Contribuir con una corrección o volcadoCONTRIBUTING.mddocs/START_HERE.md
ÁreaEstadoDónde
Kernels didácticos SM120Completos a través de kernels 01-12corpus/basics/01_vector_add/ a corpus/math_and_spills/12_register_spill/
Estudios de núcleos tensorialesCompletos hasta el Kernel 25corpus/tensor_cores/
Hallazgos globalesFuente de verdad activaknowledge/FINDINGS.md
Glosario de instrucciones SM120Activo, respaldado por evidenciaknowledge/SASS_INSTRUCTIONS_SM120.md
Pilotos de codificaciónIniciados con LDSM, STSM, QMMAknowledge/encoding/
Validación cruzada con denvdisPasada inicial completa; persisten brechas más profundas en códigos de controlknowledge/DENVDIS_INTEGRATION.md
Biblioteca de patronesBiblioteca inicial de la Fase 3 completapatterns/
Auditorías de producciónSiguiente faseproduction/
Familia de patronesEjemplosDónde
Cómputo de núcleos tensorialesCadenas de acumuladores HMMA, QMMA, OMMA; metadatos dispersos; fragmentos estrechospatterns/02-* a patterns/04-*, patterns/10-*, patterns/21-*
Memoria de matrices y epílogosLDSM, STSM, pipelines de copia asíncrona, epílogos de reducción REDGpatterns/05-*, patterns/06-*, patterns/07-*, patterns/28-*
Flujo de controldivergencia/reconvergencia, bordes de retroceso de bucle, salidas predicadas, trampas frías, CALLs localespatterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-*
Memoria y registrosmemoria global vectorizada, spills, staging en memoria compartida, descriptores, flujo de registros uniformespatterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-*
Aritmética y planificaciónFusión FFMA, constantes, caminos lentos MUFU, scoreboards, reciclaje de vida útilpatterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-*
Colectivos de warpreducciones warp, primitivas shuffle/vote/match/syncpatterns/01-*, patterns/25-*
EtiquetaSignificado
[OBS]Observado directamente en un volcado, registro, salida de ejecución o perfil.
[INF]Inferido de la evidencia observada.
[HYP]Plausible pero no confirmado.
[RES]Una hipótesis anterior resuelta por evidencia posterior.
[GAP]Pregunta abierta documentada explícitamente.
BSYNC
WARPSYNC
  • SHFL, VOTE, REDUX
  • flujo de registros uniformes: S2UR, R2UR, UMOV, ULEA, LDCU
  • FaseEstadoSalidaPor qué importa
    1. Kernels didácticosTerminadocorpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/Establece el vocabulario de lectura a partir de experimentos controlados CUDA a SASS.
    2. Corpus de núcleos tensoriales SM120Terminadocorpus/tensor_cores/13_hmma_fp16/ a 25_stsm_epilogue/Captura el primer conjunto de evidencia SM120 / SM120a de núcleos tensoriales, memoria de matrices, flujo de control y epílogo.
    2.5. Validación cruzada denvdisPasada inicial completaknowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.mdUsa denvdis como verificación cruzada a nivel de bits sin reemplazar la evidencia de volcado local. La colocación completa de bits de stall/yield sigue abierta.
    3. Biblioteca de patronesBiblioteca inicial completapatterns/Convierte estructuras repetidas del compilador/SASS en firmas reutilizables.
    4. Auditorías de producciónSiguienteproduction/Prueba si los patrones del corpus explican kernels reales de bibliotecas de producción.
    5. Herramienta de auditoríaPlanificadoPipeline cubin a informeHace que la capa de patrones sea programable y repetible.
    6. Reproducción entre arquitecturasPlanificadoComparaciones SM80, SM86, SM89, SM90a, SM100a, SM120Separa hechos específicos de arquitectura del comportamiento general de SASS de NVIDIA.
    KernelTema
    13Línea base HMMA, asignación de registros, encadenamiento de acumuladores
    14Línea base QMMA FP8 / FP6 / FP4
    15Variantes MMA estrechas
    16Pico FP4 y OMMA/QMMA con escalado por bloque
    17Comportamiento de LDSM y carga de matrices
    18Baldosa MMA en pipeline y staging de copia asíncrona
    19Metadatos de MMA dispersa
    20Flujo de control y detección de bordes de retroceso
    21Divergencia y reconvergencia
    22Comportamiento de almacenamiento de matrices STSM
    23Sondas de diseño de fragmentos FP4 / FP6
    24Auditoría mini-GEMM de producción
    25Diseño de epílogo STSM y semántica de almacenamiento de retorno
    ArquitecturaGPU representativaPor qué
    SM80A100Línea base Ampere para centros de datos
    SM86RTX 3090Corpus Ampere de consumo
    SM89RTX 4090Tarjeta de inferencia de consumo común
    SM90aH100TMA, WGMMA, especialización warp, clústeres
    SM100aB200tcgen05.mma, TMEM
    SM120RTX 5070 Ti / 5090Punto de partida Blackwell de consumo