
Обратное проектирование словаря инструкций NVIDIA SASS, аудит ядер и распознавание шаблонов для различных архитектур GPU.
Обратная разработка NVIDIA SASS: от контролируемых ядер до продукционных аудитов.
Статья 1 · Статья 2 · База знаний · Библиотека шаблонов · Глоссарий инструкций SM120 · Заметки по кодированию · Начать отсюда · Структура проекта · Главы о тензорных ядрах · Участие
SASS King — это проект систематической обратной разработки NVIDIA SASS, родного набора инструкций GPU, генерируемого внутри скомпилированных двоичных файлов CUDA. Проект начинается с потребительского оборудования Blackwell SM120 / SM120a и со временем расширяется до полной кроссплатформенной библиотеки ISA и шаблонов.
Цель — практическая: помочь инженеру по ядрам открыть дамп SASS, распознать шаблоны компилятора, выявить значимые для производительности структуры и связать двоичный код с решениями по оптимизации на уровне исходного кода.
Проект завершил начальную библиотеку шаблонов Фазы 3: 29 повторно используемых сигнатур SASS теперь формализованы в patterns/; knowledge/FINDINGS.md сохраняется как полная цепочка доказательств. Следующий крупный шаг — Фаза 4: применение этих шаблонов к реальным продукционным ядрам.
Репозиторий организован как конвейер доказательств:
corpus/ контролируемые ядра и исходные дампы SASS
knowledge/ общепроектные выводы, заметки об инструкциях и кодировании
patterns/ повторно используемые сигнатуры аудита Фазы 3
production/ аудиты реальных ядер Фазы 4
Последняя широкая публичная работа по обратной разработке SASS, сопоставимая по духу, была работа Jia et al. по Volta и Turing в 2018 году. Ampere, Hopper и Blackwell значительно изменили набор инструкций: пути асинхронного копирования, семейства тензорных ядер, инструкции матричной загрузки/сохранения, разреженные и масштабированные формы MMA, а также новые потоки униформных регистров.
SASS King заполняет этот пробел, сочетая контролируемые микроядра, чтение сырого SASS, время выполнения экспериментов и аудиты продукционных ядер.
Формальная библиотека шаблонов является основным результатом Фазы 3. Она превращает доказательства из глав в повторно используемые сигнатуры аудита, чтобы аудит мог ссылаться на именованный шаблон вместо того, чтобы каждый раз переписывать полную цепочку исследований.
Фаза 3 считается завершённой, потому что:
knowledge/FINDINGS.md;patterns/README.md;Каждая страница шаблона включает:
Используйте patterns/README.md в качестве индекса для аудита. Используйте knowledge/FINDINGS.md, когда вам нужен более длинный исследовательский контекст, стоящий за шаблоном.
Фаза 3 не утверждает, что всё поведение NVIDIA SASS декодировано. Она устанавливает повторно используемый уровень шаблонов SM120 / SM120a, достаточный для начала ручных продукционных аудитов. Декодирование макета времени выполнения, полное размещение битов управляющих кодов, автоматическое отчётность по cubin и кроссплатформенное воспроизведение остаются будущей работой.
Публичные статьи:
Контролируемое варьирование. Два ядра отличаются ровно одной переменной: тип данных, порядок операндов, коэффициент развёртки, макет памяти или цель компиляции. Разница в SASS изолирует решение компилятора.
Строгие теги утверждений. Каждое техническое утверждение использует тег:
Сверху вниз и снизу вверх вместе. Микроядра изолируют отдельные инструкции и решения компилятора. Ядра, похожие на продукционные, показывают, какие шаблоны важны в реальном коде.
Аудиты на основе шаблонов. Продукционный аудит должен ссылаться на формальную страницу PATTERN-NN только после сопоставления видимой сигнатуры SASS и переноса её пределов достоверности, анти-шаблонов и открытых пробелов.
Первый проход фокусируется на конвейере тензорных ядер и памяти SM120:
HMMA, QMMA, OMMALDSM, STSMLDGSTS, LDGDEPBAR, DEPBARLDG, STG, LDS, STS, REDGBRA, EXIT, BSSY, , Проект не претендует на то, что ISA уже завершена. Публичный глоссарий отслеживает то, что наблюдается и объяснено; более глубокие страницы в knowledge/encoding/ отслеживают семейства с достаточным количеством доказательств для документирования в стиле шаблонов.
SASS King не конкурирует с дизассемблерами SASS на уровне битов. Проект использует локальные дампы в качестве основного доказательства и может использовать redplait/denvdis в качестве перекрёстной проверки для полей инструкций, таблиц планирования, предикатов и отслеживания регистров. denvdis может проверять интерпретации кодирования низкого уровня; SASS King владеет доказательствами контролируемого варьирования, семантическим слоем шаблонов и интерпретацией продукционного аудита.
flowchart LR
P1["Фаза 1<br/>Обучающие ядра<br/>01-12"] --> P2["Фаза 2<br/>Корпус тензорных ядер SM120<br/>13-25"]
P2 --> P25["Фаза 2.5<br/>Кросс-валидация denvdis<br/>Битовый бэкенд"]
P25 --> P3["Фаза 3<br/>Библиотека шаблонов<br/>Сигнатуры компилятора"]
P3 --> P4["Фаза 4<br/>Продукционные аудиты<br/>Реальные ядра"]
P4 --> P5["Фаза 5<br/>Инструмент аудита<br/>Отчёты cubin"]
P5 --> P6["Фаза 6<br/>Кроссплатформенное воспроизведение<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;
Ядра 01–12 устанавливают базовые концепции SASS: слияние FMA, поведение scoreboard, понижение циклов, общая память, глобальная память, варп-примитивы, медленные пути математических операций и разливы локальной памяти.
Ядра 13–25 охватывают текущий путь тензорных ядер SM120:
Проверить redplait/denvdis в качестве битового бэкенда перекрёстной проверки для SM120 / SM120a перед тем, как продукционные аудиты будут полагаться на библиотеку шаблонов. Проход выполняет nvd -O, nvd -S, nvd -p и, где полезно, nvd -T на репрезентативных локальных cubin или дампах, охватывающих HMMA, QMMA, QMMA.SF, QMMA.SP, OMMA, LDSM, STSM b16/b8, LDGSTS, DEPBAR и маркеры расхождения.
Результат — knowledge/DENVDIS_INTEGRATION.md: таблица совместимости по фактам от семейства до статуса распознавания denvdis, покрытие модификаторов, раскрытые поля управляющих кодов и действие SASS King. Вывод denvdis является вспомогательным доказательством, а не заменой наблюдений из локальных дампов.
Формализованные повторяющиеся структуры в повторно используемые сигнатуры:
LDGSTS -> DEPBAR -> LDSM -> MMAHMMA / QMMA / OMMASTSM -> BAR -> LDS -> STGНачальная библиотека Фазы 3 содержит 29 страниц шаблонов в patterns/. knowledge/FINDINGS.md остаётся журналом исследований и источником истины; patterns/ — точка входа для аудита.
Фаза 3 завершена на уровне начальной библиотеки. Оставшиеся пункты, такие как декодирование макета времени выполнения, полное размещение битов управляющих кодов и кроссплатформенное воспроизведение, отслеживаются как пробелы или будущие фазы, а не как блокираторы для начала продукционных аудитов Фазы 4.
Применить библиотеку шаблонов к реальным ядрам из библиотек, таких как FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad и связанные проекты. Цель — репрезентативное покрытие по алгоритмическому шаблону, а не один файл markdown на ядро.
Первый результат Фазы 4 должен быть ручным отчётом аудита, который:
PATTERN-NN;Создать конвейер, который принимает cubin, обнаруживает известные шаблоны и выдаёт отчёт, ориентированный на оптимизацию.
Воспроизвести методологию на дополнительных целях:
.
├── corpus/ # Контролируемые ядра, дампы и главы с выводами
│ ├── basics/ # Ядра 01-08: основы скаляров/векторов и памяти
│ ├── warp_collectives/ # Ядра 09-10: shuffle, vote, редукция
│ ├── math_and_spills/ # Ядра 11-12: медленные пути и разливы
│ └── tensor_cores/ # Ядра 13-25: исследования тензорных ядер
├── knowledge/ # Выводы, глоссарий, заметки по кодированию
│ ├── FINDINGS.md
│ ├── SASS_INSTRUCTIONS_SM120.md
│ └── encoding/
├── patterns/ # Формальная библиотека шаблонов Фазы 3
├── production/ # Аудиты продукционных ядер Фазы 4
├── docs/ # Введение, структура и заметки для выпусков
└── guide/ # Внешний подмодуль руководства по чтению SASS
Каждая папка главы содержит исходные ядра, скомпилированные артефакты (если применимо), дампы SASS (если они являются частью проверенного набора доказательств) и файл conclusion<N>.md с выводами.
Для более полного объяснения того, что находится в каждом каталоге, прочитайте Структура проекта.
cuobjdump --dump-sass для сырой дизассемблеции.gpuasm.com для scoreboards, stalls, давления и стрелок зависимостей.%clock для зондов задержки инструкций.nvcc -Xptxas -v для метаданных о регистрах и разливах.SASS King работает на уровне алгоритмических шаблонов: распознаёт, как структурированы скомпилированные ядра, и связывает эти структуры с решениями по оптимизации на уровне исходного кода.
Вклад приветствуется, особенно:
См. CONTRIBUTING.md для ожидаемого формата метаданных и стандарта написания.
Флориан Маттана. florianmattana.com
| Если вы хотите... | Начать отсюда | Затем прочитайте |
|---|
| Понять проект за 10 минут | docs/README.md | docs/START_HERE.md, затем docs/PROJECT_STRUCTURE.md |
| Воспроизвести доказательства | corpus/README.md | одну главу conclusion*.md, затем её дамп .sass |
| Найти источник истины | knowledge/FINDINGS.md | knowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md |
| Распознать шаблон в новом дампе | patterns/README.md | соответствующую страницу patterns/NN-*.md |
| Начать продукционный аудит | production/README.md | соответствующие страницы PATTERN-NN и исходные доказательства |
| Внести исправление или дамп | CONTRIBUTING.md | docs/START_HERE.md |
| Область | Статус | Где |
|---|
| Обучающие ядра SM120 | Завершены ядра 01–12 | corpus/basics/01_vector_add/ — corpus/math_and_spills/12_register_spill/ |
| Исследования тензорных ядер | Завершены до ядра 25 | corpus/tensor_cores/ |
| Глобальные выводы | Активный источник истины | knowledge/FINDINGS.md |
| Глоссарий инструкций SM120 | Активный, основанный на доказательствах | knowledge/SASS_INSTRUCTIONS_SM120.md |
| Пилоты по кодированию | Начаты с LDSM, STSM, QMMA | knowledge/encoding/ |
| Кросс-валидация denvdis | Первый проход завершён; остаются пробелы в размещении управляющих кодов | knowledge/DENVDIS_INTEGRATION.md |
| Библиотека шаблонов | Начальная библиотека Фазы 3 завершена | patterns/ |
| Продукционные аудиты | Следующая фаза | production/ |
| Семейство шаблонов | Примеры | Где |
|---|
| Вычисления на тензорных ядрах | Цепочки аккумуляторов HMMA, QMMA, OMMA; разреженные метаданные; узкие фрагменты | patterns/02-* — patterns/04-*, patterns/10-*, patterns/21-* |
| Матричная память и эпилоги | LDSM, STSM, конвейеры асинхронного копирования, эпилоги редукции REDG | patterns/05-*, patterns/06-*, patterns/07-*, patterns/28-* |
| Поток управления | расхождение/схождение, обратные рёбра циклов, предикатные выходы, холодные ловушки, локальные CALL | patterns/08-*, patterns/14-*, patterns/16-*, patterns/26-*, patterns/29-* |
| Память и регистры | векторизованная глобальная память, разливы, общая память с стадиированием, дескрипторы, поток униформных регистров | patterns/09-*, patterns/11-*, patterns/17-*, patterns/19-*, patterns/20-* |
| Арифметика и планирование | слияние FFMA, константы, медленные пути MUFU, scoreboards, переработка жизненного цикла | patterns/12-*, patterns/18-*, patterns/22-*, patterns/23-*, patterns/24-* |
| Варп-коллективы | варп-редукции, shuffle/vote/match/sync примитивы | patterns/01-*, patterns/25-* |
| Тег | Значение |
|---|
[OBS] | Непосредственно наблюдалось в дампе, журнале, выводе времени выполнения или профиле. |
[INF] | Выведено из наблюдаемых доказательств. |
[HYP] | Правдоподобно, но не подтверждено. |
[RES] | Предыдущая гипотеза, разрешённая более поздними доказательствами. |
[GAP] | Открытый вопрос, задокументированный явно. |
BSYNCWARPSYNCSHFL, VOTE, REDUXS2UR, R2UR, UMOV, ULEA, LDCU| Фаза | Статус | Результат | Почему это важно |
|---|
| 1. Обучающие ядра | Завершено | corpus/basics/, corpus/warp_collectives/, corpus/math_and_spills/ | Устанавливает словарь чтения на основе контролируемых экспериментов CUDA → SASS. |
| 2. Корпус тензорных ядер SM120 | Завершено | corpus/tensor_cores/13_hmma_fp16/ — 25_stsm_epilogue/ | Захватывает первый набор доказательств тензорных ядер, матричной памяти, потока управления и эпилогов SM120 / SM120a. |
| 2.5. Кросс-валидация denvdis | Первый проход завершён | knowledge/DENVDIS_INTEGRATION.md, knowledge/encoding/CONTROL_CODE.md | Использует denvdis в качестве битовой перекрёстной проверки без замены локальных дампов. Полное размещение битов stall/yield остаётся открытым. |
| 3. Библиотека шаблонов | Начальная библиотека завершена | patterns/ | Превращает повторяющиеся структуры SASS/компилятора в повторно используемые сигнатуры. |
| 4. Продукционные аудиты | Следующая | production/ | Проверяет, объясняют ли шаблоны корпуса реальные ядра из продукционных библиотек. |
| 5. Инструмент аудита | Запланировано | конвейер cubin → отчёт | Делает слой шаблонов скриптуемым и повторяемым. |
| 6. Кроссплатформенное воспроизведение | Запланировано | Сравнение SM80, SM86, SM89, SM90a, SM100a, SM120 | Отделяет факты, специфичные для архитектуры, от общего поведения NVIDIA SASS. |
| Ядро | Тема |
|---|
| 13 | HMMA базовый, выделение регистров, цепочки аккумуляторов |
| 14 | QMMA FP8 / FP6 / FP4 базовый |
| 15 | Узкие варианты MMA |
| 16 | Пик FP4 и блочно-масштабированные OMMA/QMMA |
| 17 | LDSM и поведение матричной загрузки |
| 18 | Конвейерная плитка MMA и стадиирование асинхронного копирования |
| 19 | Разреженные метаданные MMA |
| 20 | Поток управления и обнаружение обратных рёбер |
| 21 | Расхождение и схождение |
| 22 | Поведение матричного сохранения STSM |
| 23 | Зонды макетов фрагментов FP4 / FP6 |
| 24 | Аудит продукционного мини-GEMM |
| 25 | Эпилог STSM: макет и семантика обратной записи |
| Архитектура | Репрезентативный GPU | Зачем |
|---|
| SM80 | A100 | Базовый датацентр Ampere |
| SM86 | RTX 3090 | Потребительский корпус Ampere |
| SM89 | RTX 4090 | Распространённая потребительская карта для инференса |
| SM90a | H100 | TMA, WGMMA, специализация варпов, кластеры |
| SM100a | B200 | tcgen05.mma, TMEM |
| SM120 | RTX 5070 Ti / 5090 | Потребительская отправная точка Blackwell |