Skip to content
KitploitKITPLOIT
ИнструментыБлог
Отправить
ИнструментыБлог
Отправить

Инструменты для хакинга, пентеста и кибербезопасности — ваш арсенал защиты!

Kitploit — это каталог инструментов для хакинга, кибербезопасности и пентестинга. Находите последние обновления проектов для поиска уязвимостей, анализа систем, автоматизации тестирования и усиления вашей безопасности.

··Ленты·Контакты·Конфиденциальность·© 2026 Kitploit

Каталог инструментов

Категории

Все категории
Loading categories
sass-king — Обратное проектирование словаря инструкций NVIDIA SASS, аудит ядер и распознавание шаблонов для различных архитектур GPU. | Kitploit
Инструменты/GitHubGitHub/florianmattana/sass-king
Безопасность встроенных системСтатический анализАнализ КодаОбратная инженерияАппаратная БезопасностьБезопасность оборудования и IoTАнализ Бинарных ФайловСтатьи и ИсследованияОбучение и ОбразованиеПодобранные РесурсыАнализ Прошивок
317153 месяцев назадПроверено Kitploit

Популярное

Смотреть все →

Откройте для себя самые используемые инструменты нашего сообщества.

Изучить все инструменты

Просмотрите нашу коллекцию инструментов

Смотреть все инструменты →
Поделиться
GitHub
florianmattana/sass-king

sass-king

Обратное проектирование словаря инструкций NVIDIA SASS, аудит ядер и распознавание шаблонов для различных архитектур GPU.

РепозиторийСайт

SASS King logo

SASS King

Обратная разработка NVIDIA SASS: от контролируемых ядер до продукционных аудитов.

Статья 1 · Статья 2 · База знаний · Библиотека шаблонов · Глоссарий инструкций SM120 · Заметки по кодированию · Начать отсюда · Структура проекта · Главы о тензорных ядрах · Участие

Архитектура Статус Лицензия

SASS King — это проект систематической обратной разработки NVIDIA SASS, родного набора инструкций GPU, генерируемого внутри скомпилированных двоичных файлов CUDA. Проект начинается с потребительского оборудования Blackwell SM120 / SM120a и со временем расширяется до полной кроссплатформенной библиотеки ISA и шаблонов.

Цель — практическая: помочь инженеру по ядрам открыть дамп SASS, распознать шаблоны компилятора, выявить значимые для производительности структуры и связать двоичный код с решениями по оптимизации на уровне исходного кода.

Проект завершил начальную библиотеку шаблонов Фазы 3: 29 повторно используемых сигнатур SASS теперь формализованы в patterns/; knowledge/FINDINGS.md сохраняется как полная цепочка доказательств. Следующий крупный шаг — Фаза 4: применение этих шаблонов к реальным продукционным ядрам.

Быстрая навигация

Репозиторий организован как конвейер доказательств:

root@kitploit:~
corpus/      контролируемые ядра и исходные дампы SASS
knowledge/   общепроектные выводы, заметки об инструкциях и кодировании
patterns/    повторно используемые сигнатуры аудита Фазы 3
production/  аудиты реальных ядер Фазы 4

Зачем это существует

Последняя широкая публичная работа по обратной разработке SASS, сопоставимая по духу, была работа Jia et al. по Volta и Turing в 2018 году. Ampere, Hopper и Blackwell значительно изменили набор инструкций: пути асинхронного копирования, семейства тензорных ядер, инструкции матричной загрузки/сохранения, разреженные и масштабированные формы MMA, а также новые потоки униформных регистров.

SASS King заполняет этот пробел, сочетая контролируемые микроядра, чтение сырого SASS, время выполнения экспериментов и аудиты продукционных ядер.

Текущее состояние

Результат Фазы 3

Формальная библиотека шаблонов является основным результатом Фазы 3. Она превращает доказательства из глав в повторно используемые сигнатуры аудита, чтобы аудит мог ссылаться на именованный шаблон вместо того, чтобы каждый раз переписывать полную цепочку исследований.

