PTX 与 SASS：GPU 内核的中间表示与微调方法
说明：本篇为综述/无第一手实验数据，采用公开编译链文档与架构解析路线。
实验目标
08 章的 CUDA 编译管线篇讲了 nvcc 怎么把 .cu 编成可执行文件；本篇继续往深一层：编译中间到底生成了什么，以及工程师何时需要"看汇编"来调优。目标：搞清 PTX 与 SASS 两种中间表示的分工与生成路径、用 cuobjdump/nvdisasm 把内核反汇编出来的方法、以及三类"看汇编才发现的优化"（指令选择、寄存器压力、尾效应），并明确"什么时候该看、什么时候别浪费时间"。
代码逐段拆解
第一段：从 .cu 到 SASS 的编译链
bash
# 常规编译：nvcc 生成 fatbin（内含 PTX 与各架构 SASS）
nvcc -arch=sm_90 -O3 -o kernel kernel.cu
# 只看 PTX（虚拟中间表示，跨架构）
nvcc -arch=compute_90 -ptx -o kernel.ptx kernel.cu
# 反汇编 SASS（真实机器码，绑定具体架构）
cuobjdump -sass kernel > kernel.sass
# 或：nvdisasm 直接反汇编 cubin
nvdisasm -c kernel.cubin > kernel.sass要点：
- PTX 是"虚拟 ISA"，SASS 是"真机 ISA"：PTX 面向
compute_XX（虚拟架构），跨代通用，是 JIT 与驱动兼容的基石；SASS 面向sm_XX（真实架构），是最终执行的机器码。-arch=sm_90时 nvcc 既嵌 PTX 又嵌 SASS。 cuobjdump -sass是看真实执行码的入口：优化器做的寄存器分配、指令调度、bank conflict 消除都在 SASS 层可见——PTX 看不到。- 反汇编是"读"，不是"改"：工程师不手写 SASS（除非为特定架构做极致的 FFMA 合并），反汇编的价值在于确认编译器是否做了期望的优化，以及发现它没做的优化。
第二段：SASS 里值得看的四类信号（简化示例）
sass
/* 观察点 1：FFMA 是否合并 —— FFMA R6, R4, R5, R6 一条指令完成乘加 */
/* 观察点 2：寄存器压力 —— 一个 warp 用的寄存器越多，占用率越低 */
/* 观察点 3：LDL/STL 访存指令 —— 是否向量化（LDG.E.128 一次 16B） */
/* 观察点 4：Bank 冲突 —— SASS 层面看共享内存访问的 bank 分布 */解读方法：
- 乘加融合（FFMA）：
a*b+c在 SASS 里应是单条 FFMA；如果编译器生成了 FMUL + FADD 两条，说明没吃到融合，代码写法可能阻碍了它（如中间变量写法）。 - 访存向量化：连续的 float4 加载应表现为
LDG.E.128；如果退化成 4 条LDG.E.32，访存效率掉 4 倍，改代码对齐可修复。 - 寄存器压力：SASS 里每条指令的寄存器编号跨度能看出每个 thread 用了多少寄存器；超过阈值（如 128 个）会压降占用率，对应 Nsight 里的低 Occupancy。
原理讲解
为什么有两层中间表示
- PTX（虚拟 ISA）存在的理由：兼容性。同一份 PTX 能在新架构上 JIT 出新 SASS，老程序不重编译就能跑新卡——这是 CUDA"向前兼容"承诺的载体。
- SASS（真实 ISA）存在的理由：性能。每个 sm_XX 架构的指令集、寄存器文件、内存层次都不同，SASS 是按架构特化的最优码。
- 两层之间的分工：编译器在前端生成 PTX（语言层优化），后端把 PTX 变成 SASS（架构层优化：寄存器分配、指令调度）。工程师看性能问题，主要看后端产物 SASS。
三类"看汇编才发现的优化"
- 指令选择问题：编译器选了低效指令序列（没融合的乘加、多余的 mov、未向量化的访存）。修法通常在源代码——改变写法（
float4结构、#pragma unroll、消除中间变量）让编译器"看见"融合机会。 - 寄存器压力问题：一个 thread 寄存器太多 → 占用率低 → 延迟无法掩盖。SASS 里确认后，用
__launch_bounds__限制寄存器上限或重构 kernel。 - 尾效应与分支发散：SASS 里的
BRA与 predicated 指令能看出 warp 内是否发散。发散会让 warp 串行执行分支，是"占满利用率但效率低"的典型来源。
何时看汇编（判断准则）
| 场景 | 是否值得看 SASS |
|---|---|
| 内核已是访存/计算瓶颈（Nsight 判定） | 看：确认访存是否向量化、计算是否融合 |
| 占用率低但代码看起来正常 | 看：确认寄存器压力 |
| 一般业务代码、非热点内核 | 不看：编译器的默认优化已够好 |
| 性能调优的最后 20% | 看：这是收尾手段，不是起点 |
关键认知：SASS 调优是"证据确认"而不是"猜测优化"——先用 Nsight（上一篇）定性瓶颈，再看 SASS 确认机制，最后改源代码。跳过定性直接翻汇编，容易在大海捞针。
实验结果
（估算口径，非实测）
| 观察信号 | 对应问题 | 修复动作 | 典型收益 |
|---|---|---|---|
| 未融合 FMUL+FADD | 乘法加法未融合 | 改写法让编译器融合 | 1.1-1.3× |
| LDG.E.32 ×4 | 访存未向量化 | float4 对齐 | 1.5-3× |
| 寄存器 >128/thread | 占用率被压 | __launch_bounds__ | 1.2-1.5× |
| 分支发散明显 | warp 串行化 | 数据重排消除发散 | 1.3-2× |
工程启示
- SASS 是确认工具，不是优化起点：先 Nsight 定性（上篇），再看 SASS 确认（本篇），最后改源码——顺序反了就是事倍功半。
- 修代码而不是改汇编：99% 的情况应该改源代码让编译器生成更好的 SASS，手改 SASS 只在极端收尾场景有意义。
- 把反汇编结果留档：性能敏感内核的 SASS 与 ncu 报告一起存基线，升级编译器/驱动后 diff 一次，防止无声的性能回退。
- 编译器版本影响 SASS：同一源码不同 CUDA 版本生成的 SASS 可能差 10%-20%——升级 CUDA 后做回归对比是必须的纪律。
下一篇预告：编译器与内核链路补完后，回到框架层——"DeepSpeed ZeRO：训练显存优化的三阶段原理与实践"。