Skip to content
KitploitKITPLOIT
工具博客
提交
工具博客
提交

黑客、渗透测试和网络安全工具,武装您的安全武器库!

Kitploit 是一个黑客、网络安全和渗透测试工具的目录。发现最新的项目更新,查找漏洞、分析系统、自动化测试并加强你的安全。

··订阅源·联系·隐私·© 2026 Kitploit

工具目录

分类

查看所有分类
Loading categories
工具/GitHubGitHub/florianmattana/sass-king
嵌入式系统安全静态分析代码分析逆向工程硬件安全硬件与物联网安全二进制分析论文与研究学习与教育精选资源固件分析
GitHubflorianmattana/sass-king

sass-king

逆向工程NVIDIA SASS指令字典,跨GPU架构的内核审计与模式识别。

查看仓库网站
3171563个月前Kitploit 审核通过

最受欢迎

查看全部 →

发现我们社区最常用的工具。

探索所有工具

浏览我们的工具集合

查看所有工具 →
分享

SASS King 标志

SASS King

从受控内核到生产审计的 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.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

该仓库组织为一个证据管线:

root@kitploit:~
corpus/      受控内核和原始 SASS 证据
knowledge/   项目范围的发现、指令注释和编码注释
patterns/    可重用的第三阶段审计签名
production/  第四阶段真实内核审计

为什么存在

上一次在精神上可比的广泛公共 SASS 逆向工程工作是 2018 年 Jia 等人关于 Volta 和 Turing 的研究。Ampere、Hopper 和 Blackwell 显著改变了指令组合:异步拷贝路径、张量核心家族、矩阵加载/存储指令、稀疏和缩放 MMA 形式,以及新的统一寄存器流。

SASS King 通过结合受控微内核、原始 SASS 阅读、运行时探测和生产内核审计来填补这一空白。

当前状态

领域状态位置
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、QMMA 开始knowledge/encoding/
denvdis 交叉验证初始遍已完成;更深层次的控制码空白仍然存在knowledge/DENVDIS_INTEGRATION.md
模式库初始第三阶段库已完成patterns/
生产审计下一阶段production/

第三阶段交付物

正式的模式库是第三阶段的主要输出。它将各章节的证据转化为可重用的审计签名,使得审计可以引用一个命名的模式,而不是每次都重写完整的研究轨迹。

第三阶段被认为已完成,因为:

  • 在章节 01-25 中发现的重复结构已被提升为 29 个命名模式页面;
  • 每个模式都有通俗易懂的解释、SASS 签名、变体、反模式、开放缺口和置信级别;
  • 声明标签仍限于 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-*
控制流散开/重新汇聚、循环回边、谓词退出、冷陷阱、本地 CALLpatterns/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-*

每个模式页面包括:

  • 通俗易懂的含义;
  • SASS 签名;
  • 观察到的变体;
  • 解释边界;
  • 反模式;
  • 开放缺口;
  • 置信级别。

使用 patterns/README.md 作为面向审计的索引。当需要某个模式背后的更长研究背景时,使用 knowledge/FINDINGS.md。

第三阶段并不声称每个 NVIDIA SASS 行为都已被解码。它建立了一个可重用的 SM120 / SM120a 模式层,足以开始手动生产审计。运行时布局解码、完整控制码位布局、自动化 cubin 报告和跨架构重放仍是未来工作。

从这里开始

  • 项目新手:阅读从这里开始。
  • 想要项目范围地图:阅读知识库索引。
  • 想要当前指令映射:阅读SM120 / SM120a 上的 SASS 指令。
  • 想要编码笔记:阅读编码笔记。
  • 想要原始事实来源:阅读发现。
  • 想要可重用的审计签名:阅读模式库。
  • 想要张量核心证据:从张量核心章节开始。
  • 想要贡献转储或修正:阅读贡献指南。
  • 想要 v0.1 边界:阅读发布说明。

公开文章:

  • 第 1 部分 - 从第一原理阅读 NVIDIA SASS
  • 第 2 部分 - 读懂编译器的心思

方法论

受控变异。 两个内核恰好一个变量不同:dtype、操作数顺序、展开因子、内存布局或编译目标。SASS 差异隔离出编译器的决策。

严格声明标签。 每个技术声明都使用一个标签:

标签含义
[OBS]直接在转储、日志、运行时输出或性能分析中观察到。
[INF]从观察到的证据推断得出。
[HYP]合理但未确认。
[RES]先前的假设被后来的证据解决。
[GAP]明确记录待解答的问题。

自上而下和自下而上相结合。 微内核隔离单个指令和编译器决策。类似生产环境的内核显示哪些模式在实际代码中重要。