Фаза 3 считается завершённой, потому что:

  • повторяющиеся структуры, найденные в главах 01–25, были повышены до 29 именованных страниц шаблонов;
  • каждый шаблон содержит объяснение на простом английском, сигнатуру SASS, варианты, анти-шаблоны, открытые пробелы и уровень достоверности;
  • теги утверждений остаются привязанными к исходным доказательствам в knowledge/FINDINGS.md;
  • навигация для аудита теперь начинается с patterns/README.md;
  • нерешённые вопросы явно переносятся как пробелы, а не скрываются внутри текста шаблона.

Каждая страница шаблона включает:

  • значение на простом английском;
  • сигнатуру SASS;
  • наблюдаемые варианты;
  • границы интерпретации;
  • анти-шаблоны;
  • открытые пробелы;
  • уровень достоверности.

Используйте patterns/README.md в качестве индекса для аудита. Используйте knowledge/FINDINGS.md, когда вам нужен более длинный исследовательский контекст, стоящий за шаблоном.

Фаза 3 не утверждает, что всё поведение NVIDIA SASS декодировано. Она устанавливает повторно используемый уровень шаблонов SM120 / SM120a, достаточный для начала ручных продукционных аудитов. Декодирование макета времени выполнения, полное размещение битов управляющих кодов, автоматическое отчётность по cubin и кроссплатформенное воспроизведение остаются будущей работой.

Начать отсюда

  • Новичок в проекте: прочитайте Start Here.
  • Хотите общепроектную карту: прочитайте индекс базы знаний.
  • Хотите текущую карту инструкций: прочитайте Инструкции SASS на SM120 / SM120a.
  • Хотите заметки по кодированию: прочитайте заметки по кодированию.
  • Хотите исходный источник истины: прочитайте выводы.
  • Хотите повторно используемые сигнатуры аудита: прочитайте библиотеку шаблонов.
  • Хотите доказательства тензорных ядер: начните с глав о тензорных ядрах.
  • Хотите внести дампы или исправления: прочитайте участие.
  • Хотите границы версии v0.1: прочитайте примечания к выпуску.

Публичные статьи:

  • Часть 1 — Чтение NVIDIA SASS с первых принципов
  • Часть 2 — Чтение мыслей компилятора

Методология

Контролируемое варьирование. Два ядра отличаются ровно одной переменной: тип данных, порядок операндов, коэффициент развёртки, макет памяти или цель компиляции. Разница в SASS изолирует решение компилятора.

Строгие теги утверждений. Каждое техническое утверждение использует тег:

Сверху вниз и снизу вверх вместе. Микроядра изолируют отдельные инструкции и решения компилятора. Ядра, похожие на продукционные, показывают, какие шаблоны важны в реальном коде.

Аудиты на основе шаблонов. Продукционный аудит должен ссылаться на формальную страницу PATTERN-NN только после сопоставления видимой сигнатуры SASS и переноса её пределов достоверности, анти-шаблонов и открытых пробелов.

Что охвачено

Первый проход фокусируется на конвейере тензорных ядер и памяти SM120:

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

Проект не претендует на то, что ISA уже завершена. Публичный глоссарий отслеживает то, что наблюдается и объяснено; более глубокие страницы в knowledge/encoding/ отслеживают семейства с достаточным количеством доказательств для документирования в стиле шаблонов.

SASS King не конкурирует с дизассемблерами SASS на уровне битов. Проект использует локальные дампы в качестве основного доказательства и может использовать redplait/denvdis в качестве перекрёстной проверки для полей инструкций, таблиц планирования, предикатов и отслеживания регистров. denvdis может проверять интерпретации кодирования низкого уровня; SASS King владеет доказательствами контролируемого варьирования, семантическим слоем шаблонов и интерпретацией продукционного аудита.

Дорожная карта

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

Фаза 1 — Обучающие ядра

Ядра 01–12 устанавливают базовые концепции SASS: слияние FMA, поведение scoreboard, понижение циклов, общая память, глобальная память, варп-примитивы, медленные пути математических операций и разливы локальной памяти.

Фаза 2 — Охват тензорных ядер и SM120

