GPU系统拆解-07-PTX、SASS 与编译链：一段 CUDA 代码如何变成指令

本文是「GPU系统拆解」系列第 07 篇。
系列导读：GPU系统拆解-00-导读：从架构认知到推理系统的学习路线
上一篇：GPU系统拆解-06-内存层次与访存主线：GPU 性能为什么常常卡在数据
下一篇：GPU系统拆解-08-Tensor Core、GEMM 与 FlashAttention：AI 计算热点为什么这样组织

这一篇要解决的是一个很容易被忽略的问题：你写下的 CUDA 源码，最终到底变成了什么。学完这一篇之后，你不需要会手写 SASS，但应该知道什么时候该看 PTX、什么时候该看 SASS，以及它们为什么能帮助你解释 kernel 的真实性能表现。

1. 先给结论

  • CUDA 源码不会一步直接变成 GPU 机器码，中间通常还要经过 PTX 这一层。
  • PTX 是虚拟 ISA 和中间表示，SASS 才是某一具体 GPU 架构真正执行的机器指令。
  • 同一份 .cu 源码，在不同 GPU 上最终生成的机器码可能不同，所以性能也可能不同。
  • 很多性能问题不是算法本身的问题，而是编译结果不理想，例如寄存器过多、spill、指令膨胀、没有走到预期的矩阵路径。
  • PTX / SASS 不是日常开发起点，而是定位复杂性能问题时的重要证据层。
  • 对 AI infra / 推理 来说，Tensor Core 没用上、小 kernel 底层效率差、变长序列导致控制流复杂、寄存器压力过高，这些都很值得从 PTX / SASS 角度去看。

2. 先把整条编译链记住

可以先把 CUDA 代码的大致链路记成这样：

1
2
3
4
5
6
7
8
CUDA C++ 源码
  -> 前端编译
  -> PTX
  -> ptxas / 后端
  -> cubin / fatbin
  -> 驱动加载
  -> 必要时 JIT
  -> GPU 执行

这条链最重要的意义不是记名词，而是建立一个判断：

你写的源码只是最上层意图，真正落到 GPU 上执行的是更底层的结果。

2.1 前端在做什么

前端主要负责把 CUDA C++ 里的这些东西展开和整理：

  • __global__、__device__
  • 模板
  • 内联
  • 常量传播
  • 循环展开
  • 各种语义分析

到这一步，代码已经开始从“人类写的程序”往“GPU 可执行程序”靠近。

2.2 PTX 在整条链里的位置

PTX 不是最终机器码，而是一层更接近 GPU 的中间表示。

它已经会出现很多底层概念：

  • 寄存器
  • 地址空间
  • load/store
  • 谓词
  • 分支
  • 某些矩阵路径相关操作

所以 PTX 已经不是“算法描述”，而是“GPU 视角下的程序描述”。

2.3 cubin、fatbin、JIT 应该怎么理解

可以先用最实用的方式理解：

  • cubin：某个具体架构可直接执行的设备二进制
  • fatbin：打包了多种设备代码或中间代码的容器
  • JIT：驱动在运行时根据当前 GPU，把中间表示进一步编译成适合本机的机器码

你不用在这一篇里把所有工具链细节背全，但至少要知道：

  • 有时程序会直接带着某些架构的已编译结果
  • 有时也会保留 PTX，留给驱动在运行时再做最后一段编译

3. PTX 和 SASS 到底分别是什么

3.1 PTX：虚拟 ISA 和中间层

PTX 的定位是：

介于高层 CUDA 源码和最终机器码之间的虚拟指令集。

它的价值在于两点：

  • 给上层提供相对稳定的抽象
  • 给下层后端和驱动保留优化空间

这就是为什么 NVIDIA 不直接要求大家面向具体 GPU 机器码编程。

3.2 SASS：真正面向具体架构的机器指令

如果说 PTX 是“还没完全落地的中间表示”，那么 SASS 就是某一具体架构真正执行的机器指令表示。

它更贴近这些问题：

  • 这条 load/store 最终怎么发射
  • 这个分支最后怎么实现
  • 寄存器怎么分配
  • 某些架构特性有没有真正被用上

所以：

  • 想看“程序大致被翻译成了什么”，PTX 往往就很有帮助
  • 想看“某张具体卡上最后到底跑了什么”，SASS 更有解释力

4. 为什么 AI infra / 推理工程师需要理解它们

先纠正一个误区：

不是只有写汇编的人才需要懂 PTX / SASS。

大多数 GPU 工程师并不会手写 SASS。主流路径仍然是：

  • CUDA C++
  • 各种模板和库
  • CUTLASS / Triton / TVM 等生成代码

但你仍然要懂 PTX / SASS 的意义，因为很多关键问题只看源码解释不通。

例如：

  • 为什么加了一个看似无害的临时变量后 kernel 变慢了
  • 为什么寄存器数突然暴涨
  • 为什么 occupancy 明明不低，性能还是很差
  • 为什么你以为会走 Tensor Core，结果没走上
  • 为什么相同源码在 Ada 和 Hopper 上表现差异明显

