
Ingeniería inversa del diccionario de instrucciones NVIDIA SASS, auditorías de kernels y reconocimiento de patrones en todas las arquitecturas de GPU.
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
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.
El repositorio está organizado como un pipeline de evidencia:
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
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.
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:
knowledge/FINDINGS.md;patterns/README.md;Cada página de patrón incluye:
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.
Artículos públicos:
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.
El primer pase se centra en el pipeline de núcleos tensoriales y memoria SM120:
HMMA, QMMA, OMMALDSM, STSMLDGSTS, LDGDEPBAR, DEPBARLDG, STG, LDS, STS, REDGBRA, 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.
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;
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.
Los kernels 13-25 cubren la ruta actual de núcleos tensoriales SM120:
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.
Estructuras recurrentes formalizadas en firmas reutilizables:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMA encadenadosSTSM -> BAR -> LDS -> STGLa 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.
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:
PATTERN-NN coincidentes;Construye un pipeline que toma un cubin, detecta patrones conocidos y emite un informe orientado a la optimización.
Reproduce la metodología en objetivos adicionales:
.
├── 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.
cuobjdump --dump-sass para desensamblado en bruto.gpuasm.com para scoreboards, stalls, presión y flechas de dependencia.%clock para sondas de latencia de instrucciones.nvcc -Xptxas -v para metadatos de registros y spills.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.
Las contribuciones son bienvenidas, especialmente:
Consulta CONTRIBUTING.md para conocer los metadatos esperados y el estándar de escritura.
Florian Mattana. florianmattana.com
| Si quieres... | Empieza aquí | Luego lee |
|---|
| Entender el proyecto en 10 minutos | docs/README.md | docs/START_HERE.md, luego docs/PROJECT_STRUCTURE.md |
| Reproducir la evidencia | corpus/README.md | un capítulo conclusion*.md, luego su volcado .sass |
| Encontrar la fuente de verdad | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md |
| Reconocer un patrón en un volcado nuevo | patterns/README.md | la página patterns/NN-*.md correspondiente |
| Iniciar una auditoría de producción | production/README.md | páginas PATTERN-NN correspondientes y evidencia fuente |
| Contribuir con una corrección o volcado | CONTRIBUTING.md | docs/START_HERE.md |
| Área | Estado | Dónde |
|---|
| Kernels didácticos SM120 | Completos a través de kernels 01-12 | corpus/basics/01_vector_add/ a corpus/math_and_spills/12_register_spill/ |
| Estudios de núcleos tensoriales | Completos hasta el Kernel 25 | corpus/tensor_cores/ |
| Hallazgos globales | Fuente de verdad activa | knowledge/FINDINGS.md |
| Glosario de instrucciones SM120 | Activo, respaldado por evidencia | knowledge/SASS_INSTRUCTIONS_SM120.md |
| Pilotos de codificación | Iniciados con LDSM, STSM, QMMA | knowledge/encoding/ |
| Validación cruzada con denvdis | Pasada inicial completa; persisten brechas más profundas en códigos de control | knowledge/DENVDIS_INTEGRATION.md |
| Biblioteca de patrones | Biblioteca inicial de la Fase 3 completa | patterns/ |
| Auditorías de producción | Siguiente fase | production/ |
| Familia de patrones | Ejemplos | Dónde |
|---|
| Cómputo de núcleos tensoriales | Cadenas de acumuladores HMMA, QMMA, OMMA; metadatos dispersos; fragmentos estrechos | patterns/02-* a patterns/04-*, patterns/10-*, patterns/21-* |
| Memoria de matrices y epílogos | LDSM, STSM, pipelines de copia asíncrona, epílogos de reducción REDG | patterns/05-*, patterns/06-*, patterns/07-*, patterns/28-* |
| Flujo de control | divergencia/reconvergencia, bordes de retroceso de bucle, salidas predicadas, trampas frías, CALLs locales | patterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-* |
| Memoria y registros | memoria global vectorizada, spills, staging en memoria compartida, descriptores, flujo de registros uniformes | patterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-* |
| Aritmética y planificación | Fusión FFMA, constantes, caminos lentos MUFU, scoreboards, reciclaje de vida útil | patterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-* |
| Colectivos de warp | reducciones warp, primitivas shuffle/vote/match/sync | patterns/01-*, patterns/25-* |
| Etiqueta | Significado |
|---|
[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. |
BSYNCWARPSYNCSHFL, VOTE, REDUXS2UR, R2UR, UMOV, ULEA, LDCU| Fase | Estado | Salida | Por qué importa |
|---|
| 1. Kernels didácticos | Terminado | corpus/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 SM120 | Terminado | corpus/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 denvdis | Pasada inicial completa | knowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.md | Usa 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 patrones | Biblioteca inicial completa | patterns/ | Convierte estructuras repetidas del compilador/SASS en firmas reutilizables. |
| 4. Auditorías de producción | Siguiente | production/ | Prueba si los patrones del corpus explican kernels reales de bibliotecas de producción. |
| 5. Herramienta de auditoría | Planificado | Pipeline cubin a informe | Hace que la capa de patrones sea programable y repetible. |
| 6. Reproducción entre arquitecturas | Planificado | Comparaciones SM80, SM86, SM89, SM90a, SM100a, SM120 | Separa hechos específicos de arquitectura del comportamiento general de SASS de NVIDIA. |
| Kernel | Tema |
|---|
| 13 | Línea base HMMA, asignación de registros, encadenamiento de acumuladores |
| 14 | Línea base QMMA FP8 / FP6 / FP4 |
| 15 | Variantes MMA estrechas |
| 16 | Pico FP4 y OMMA/QMMA con escalado por bloque |
| 17 | Comportamiento de LDSM y carga de matrices |
| 18 | Baldosa MMA en pipeline y staging de copia asíncrona |
| 19 | Metadatos de MMA dispersa |
| 20 | Flujo de control y detección de bordes de retroceso |
| 21 | Divergencia y reconvergencia |
| 22 | Comportamiento de almacenamiento de matrices STSM |
| 23 | Sondas de diseño de fragmentos FP4 / FP6 |
| 24 | Auditoría mini-GEMM de producción |
| 25 | Diseño de epílogo STSM y semántica de almacenamiento de retorno |
| Arquitectura | GPU representativa | Por qué |
|---|
| SM80 | A100 | Línea base Ampere para centros de datos |
| SM86 | RTX 3090 | Corpus Ampere de consumo |
| SM89 | RTX 4090 | Tarjeta de inferencia de consumo común |
| SM90a | H100 | TMA, WGMMA, especialización warp, clústeres |
| SM100a | B200 | tcgen05.mma, TMEM |
| SM120 | RTX 5070 Ti / 5090 | Punto de partida Blackwell de consumo |