Ядра 13–25 охватывают текущий путь тензорных ядер SM120:

Фаза 2.5 — Кросс-валидация denvdis

Проверить 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 является вспомогательным доказательством, а не заменой наблюдений из локальных дампов.

Фаза 3 — Библиотека шаблонов

Формализованные повторяющиеся структуры в повторно используемые сигнатуры:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • цепочки HMMA / QMMA / OMMA
  • STSM -> BAR -> LDS -> STG
  • варп-редукции и межлинейные коллективы
  • сигнатуры разливов регистров
  • скалярные и униформные шаблоны потока управления

Начальная библиотека Фазы 3 содержит 29 страниц шаблонов в patterns/. knowledge/FINDINGS.md остаётся журналом исследований и источником истины; patterns/ — точка входа для аудита.

Фаза 3 завершена на уровне начальной библиотеки. Оставшиеся пункты, такие как декодирование макета времени выполнения, полное размещение битов управляющих кодов и кроссплатформенное воспроизведение, отслеживаются как пробелы или будущие фазы, а не как блокираторы для начала продукционных аудитов Фазы 4.

Фаза 4 — Продукционные аудиты

Применить библиотеку шаблонов к реальным ядрам из библиотек, таких как FlashAttention, CUTLASS, xFormers, Transformer Engine, FlashInfer, llama.cpp / ggml, tinygrad и связанные проекты. Цель — репрезентативное покрытие по алгоритмическому шаблону, а не один файл markdown на ядро.

Первый результат Фазы 4 должен быть ручным отчётом аудита, который:

  • сегментирует одно реальное ядро на регионы SASS;
  • ссылается на соответствующие страницы PATTERN-NN;
  • присваивает уровни достоверности каждому выводу;
  • записывает необъяснённые регионы как новые пробелы;
  • избегает создания инструмента аудита до тех пор, пока хотя бы один ручной отчёт не станет стабильным.

Фаза 5 — Инструмент аудита

Создать конвейер, который принимает cubin, обнаруживает известные шаблоны и выдаёт отчёт, ориентированный на оптимизацию.

Фаза 6 — Кроссплатформенное воспроизведение

Воспроизвести методологию на дополнительных целях:

Карта репозитория

root@kitploit:~
.
├── 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, давления и стрелок зависимостей.
  • Nsight Compute для профилирования и атрибуции stalls.
  • Микробенчмарки %clock для зондов задержки инструкций.
  • nvcc -Xptxas -v для метаданных о регистрах и разливах.

Связанные работы

  • Jia et al. 2018, "Dissecting the NVIDIA Volta GPU Architecture via Microbenchmarking" — дисциплина эмпирического микробенчмаркинга, лежащая в основе проверки задержки, пропускной способности и зависимостей.
  • kuterdinel.com/nv_isa — работа по фаззингу кодирования ISA NVIDIA, особенно идея вывода машиночитаемых правил кодирования из поведения дизассемблера.
  • redplait/denvdis — таблицы opcode, битовая дизассемблеция, проверка полей кодирования, анализ планирования, отслеживание регистров и манипуляции с cubin. Извлечённые таблицы SM120 data12 используются как низкоуровневая перекрёстная проверка, в то время как локальные дампы остаются первичным доказательством.
  • Инструментарий и заметки Redplait: cubin editor ced, привязки Perl для дизассемблера SASS, анализ задержки SASS и заметки о libcuda/nvasm_internal.
  • Huerta et al. 2025 — обратная разработка планирования, управляемого компилятором, управляющих кодов, счётчиков зависимостей, флагов повторного использования и поведения yield.
  • Yan et al. 2026 — анализ уровня драйвера: запуск и pushbuffer ниже SASS.
  • MaxAS и TuringAS — предыдущие публичные усилия по созданию ассемблера SASS для старых архитектур NVIDIA.
  • Документация NVIDIA CUDA Binary Utilities — официальная документация по работе с cubin, fatbin и дизассемблированию.

SASS King работает на уровне алгоритмических шаблонов: распознаёт, как структурированы скомпилированные ядра, и связывает эти структуры с решениями по оптимизации на уровне исходного кода.