这些问题的共同点是：

你必须知道编译器和后端最终把源码翻译成了什么。

5. 从源码到 PTX，思维方式为什么会变

先看一个最简单的例子：

1
2
3
4
5
6
__global__ void vec_add(const float* a, const float* b, float* c, int n) {
    int idx = blockIdx.x * blockDim.x + threadIdx.x;
    if (idx < n) {
        c[idx] = a[idx] + b[idx];
    }
}

从源码看，你只看到了：

  • 索引计算
  • 一个边界判断
  • 两次读取
  • 一次加法
  • 一次写回

但从 PTX / 底层视角看，这段逻辑通常会拆成更多动作：

  1. 读取内建索引寄存器
  2. 做整数乘加，计算 idx
  3. 比较 idx < n
  4. 形成谓词
  5. 计算 a[idx]、b[idx]、c[idx] 的地址
  6. 发出 load
  7. 发出浮点加法
  8. 发出 store

这里要建立的核心直觉是：

源码里一句很普通的数组访问，落到 GPU 上往往已经拆成了一串地址计算、寄存器分配、load/store 和控制流处理。

所以很多性能问题并不是“加法慢”，而是：

  • 地址算得太复杂
  • 多了额外的 load/store
  • 对齐不好
  • 谓词或分支让活动线程模式变差
  • 寄存器占用过高

6. 谓词、分支和“看似简单的 if”

GPU 控制流里一个很重要的概念是 predicate。

很多高层的 if，在底层不一定变成你想象中的粗暴跳转，也可能变成“带条件执行”的指令。

例如：

1
2
3
if (idx < n) {
    x = y + z;
}

底层可以粗略理解成：

1
2
3
p = (idx < n)
@p 执行 add
@p 执行 store

这意味着两件事：

  • 即使你只写了一个很小的边界判断，底层也确实存在控制流成本
  • 在 warp 级执行模型下，条件是否整齐、活动线程是否集中，会影响实际执行效率

这类问题在推理里很常见，尤其是：

  • 变长序列
  • ragged batch
  • 边界 tile
  • decode 阶段 token 不齐
  • padding 和 mask 处理

所以很多“逻辑很简单”的 kernel，底层未必简单。

7. 寄存器为什么总是第一嫌疑人

很多人一开始学 GPU 优化只盯 shared memory，后面才会意识到：

寄存器往往才是最容易不知不觉把性能搞坏的资源。

7.1 为什么寄存器这么重要

寄存器是线程私有、速度最快的存储层。中间结果如果能一直留在寄存器里，通常会有：

  • 更低延迟
  • 更少额外访存
  • 更好的算术流水

所以编译器会尽量把活跃变量放在寄存器里。

7.2 为什么寄存器又会变成问题

因为寄存器不是无限的。

如果每线程用的寄存器太多，会直接带来两个结果：

  • 同一个 SM 上能驻留的 warp 和 block 更少
  • occupancy 下降

更糟的是，当寄存器放不下时，会发生 spill。

7.3 spill 为什么危险

当变量 spill 到 local memory 时，名字看起来像“本地内存”，但它通常并不是真正意义上的快。

更实用的理解是：

spill 往往意味着本来应该待在寄存器里的数据，被迫变成了更贵的内存访问。

于是就会出现这种常见链条：

1
2
3
4
5
6
代码复杂一点
  -> 活跃变量增多
  -> 寄存器数上升
  -> spill
  -> 额外 load/store 增多
  -> kernel 变慢

这类问题单看源码经常看不出来，必须结合编译结果和 profile 一起判断。

8. 为什么“源码更优雅”不一定“机器码更高效”

这是 GPU 学习里必须尽快接受的一件事。

软件工程视角里更优雅的写法，未必能生成更高效的底层代码。

常见原因包括：

  • 引入更多临时对象
  • 把索引封装得太深
  • 看似更通用的函数拆分导致内联不理想
  • 模板层次更复杂，但没有换来更好的静态展开
  • 为了代码复用，引入了额外地址计算和控制流

这些都可能导致：

  • 指令条数增多
  • 寄存器占用升高
  • load/store 更多
  • 编译器更难做激进优化

所以高性能 GPU 代码经常看起来没有那么“优雅”，不是作者不会写，而是优先级不同：

高性能 kernel 首先是为了生成好代码，其次才是源码形式漂亮。

9. 看 PTX / SASS 时，先找哪些“坏味道”

工程实践里，最有效的方式不是逐条啃底层指令，而是先找明显异常。

9.1 指令条数异常多

如果一个逻辑很简单的 kernel，底层却膨胀出很多地址计算、移动、类型转换，通常说明：

  • 索引表达式太复杂
  • 冗余操作没消掉
  • 编译器没能很好地简化代码

9.2 大量 local memory 访问

这通常是 spill 或对象无法留在寄存器里的信号。

9.3 load/store 粒度不理想

如果你预期应该是规整访问，但底层变成很多零散读写，通常说明访存路径没有走顺。

9.4 没走到预期的矩阵路径

