florianmattana/sass-king
GitHub: florianmattana/sass-king
系统化反向工程 NVIDIA SASS 并关联 NCU 性能分析，填补 GPU 指令级审计与性能映射的空白。
Stars: 89 | Forks: 3
# SASS King
系统化地对 NVIDIA SASS 进行反向工程，跨越不同架构。
每一项研究都将 **SASS 阅读**（`cuobjdump` + [gpuasm.com](https://gpuasm.com/)）与 **NCU 性能分析** 相结合，以关联指令与实测性能。
完整背景：[第一部分 — 从第一性原理阅读 NVIDIA SASS](https://florianmattana.com/posts/sass_king/)
## 路线图
### 阶段 1 — 在 SM120 上教学内核（可控变化）
* [x] **第一部分** — 内核 01 至 04：基线、FMA 融合、记分板、展开级联。[阅读](https://florianmattana.com/posts/sass_king/)
* [x] 内核 05 — 循环次数较少且固定的循环
* [x] 内核 06 — 共享内存标量（`LDS`、`STS`、`BAR.SYNC`），运行时取模 `CALL`
* [x] 内核 07 — 共享内存模式（库冲突、填充、多缓冲）
* [x] 内核 08 — 向量化全局内存（`LDG.E.128`、`LDG.E.ENL2.256`、FP64）
* [x] 内核 09 — 战线原语（`SHFL.BFLY`、`VOTE`、`MATCH`）
* [x] 内核 10 — 战线规约模式（`REDUX`、蝴蝶、lane-zero）
* [x] 内核 11 — 慢路径算术（`MUFU.RCP/LG2/EX2/RSQ`、内联除法、`log2f`、`expf`、`sinf`、`sqrtf`、Payne-Hanek）
* [x] 内核 12 — 寄存器溢出与局部内存（`STL`、`LDL`、`LDL.LU`、`STL.128`、栈帧、`R2UR`）
* [ ] 内核 13+ — 张量核心（`HMMA`、`QMMA`、`OMMA`）
### 阶段 2 — 经典算法
* [ ] 向量加法
* [ ] 前缀和（扫描）
* [ ] SGEMM
* [ ] 规约（求和、最大值）
* [ ] Softmax
* [ ] LayerNorm / RMSNorm
### 阶段 3 — 库审计
对实际生产内核进行端到端注释。目标如下；内核数量反映 gpuasm.com 上的可用情况。
| 库 | 内核数 | 状态 |
|------------------|--------|--------|
| flash_attn2 | 138 | Planned |
| flash_attn4 | 49 | Planned |
| cutlass (SM120a) | 113 | Planned |
| cute-tutorial | 13 | Planned |
| xformers | 36 | Planned |
| transformer_engine | 109 | Planned |
| flashinfer | 36 | Planned |
| flashmla | 9 | Planned |
| deepep | 2 | Planned |
| llamacpp / ggml | 218 | Planned |
| sglang | 14 | Planned |
| llmc | 8 | Planned |
| tinygrad | 12 | Planned |
| nunchaku | 37 | Planned |
| fouroversix | 57 | Planned |
| bitsandbytes | 2 | Planned |
| arcquant | 24 | Planned |
| qerl | 6 | Planned |
| sgemm | 60 | Planned |
| quack | — | Planned |
### 阶段 4 — 跨架构
在同一研究基础上复现于以下架构：
* [ ] SM80（A100）
* [ ] SM86（RTX 3090）—— gpuasm 示例语料库
* [ ] SM89（RTX 4090）
* [ ] SM90a（H100）
* [ ] SM100a（B200）
### 阶段 5 — 参考
* [ ] 每条指令的 SASS 参考，每条操作码一页，包含每架构的实测延迟、吞吐量、流水线以及双发射规则。
## 目标架构
| 架构 | GPU | 原因 |
|--------|------------------|-----------------------------------------------------|
| SM80 | A100 | 数据中心 Ampere 基线 |
| SM86 | RTX 3090 | 消费级 Ampere，gpuasm 示例语料库 |
| SM89 | RTX 4090 | 最常见的消费级推理卡 |
| SM90a | H100 | TMA、WGMMA、战线专业化、mbarrier、集群 |
| SM100a | B200 | `tcgen05.mma`、TMEM |
| SM120 | RTX 5070 Ti/5090 | 混合 SM90/SM100 ISA，`mma.sync` 搭配 `mxf8f6f4` |
工作从 SM120 开始（直接硬件访问）。其他架构通过公开转储和贡献者进行。
## 工具
* `cuobjdump --dump-sass` 用于原始反汇编
* [gpuasm.com](https://gpuasm.com/) 用于查看记分板、停顿、压力与依赖箭头
* Nsight Compute（NCU）用于每内核性能分析、SASS 到源码映射以及停顿归因
## 作者
Florian Mattana。 [florianmattana.com](https://florianmattana.com)
标签：cuobjdump, cute-tutorial, cutlass, flash_attn2, flash_attn4, FMA 融合, FP64, gpuasm.com, GPU 指令级分析, GPU 架构, HMMA, LDG.E.128, LDL, MATCH, MUFU, NCU 性能分析, NVIDIA SASS, OMMA, QMMA, REDUX, SASS 阅读, SHFL.BFLY, SM120, STL, transformer_engine, Vectored Exception Handling, VOTE, warp 原语, warp 约简, Waymore结果处理, xformers, 共享内存, 内核审计, 反向工程, 向量化全局内存, 寄存器溢出, 张量核心, 性能调优, 本地内存, 模式识别, 生产内核审计, 银行冲突