Участие

Вклад приветствуется, особенно:

  • сырые дампы SASS с оборудования, недоступного здесь напрямую;
  • контролируемые исследования ядер, изолирующие одно решение компилятора;
  • исправления существующих наблюдений;
  • новые предложения шаблонов для продукционных ядер;
  • кроссплатформенные сравнения.

См. CONTRIBUTING.md для ожидаемого формата метаданных и стандарта написания.

Автор

Флориан Маттана. florianmattana.com

Скачать инструмент
Если вы хотите...Начать отсюдаЗатем прочитайте
Понять проект за 10 минутdocs/README.mddocs/START_HERE.md, затем docs/PROJECT_STRUCTURE.md
Воспроизвести доказательстваcorpus/README.mdодну главу conclusion*.md, затем её дамп .sass
Найти источник истиныknowledge/FINDINGS.mdknowledge/SASS_INSTRUCTIONS_SM120.md, knowledge/encoding/README.md
Распознать шаблон в новом дампеpatterns/README.mdсоответствующую страницу patterns/NN-*.md
Начать продукционный аудитproduction/README.mdсоответствующие страницы PATTERN-NN и исходные доказательства
Внести исправление или дампCONTRIBUTING.mddocs/START_HERE.md
ОбластьСтатусГде
Обучающие ядра SM120Завершены ядра 01–12corpus/basics/01_vector_add/ — corpus/math_and_spills/12_register_spill/
Исследования тензорных ядерЗавершены до ядра 25corpus/tensor_cores/
Глобальные выводыАктивный источник истиныknowledge/FINDINGS.md
Глоссарий инструкций SM120Активный, основанный на доказательствахknowledge/SASS_INSTRUCTIONS_SM120.md
Пилоты по кодированиюНачаты с LDSM, STSM, QMMAknowledge/encoding/
Кросс-валидация denvdisПервый проход завершён; остаются пробелы в размещении управляющих кодовknowledge/DENVDIS_INTEGRATION.md
Библиотека шаблоновНачальная библиотека Фазы 3 завершенаpatterns/
Продукционные аудитыСледующая фазаproduction/
Семейство шаблоновПримерыГде
Вычисления на тензорных ядрахЦепочки аккумуляторов HMMA, QMMA, OMMA; разреженные метаданные; узкие фрагментыpatterns/02-* — patterns/04-*, patterns/10-*, patterns/21-*
Матричная память и эпилогиLDSM, STSM, конвейеры асинхронного копирования, эпилоги редукции REDGpatterns/05-*, patterns/06-*, patterns/07-*, patterns/28-*
Поток управлениярасхождение/схождение, обратные рёбра циклов, предикатные выходы, холодные ловушки, локальные CALLpatterns/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]Открытый вопрос, задокументированный явно.
BSYNC
WARPSYNC
  • SHFL, VOTE, REDUX
  • поток униформных регистров: S2UR, 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.
    ЯдроТема
    13HMMA базовый, выделение регистров, цепочки аккумуляторов
    14QMMA FP8 / FP6 / FP4 базовый
    15Узкие варианты MMA
    16Пик FP4 и блочно-масштабированные OMMA/QMMA
    17LDSM и поведение матричной загрузки
    18Конвейерная плитка MMA и стадиирование асинхронного копирования
    19Разреженные метаданные MMA
    20Поток управления и обнаружение обратных рёбер
    21Расхождение и схождение
    22Поведение матричного сохранения STSM
    23Зонды макетов фрагментов FP4 / FP6
    24Аудит продукционного мини-GEMM
    25Эпилог STSM: макет и семантика обратной записи
    АрхитектураРепрезентативный GPUЗачем
    SM80A100Базовый датацентр Ampere
    SM86RTX 3090Потребительский корпус Ampere
    SM89RTX 4090Распространённая потребительская карта для инференса
    SM90aH100TMA, WGMMA, специализация варпов, кластеры
    SM100aB200tcgen05.mma, TMEM
    SM120RTX 5070 Ti / 5090Потребительская отправная точка Blackwell