
从受控内核到生产审计的 NVIDIA SASS 逆向工程。
文章 1 · 文章 2 · 知识库 · 模式库 · SM120 指令词汇表 · 编码笔记 · 从这里开始 · 项目结构 · 张量核心章节 · 贡献指南
SASS King 是一个针对 NVIDIA SASS 的系统性逆向工程项目,SASS 是编译后的 CUDA 二进制文件中发出的原生 GPU 指令集。该项目从 SM120 / SM120a 消费级 Blackwell 硬件开始,并逐步向完整的跨架构 ISA 和模式库扩展。
目标是务实的:帮助内核工程师打开一个 SASS 转储,识别编译器模式,识别与性能相关的结构,并将二进制文件与源代码级别的优化决策联系起来。
该项目已完成初始的第三阶段模式库:patterns/ 下已规范化了 29 个可重用的 SASS 签名,knowledge/FINDINGS.md 作为完整的证据链。下一步是第四阶段:将这些模式应用于实际的生产内核。
| 如果你想要…… | 从这里开始 | 然后阅读 |
|---|---|---|
| 在 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 |
该仓库组织为一个证据管线:
corpus/ 受控内核和原始 SASS 证据
knowledge/ 项目范围的发现、指令注释和编码注释
patterns/ 可重用的第三阶段审计签名
production/ 第四阶段真实内核审计
上一次在精神上可比的广泛公共 SASS 逆向工程工作是 2018 年 Jia 等人关于 Volta 和 Turing 的研究。Ampere、Hopper 和 Blackwell 显著改变了指令组合:异步拷贝路径、张量核心家族、矩阵加载/存储指令、稀疏和缩放 MMA 形式,以及新的统一寄存器流。
SASS King 通过结合受控微内核、原始 SASS 阅读、运行时探测和生产内核审计来填补这一空白。
| 领域 | 状态 | 位置 |
|---|---|---|
| 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 |
| 模式库 | 初始第三阶段库已完成 | patterns/ |
| 生产审计 | 下一阶段 | production/ |
正式的模式库是第三阶段的主要输出。它将各章节的证据转化为可重用的审计签名,使得审计可以引用一个命名的模式,而不是每次都重写完整的研究轨迹。
第三阶段被认为已完成,因为:
knowledge/FINDINGS.md 中的源证据;patterns/README.md 开始;| 模式族 | 示例 | 位置 |
|---|---|---|
| 张量核心计算 | 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 慢路径、记分板、生命周期回收 | patterns/12-*,patterns/18-*,patterns/22-*,patterns/23-*,patterns/24-* |
| Warp 集合 | warp 归约、洗牌/投票/匹配/同步原语 | patterns/01-*,patterns/25-* |
每个模式页面包括:
使用 patterns/README.md 作为面向审计的索引。当需要某个模式背后的更长研究背景时,使用 knowledge/FINDINGS.md。
第三阶段并不声称每个 NVIDIA SASS 行为都已被解码。它建立了一个可重用的 SM120 / SM120a 模式层,足以开始手动生产审计。运行时布局解码、完整控制码位布局、自动化 cubin 报告和跨架构重放仍是未来工作。
公开文章:
受控变异。 两个内核恰好一个变量不同:dtype、操作数顺序、展开因子、内存布局或编译目标。SASS 差异隔离出编译器的决策。
严格声明标签。 每个技术声明都使用一个标签:
| 标签 | 含义 |
|---|---|
[OBS] | 直接在转储、日志、运行时输出或性能分析中观察到。 |
[INF] | 从观察到的证据推断得出。 |
[HYP] | 合理但未确认。 |
[RES] | 先前的假设被后来的证据解决。 |
[GAP] | 明确记录待解答的问题。 |
自上而下和自下而上相结合。 微内核隔离单个指令和编译器决策。类似生产环境的内核显示哪些模式在实际代码中重要。
模式优先的审计。 生产审计应在匹配可见的 SASS 签名并继承其置信界限、反模式和开放缺口后,才引用正式的 PATTERN-NN 页面。
第一遍专注于 SM120 张量核心和内存管线:
HMMA、QMMA、OMMALDSM、STSMLDGSTS、LDGDEPBAR、DEPBARLDG、STG、LDS、STS、REDGBRA、EXIT、BSSY、BSYNC、WARPSYNCSHFL、VOTE、REDUXS2UR、R2UR、UMOV、ULEA、LDCU该项目并不假装 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;| 阶段 | 状态 | 输出 | 为何重要 |
|---|---|---|---|
| 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 作为位级交叉检查,而不取代本地转储证据。完整的停顿/产生位放置仍待解决。 |
| 3. 模式库 | 初始库已完成 | patterns/ | 将重复的编译器/SASS 结构转化为可重用签名。 |
| 4. 生产审计 | 下一步 | production/ | 测试语料库模式是否能解释来自生产库的真实内核。 |
| 5. 审计工具 | 规划中 | cubin 到报告管线 | 使模式层可脚本化和可重复。 |
| 6. 跨架构重放 | 规划中 | SM80、SM86、SM89、SM90a、SM100a、SM120 比较 | 将特定架构的事实与通用 NVIDIA SASS 行为区分开。 |
内核 01-12 建立基线 SASS 概念:FMA 融合、记分板行为、循环降级、共享内存、全局内存、warp 原语、慢路径数学和本地内存溢出。
内核 13-25 覆盖当前的 SM120 张量核心路径:
| 内核 | 主题 |
|---|---|
| 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 尾声布局和存储回语义 |
验证 redplait/denvdis 作为 SM120 / SM120a 的位级交叉检查后端,然后再让生产审计依赖于模式库。该遍在代表性的本地 cubin 或转储上运行 nvd -O、nvd -S、nvd -p,以及在有用时运行 nvd -T,覆盖 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 库包含 patterns/ 下的 29 个模式页面。knowledge/FINDINGS.md 仍然是研究日志和事实来源;patterns/ 是面向审计的入口点。
阶段 3 在初始库级别已完成。剩余项目如运行时布局解码、完整控制码位放置和跨架构重放被作为缺口或未来阶段跟踪,而不是启动阶段 4 生产审计的障碍。
将模式库应用于来自 FlashAttention、CUTLASS、xFormers、Transformer Engine、FlashInfer、llama.cpp / ggml、tinygrad 及相关项目的真实内核。目标是通过算法模式实现代表性覆盖,而不是每个内核一个 Markdown 文件。
阶段 4 的第一个交付物应是一份手动审计报告,该报告:
PATTERN-NN 页面;构建一个管线,接收 cubin,检测已知模式,并发出一个面向优化的报告。
在更多目标上重放方法论:
| 架构 | 代表性 GPU | 原因 |
|---|---|---|
| SM80 | A100 | 数据中心 Ampere 基线 |
| SM86 | RTX 3090 | 消费级 Ampere 语料库 |
| SM89 | RTX 4090 | 常见消费级推理卡 |
| SM90a | H100 | TMA、WGMMA、warp 专业化、集群 |
| SM100a | B200 | tcgen05.mma、TMEM |
| SM120 | RTX 5070 Ti / 5090 | 消费级 Blackwell 起点 |
.
├── corpus/ # 受控内核、转储和章节文章
│ ├── basics/ # 内核 01-08:标量/向量和内存基础
│ ├── warp_collectives/ # 内核 09-10:洗牌、投票、归约
│ ├── 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 用于记分板、停顿、压力和依赖箭头。%clock 微基准测试用于指令延迟探测。nvcc -Xptxas -v 用于寄存器和溢出元数据。SASS King 在算法模式层运作:识别编译内核的结构方式,并将这些结构与源代码级别的优化决策联系起来。
欢迎贡献,尤其是:
请参阅 CONTRIBUTING.md 了解预期的元数据和写作标准。
Florian Mattana. florianmattana.com