模式优先的审计。 生产审计应在匹配可见的 SASS 签名并继承其置信界限、反模式和开放缺口后,才引用正式的 PATTERN-NN 页面。

涵盖范围

第一遍专注于 SM120 张量核心和内存管线:

  • HMMA、QMMA、OMMA
  • LDSM、STSM
  • LDGSTS、LDGDEPBAR、DEPBAR
  • LDG、STG、LDS、STS、REDG
  • BRA、EXIT、BSSY、BSYNC、WARPSYNC
  • SHFL、VOTE、REDUX
  • 统一寄存器流:S2UR、R2UR、UMOV、ULEA、LDCU

该项目并不假装 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. 教学内核完成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 行为区分开。

阶段 1 - 教学内核

内核 01-12 建立基线 SASS 概念:FMA 融合、记分板行为、循环降级、共享内存、全局内存、warp 原语、慢路径数学和本地内存溢出。

阶段 2 - 张量核心和 SM120 覆盖

内核 13-25 覆盖当前的 SM120 张量核心路径:

内核主题
13HMMA 基线、寄存器分配、累加器链
14QMMA FP8 / FP6 / FP4 基线
15窄 MMA 变体
16FP4 峰值和块缩放 OMMA/QMMA
17LDSM 和矩阵加载行为
18流水线化 MMA 瓦片和异步拷贝暂存
19稀疏 MMA 元数据
20控制流和回边检测
21散开和重新汇聚
22STSM 矩阵存储行为
23FP4 / FP6 分片布局探测
24生产微型 GEMM 审计
25STSM 尾声布局和存储回语义

阶段 2.5 - denvdis 交叉验证

验证 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 输出是辅助证据,不是本地转储观察的替代品。

阶段 3 - 模式库

将重复出现的结构形式化为可重用签名:

  • LDGSTS -> DEPBAR -> LDSM -> MMA
  • 链式 HMMA / QMMA / OMMA
  • STSM -> BAR -> LDS -> STG
  • warp 归约和跨通道集合
  • 寄存器溢出签名
  • 标量和统一控制流模式

初始阶段 3 库包含 patterns/ 下的 29 个模式页面。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 - 跨架构

在更多目标上重放方法论:

架构代表性 GPU原因
SM80A100数据中心 Ampere 基线
SM86RTX 3090消费级 Ampere 语料库
SM89RTX 4090常见消费级推理卡
SM90aH100TMA、WGMMA、warp 专业化、集群
SM100aB200tcgen05.mma、TMEM
SM120RTX 5070 Ti / 5090消费级 Blackwell 起点

仓库地图

root@kitploit:~
.
├── 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 用于记分板、停顿、压力和依赖箭头。
  • Nsight Compute 用于性能分析和停顿归因。
  • %clock 微基准测试用于指令延迟探测。
  • nvcc -Xptxas -v 用于寄存器和溢出元数据。

相关工作

  • Jia 等人 2018 年,“通过微基准测试剖析 NVIDIA Volta GPU 架构” 提供了延迟、吞吐量和依赖性验证的经验性微基准测试方法。
  • kuterdinel.com/nv_isa 用于模糊测试的 NVIDIA ISA 编码工作,特别是从反汇编程序行为推导机器可读编码规则的想法。
  • redplait/denvdis 用于操作码表、位级反汇编、编码字段检查、调度分析、寄存器跟踪和 cubin 操作。提取的 SM120 data12 表 用作低级交叉检查,而本地转储仍然是主要证据。
  • Redplait 工具和笔记:ced cubin 编辑器、SASS 反汇编 Perl 绑定、SASS 延迟分析 和 libcuda/nvasm_internal 笔记。
  • Huerta 等人 2025 年 用于逆向工程编译器引导的调度、控制码、依赖计数器、重用标志和产生行为。
  • Yan 等人 2026 年 用于 SASS 以下的驱动层启动和 pushbuffer 分析。
  • MaxAS 和 TuringAS 作为先前针对较旧 NVIDIA 架构的公开 SASS 汇编器工作。
  • NVIDIA CUDA 二进制工具集 文档用于官方 cubin、fatbin 和反汇编工具。

SASS King 在算法模式层运作:识别编译内核的结构方式,并将这些结构与源代码级别的优化决策联系起来。

贡献指南

欢迎贡献,尤其是:

  • 来自此处无法直接访问的硬件的原始 SASS 转储;
  • 隔离一个编译器决策的受控内核研究;
  • 对现有观察结果的修正;
  • 新的生产内核模式建议;
  • 跨架构比较。

请参阅 CONTRIBUTING.md 了解预期的元数据和写作标准。

作者

Florian Mattana. florianmattana.com

下载工具