如果你原本期待某个 GEMM 或 attention 核心部分能走到 Tensor Core 路径，但底层没体现出来，就说明：

  • 数据类型不合适
  • 布局不合适
  • tile 组织不合适
  • 调用方式没满足后端优化条件

9.5 分支和谓词过重

这在不规则 workload 中很常见，尤其是：

  • decode
  • 稀疏
  • 变长序列
  • 边界处理很多的 kernel

这也是为什么看 PTX / SASS 的目的不是“欣赏底层代码”，而是快速验证你对瓶颈的猜测。

10. CUDA 编译器到底在帮你做什么

你应该把 CUDA 编译器理解成：

在可读源码和可执行 GPU 代码之间做大量转换工作的核心后端。

它会做很多事：

  • 内联
  • 常量传播
  • 循环展开
  • 寄存器分配
  • 指令选择
  • 地址计算优化
  • 调度
  • 面向具体架构的 lowering

这带来一个很重要的工程判断：

你的源码意图，不等于最终执行结果。

所以成熟的 GPU 工程思维不是只看源码，而是一起看：

  • 源码设计
  • 编译器行为
  • 硬件执行

11. PTX / SASS 和架构演进是什么关系

前面你已经学过 Ada、Hopper 这些架构差异，这一篇要把那些差异落到“代码最终怎么执行”上。

关键点是：

  • 新架构会引入新能力
  • 这些能力不会自动无条件生效

例如：

  • 更强的 Tensor Core 路径
  • 更好的异步搬运支持
  • 更适合矩阵或 attention 的底层能力

如果你的代码组织方式不对，后端未必会把它映射到最优路径。

所以：

架构升级不等于你的 kernel 自动吃满新特性。

中间还隔着：

  • 编程模型
  • 编译器
  • 库实现
  • 数据布局
  • launch 参数
  • tile 策略

这也是为什么 CUTLASS、cuBLAS、TensorRT 这类库往往能系统性地利用底层路径，而朴素 kernel 不行。

12. 在推理场景里，哪些问题值得从 PTX / SASS 角度去看

这一部分最贴近 AI infra / 推理。

12.1 小 kernel 很多，但不只是 launch 开销问题

像这些操作经常很碎：

  • layernorm
  • softmax
  • elementwise
  • decode 里的细碎后处理

它们的问题不只是 launch 多，还可能是底层代码本身就不够紧凑。

12.2 变长序列会带来复杂控制流

一旦涉及：

  • padding
  • mask
  • ragged batch
  • decode 长度不齐

底层就更容易出现谓词和分支负担。

12.3 你以为会用到 Tensor Core，结果没用上

这类问题尤其值得看 PTX / SASS，因为它通常意味着“你的高层意图没有变成后端预期的矩阵路径”。

12.4 寄存器压力导致 occupancy 下降或 spill

这在复杂 attention kernel、fused kernel、decode kernel 里很常见。

看 profile 只能告诉你“有问题”，看 PTX / SASS 往往更能告诉你“问题怎么来的”。

13. 什么时候该看 PTX，什么时候先 profile 就够了

一个更成熟的顺序是：

13.1 先 profile，再决定要不要往下看

不要把 PTX / SASS 当作日常第一步。更合理的顺序通常是：

  1. 先 profile
  2. 先定位最慢的 kernel
  3. 再判断是否需要往编译结果这一层下钻

13.2 这些情况非常值得看

  • 某次改动后 kernel 明显变慢
  • 寄存器数异常升高
  • occupancy 意外下降
  • 预期中的 Tensor Core 路径没出现
  • 底层访存条数和你预期差很多
  • 同一份代码在不同架构上表现差别明显

13.3 这些情况通常先不用看

  • 还没做最基本的 profile
  • 还没确认瓶颈 kernel 是谁
  • 问题更像系统层调度、I/O、通信或 batch 策略问题

也就是说：

PTX / SASS 不是第一层视角，而是当你需要证据时再往下走的一层。

14. 这一篇必须记住的几句话

  • CUDA 源码不是最终执行形态，中间通常还隔着 PTX 和后端编译。
  • PTX 是虚拟 ISA 和中间表示，SASS 是某一具体 GPU 架构真正执行的机器指令。
  • 你的源码意图不等于最终生成的底层代码。
  • 高性能 CUDA 开发必须一起看源码、编译器行为和硬件执行。
  • 寄存器、spill、指令膨胀、分支和访存路径，都是底层性能问题的高频来源。
  • PTX / SASS 不是炫技工具，而是复杂问题的验证工具。

15. 精简版面试表达

如果面试官问为什么 CUDA 工程师要看 PTX / SASS，一个更成熟的回答可以是：

因为 CUDA 源码不是最终执行形态。源码会先变成 PTX，再通过后端或驱动 JIT 变成某个具体架构上的机器码。很多性能问题，比如寄存器过多、spill、指令膨胀、访存不理想，或者没走到预期的 Tensor Core 路径，只看源码很难解释，必须看编译结果才能确认问题到底出在算法设计、编译器行为，还是底层硬件映射上。

系列导航