Qingyang-Lucas-Fang/OpenAssembler
开发工具面向 sm_80 的 PTX 到 SASS 逆向编译器,替换 ptxas 代码生成并优化内核。
项目介绍
OpenAssembler 是 NVIDIA 闭源 ptxas 的第二套实现:把 ptxas、nvdisasm 和 GPU 当作黑盒逆向,为 sm_80(Ampere)实现 SASS 编解码。它复用 ptxas 生成的 cubin 作为 ELF 容器,只替换指令字、声明的寄存器数量和退出偏移表,用自研后端生成代码(source:7)。
编译流程依次是 PTX 到 IR、删除无用代码与复制合并、线性扫描寄存器分配、lowering、可选的 4 倍循环展开,以及基于硬件实测延迟的表调度,最后用自研编码器输出机器码(source:10)。前端支持 PTX、Python 子集和 CUDA C/C++,其中 CUDA C 先由 nvcc 降为 PTX(source:16)。
仓库还包含贴合 A100 行为的 SASS 模拟器、用 z3 证明的常量除法优化、541 个探针的测试语料,以及 Strassen/AlphaTensor 快速矩阵乘法方案。
快速使用
环境要求:Python 3.12+,Linux 或 macOS(不支持 Windows),以及 CUDA 12.8 中的 ptxas、nvdisasm 和可选的 nvcc。编译不需要 GPU,运行结果才需要。
git clone https://github.com/Qingyang-Lucas-Fang/OpenAssembler.git && cd OpenAssembler
python3 -m venv .venv && source .venv/bin/activate
pip install -e .
env/fetch_cuda.sh
env/cuda.sh ptxas --version # 应输出 release 12.8
Debian/Ubuntu 需要先安装 python3-venv 和 xz-utils,编译 CUDA C 还需要宿主机 g++(source:20)。macOS 用户需先用 colima 启动 Linux 虚拟机(source:25)。如果本机已有 CUDA 12.8,可设置 OPENASSEMBLER_CUDA_HOME 或把工具加入 PATH(source:27)。
把内核编译成可加载的 cubin 和可读的 SASS,并对照参考解释器做逐位验证:
openassembler compile examples/div_by_const.ptx -o out.cubin --verify
同样的命令也接受仓库中的 Python 和 CUDA C 示例(examples/vector_add.py、examples/vector_add.cu)。直接运行 openassembler 会打开交互式 shell,其中可以省略 openassembler 前缀(source:31)。
要在真实 GPU 上运行探针语料,先构建语料并用本项目后端编译,然后在本机 GPU、SSH 主机或 Modal 上运行:
openassembler corpus
openassembler check
openassembler gpu run --backend local --which ours
使用约束
- 编码器和延迟模型只针对 sm_80;sm_86/89/90/100/120 等架构暂不支持(source:18)。
- 它并非端到端替代
ptxas:仍需要ptxas生成的 cubin 作为 ELF 模板,也未覆盖全部 PTX。缺失的功能包括 f16/f64 算术、warp shuffle/vote、寄存器溢出、浮点除法与开方、ld.global.v2,以及超过 7 个活跃谓词(source:13)。 - 实测加速只针对列出的内核形状和 A100;紧凑依赖链仍比
ptxas慢 5–8%,直接 GEMM 在 n = 8192 时仍落后 cuBLAS(source:46)。 - 许可证为 Apache-2.0,
third_party/下的 AlphaTensor 分解在 NOTICE 中单独标注(source:55)。
HotRepo 点评
2026-11-02 的 Python 新项目榜记录显示,该项目在首条快照中为 104 star、3 fork,并在约一小时的统计窗口内新增 3 star 和 1 fork。观察周期很短,这只能说明早期关注度,不足以判断趋势。项目的信息密度主要来自对 sm_80 指令编码的逆向,以及 A100 上带复现命令的实测结果,而不是泛化的性能主张。
本文由 AI 根据项目 README 生成,更新于 ,以项目官方文档为准。