GPU 架构学习
第十八课：PTX、SASS 与 Compute Capability
前面讨论的是 GPU 硬件如何执行指令，但程序员通常写的是 CUDA C++：
目录
一、CUDA 源文件中其实包含两种程序二、nvcc 到底是什么三、什么是 PTX四、PTX 类似 CPU 世界中的什么五、一段简单 PTX 长什么样六、PTX 中的寄存器是真实物理寄存器吗七、PTX 是 GPU 直接执行的吗八、什么是 SASS九、PTX 与 SASS 的核心区别十、什么是 ptxas十一、什么是 Cubin十二、完整编译路径十三、什么是 Compute Capability十四、常见架构与 Compute Capability十五、compute_80 与 sm_80 有什么区别compute_80sm_80十六、为什么要分虚拟目标和真实目标十七、最简单的 nvcc 架构选项十八、如何只生成 PTX十九、如何生成特定架构 Cubin二十、什么是 Fat Binary二十一、如何编译多个架构版本二十二、运行时如何选择代码二十三、什么是 JIT 编译二十四、JIT 编译有什么代价二十五、Cubin 的兼容性二十六、为什么不只保留 PTX二十七、为什么不为所有 GPU 都生成 Cubin二十八、什么是 sm_90a 和 sm_100a二十九、为什么 sm_90 和 sm_90a 不相同三十、什么是 f 后缀三十一、如何查看 GPU 的 Compute Capability三十二、__CUDA_ARCH__ 是什么三十三、一个 PTX 指令是否对应一个 SASS 指令一对一一对多多对一完全消失被替换三十四、为什么 SASS 才能真正反映寄存器和指令三十五、如何查看可执行文件中的 PTX三十六、如何查看 SASS三十七、如何同时查看源码行和 SASS三十八、如何查看寄存器使用量三十九、什么是 Register Spill四十、为什么同一源代码在不同架构上寄存器数量不同四十一、为什么手写 PTX 不一定更快四十二、内联 PTX四十三、SASS 能否直接手写四十四、PTX 版本和 Compute Capability 不是同一概念.version.target四十五、为什么新驱动能够支持旧程序四十六、为什么库经常很大四十七、从 GPU 架构设计角度理解 ISA 分层四十八、从 RTL 角度看 SASS 指令执行四十九、实际分析一个 Kernel 的推荐步骤第一步：看 CUDA C++第二步：看编译报告第三步：看 PTX第四步：看 SASS第五步：使用性能分析器五十、本课核心结论第十九课：编译器如何优化 CUDA Kernel前面讨论的是 GPU 硬件如何执行指令，但程序员通常写的是 CUDA C++：
__global__ void vector_add(
const float* A,
const float* B,
float* C,
int n)
{
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < n) {
C[i] = A[i] + B[i];
}
}GPU 并不能直接执行这段 C++ 代码。
它需要经过：
CUDA C++
↓
PTX
↓
SASS / Cubin
↓
GPU 硬件执行理解这一过程，可以帮助分析：
- 编译器生成了什么指令；
- 为什么寄存器使用量会变化；
- 为什么同一程序能运行在不同 GPU 上；
- 为什么有时第一次启动 Kernel 较慢；
sm_80、sm_90、sm_100到底是什么意思。
一、CUDA 源文件中其实包含两种程序
一个 .cu 文件通常同时包含：
Host Code：在 CPU 上执行
Device Code：在 GPU 上执行例如：
int main()
{
// CPU 代码
cudaMalloc(...);
// 启动 GPU Kernel
vector_add<<<grid, block>>>(...);
cudaDeviceSynchronize();
}其中：
vector_add<<<grid, block>>>(...);是 Host 代码发出的 Kernel 启动命令。
而：
__global__ void vector_add(...)函数体中的代码在 GPU 上执行。
nvcc 会分离 Host 和 Device 部分：Host 代码交给主机 C++ 编译器，Device 代码由 NVIDIA 的设备编译链处理，最后再把 GPU 代码嵌入主机目标文件或可执行文件。
二、nvcc 到底是什么
nvcc 是 NVIDIA CUDA Compiler Driver。
它并不是一个单独完成所有工作的简单编译器，而是一个编译驱动程序，负责协调：
- CUDA 前端；
- Host C++ 编译器；
- PTX 生成器；
ptxas；- Device Linker；
- Host Linker；
- Fat Binary 生成过程。
因此执行：
nvcc vector_add.cu -o vector_add背后可能包含很多编译和链接步骤。
可以简化为：
图示：二、nvcc 到底是什么
nvcc 的作用就是隐藏和组织这些复杂阶段。
三、什么是 PTX
PTX 全称：
Parallel Thread Execution。
PTX 是 NVIDIA 定义的一种虚拟并行指令集。
它是一种：
低级
接近汇编
但并不绑定某一代具体 GPU的中间表示。
官方将 PTX 定义为一种低级并行线程虚拟机和虚拟指令集，目标之一就是跨越多代 GPU 架构，作为高级语言编译器和真实 GPU 指令之间的中间层。
四、PTX 类似 CPU 世界中的什么
可以类比为：
CUDA C++
↓
PTX
↓
真实 GPU 指令类似于：
Java
↓
Java Bytecode
↓
特定 CPU 机器码或者：
C/C++
↓
LLVM IR
↓
x86 / ARM 机器码但这个类比不是完全相同，只是帮助理解：
PTX 描述的是一个抽象 GPU 计算模型，而不是某块具体 GPU 的最终机器指令。
五、一段简单 PTX 长什么样
下面是经过简化的向量加法 PTX 风格代码：
.visible .entry vector_add(
.param .u64 A,
.param .u64 B,
.param .u64 C,
.param .u32 n)
{
.reg .pred %p;
.reg .b32 %r<6>;
.reg .b64 %rd<10>;
.reg .f32 %f<4>;
mov.u32 %r1, %ctaid.x;
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %tid.x;
mad.lo.s32 %r4, %r1, %r2, %r3;
setp.ge.s32 %p, %r4, %r5;
@%p bra DONE;
ld.global.f32 %f1, [%rd1];
ld.global.f32 %f2, [%rd2];
add.f32 %f3, %f1, %f2;
st.global.f32 [%rd3], %f3;
DONE:
ret;
}真实编译结果通常更复杂，但可以看出一些基本结构：
图示：五、一段简单 PTX 长什么样
六、PTX 中的寄存器是真实物理寄存器吗
不完全是。
例如：
.reg .f32 %f<4>;表示 PTX 虚拟机中声明了若干虚拟寄存器。
这些虚拟寄存器并不直接等同于 SM Register File 中的某几个固定物理寄存器。
后面的 ptxas 会进行：
- 指令选择；
- 优化；
- 寄存器分配；
- 指令调度；
- 地址模式合并；
- 常量传播；
- 死代码消除。
最后决定真正需要多少硬件寄存器。
因此：
PTX 寄存器数量
≠
最终 SASS 物理寄存器数量分析 Occupancy 时，应关注最终编译报告或 Cubin 中的资源使用量，而不是只数 PTX 中的 .reg。
七、PTX 是 GPU 直接执行的吗
通常不是。
GPU 硬件真正执行的是对应架构的原生机器指令。
PTX 还需要被转换为：
SASS这个过程可以发生在：
- 程序编译时；
- 程序安装时；
- 程序第一次运行时。
PTX 的主要价值之一，就是允许驱动在知道目标 GPU 型号后，将其转换为相应架构的原生指令。
八、什么是 SASS
SASS 通常指 NVIDIA GPU 的原生机器指令或其反汇编表示。
可以理解为：
某一代真实 GPU 硬件最终执行的指令。
例如不同架构可能出现以下类型的 SASS 指令助记符：
S2R
IMAD
IADD3
LDG
STG
FADD
FFMA
FSETP
BRA
EXIT概念上可能形成：
S2R R0, SR_CTAID.X;
S2R R1, SR_TID.X;
IMAD R2, R0, c[...], R1;
LDG.E R4, [R8];
LDG.E R5, [R10];
FADD R6, R4, R5;
STG.E [R12], R6;
EXIT;具体指令名称、编码和操作数形式会随 GPU 架构变化。
九、PTX 与 SASS 的核心区别
| 对比项 | PTX | SASS |
|---|---|---|
| 类型 | 虚拟 ISA | 原生硬件 ISA |
| 面向对象 | 一类计算能力 | 某个真实 SM 架构 |
| 可读性 | 较高 | 更接近机器指令 |
| 跨代兼容 | 较强 | 较弱 |
| 寄存器 | 虚拟寄存器 | 物理寄存器编号 |
| 生成阶段 | CUDA 前端生成 | ptxas 或 Driver JIT 生成 |
| GPU 能否直接执行 | 通常不能 | 可以 |
| 典型文件 | .ptx | .cubin 中的代码 |
最简洁的理解是：
PTX：面向虚拟 GPU 的汇编
SASS：面向真实 GPU 的汇编十、什么是 ptxas
ptxas 可以理解为：
PTX 汇编器和后端优化器。
它把 PTX 转换为特定 GPU 架构的原生机器代码。
例如：
PTX for compute_80
↓ ptxas
SASS for sm_80在这个阶段会确定：
- 使用哪些真实指令；
- 物理寄存器如何分配；
- 是否产生 Local Memory Spill；
- 指令如何调度；
- 立即数和常量如何编码；
- 某些 PTX 指令如何拆分或合并。
因此同一段 PTX 针对不同架构编译，可能生成不同 SASS。
十一、什么是 Cubin
Cubin 可以理解为：
CUDA Binary，即包含特定 GPU 原生代码的二进制文件。
它通常是 ELF 格式，除了机器指令，还可能包含：
- 符号表；
- 重定位信息；
- Kernel 元数据；
- 寄存器使用量；
- Shared Memory 使用量；
- 调试信息；
- 常量数据。
官方 CUDA Binary Utilities 文档将 Cubin 描述为包含 CUDA 可执行代码和相关元数据的 ELF 文件。
可以通过下面的命令单独生成：
nvcc -O3 -arch=sm_80 -cubin kernel.cu -o kernel_sm80.cubin这个 Cubin 面向 sm_80 架构。
十二、完整编译路径
一段 CUDA Device 代码大致经历：
CUDA C++
↓
CUDA Front End
↓
PTX
↓
ptxas
↓
SASS
↓
Cubin而主机代码经历：
CUDA / C++ Host Code
↓
GCC / Clang / MSVC
↓
x86 或 ARM 机器码最终：
CPU 机器码
+
GPU Cubin
+
可能保留的 PTX
+
CUDA Runtime 信息
↓
可执行文件十三、什么是 Compute Capability
Compute Capability，计算能力，是 NVIDIA 用来描述 GPU 计算架构功能集合的版本号。
通常写成：
例如：
7.0
7.5
8.0
8.6
8.9
9.0
10.0
12.0它描述的是：
- 支持哪些指令；
- 支持哪些数据类型；
- SM 的资源上限；
- 支持哪些同步功能；
- 是否支持 TMA；
- 是否支持 Thread Block Cluster；
- 原子操作能力；
- Tensor Core 功能。
Compute Capability 不是一个简单的性能评分。
因此：
Compute Capability 9.0不表示性能一定是：
Compute Capability 8.0 的 1.125 倍它主要表示架构功能版本。
十四、常见架构与 Compute Capability
常见对应关系可以简化为：
| Compute Capability | 常见架构或产品类别 |
|---|---|
| 7.0 | Volta |
| 7.5 | Turing |
| 8.0 | Ampere A100 类 |
| 8.6 | Ampere GA10x 类 |
| 8.9 | Ada Lovelace |
| 9.0 | Hopper |
| 10.0 | Blackwell 数据中心架构的一类 |
| 12.0 | Blackwell GeForce RTX 50 系列的一类 |
需要注意：
同一个市场宣传架构名称，不一定只对应一个 Compute Capability。
例如数据中心 Blackwell SM100 与 GeForce RTX 50 系列 SM120 具有不同的计算能力目标，使用架构专用功能编译的 SM100 代码不能直接视为与 SM120 通用。
十五、compute_80 与 sm_80 有什么区别
这是最容易混淆的地方。
compute_80
表示：
面向 Compute Capability 8.0 功能集合的虚拟 PTX 架构。
它决定 CUDA 前端可以假设哪些功能存在。
输出通常是 PTX。
sm_80
表示：
面向真实 SM 8.0 架构生成原生机器代码。
输出通常是 Cubin/SASS。
所以：
compute_80 → 虚拟目标
sm_80 → 真实硬件目标可以表示为：
CUDA C++
↓ 以 compute_80 为前端功能目标
PTX
↓ 以 sm_80 为后端机器目标
A100 类 GPU 原生代码十六、为什么要分虚拟目标和真实目标
假设使用：
compute_80编译 CUDA C++。
这意味着编译器只能使用 Compute Capability 8.0 所保证的功能。
但这份 PTX 可以进一步生成：
sm_80 原生代码
sm_86 原生代码
sm_89 原生代码前提是目标架构都能够支持这份 PTX 所假设的功能。
因此：
虚拟架构决定“程序允许使用什么功能”
真实架构决定“最终为哪块硬件生成指令”十七、最简单的 nvcc 架构选项
例如：
nvcc -O3 -arch=sm_80 kernel.cu -o app在当前 NVCC 的简写规则下，指定真实架构 sm_80 时，会使用对应的虚拟架构进行前端编译，并生成：
sm_80 Cubin
+
对应的 PTX这种简写类似于：
nvcc \
--gpu-architecture=compute_80 \
--gpu-code=sm_80,compute_80 \
kernel.cu当前 NVCC 文档明确说明，-arch=sm_XX 简写会同时加入对应原生代码和最接近的虚拟 PTX。
十八、如何只生成 PTX
可以执行：
nvcc -O3 \
-arch=compute_80 \
-ptx kernel.cu \
-o kernel.ptx得到：
kernel.ptx然后可以直接用文本编辑器查看。
这很适合观察：
- 循环是否展开；
- 是否使用向量指令；
- 是否出现原子操作；
- 编译器是否生成分支；
- 使用了哪类 Load/Store；
- 是否生成矩阵指令。
但要再次强调：
PTX 不是最终硬件执行结果。
真正做性能分析时，还需要看 SASS。
十九、如何生成特定架构 Cubin
例如为 Ampere sm_80 生成：
nvcc -O3 \
-arch=sm_80 \
-cubin kernel.cu \
-o kernel_sm80.cubin为 Hopper 生成：
nvcc -O3 \
-arch=sm_90 \
-cubin kernel.cu \
-o kernel_sm90.cubin两份 Cubin 中的原生指令可能不同，即使源代码完全相同。
二十、什么是 Fat Binary
一个程序可能需要支持多种 GPU。
例如希望同时支持：
Ampere sm_80
Ada sm_89
Hopper sm_90
Blackwell sm_100可以把多个版本的 GPU 代码放进同一个 Fat Binary：
图示：二十、什么是 Fat Binary
Fat Binary 可以容纳多个 PTX 和 Cubin 版本，运行时 CUDA Driver 会选择最合适的版本。
二十一、如何编译多个架构版本
可以使用重复的 -gencode：
nvcc -O3 kernel.cu -o app \
-gencode arch=compute_80,code=sm_80 \
-gencode arch=compute_89,code=sm_89 \
-gencode arch=compute_90,code=sm_90 \
-gencode arch=compute_100,code=sm_100 \
-gencode arch=compute_100,code=compute_100这里生成：
sm_80 原生代码
sm_89 原生代码
sm_90 原生代码
sm_100 原生代码
compute_100 PTX 后备代码--generate-code 可以重复使用，为不同虚拟架构和真实架构分别生成代码。
二十二、运行时如何选择代码
当程序启动 Kernel 时，CUDA Driver 会查看 Fat Binary。
假设当前 GPU 是 sm_90。
驱动可能按以下逻辑选择：
图示：二十二、运行时如何选择代码
直接加载 Cubin 通常启动更快。
通过 PTX JIT 则可以支持编译时尚未出现的新 GPU。
二十三、什么是 JIT 编译
JIT 是：
Just-In-Time Compilation，即运行时即时编译。
如果程序中只有 PTX，没有当前 GPU 的原生 Cubin，驱动会执行：
PTX
↓
Driver JIT Compiler
↓
当前 GPU 的 SASS/Cubin
↓
执行例如程序发布时最新 GPU 是 sm_80，但保留了 compute_80 PTX。
几年后在更新的 GPU 上运行时，新的驱动可以把这份 PTX 转换成新 GPU 的原生代码。
PTX 相对于 Cubin 具有更强的前向兼容性。
二十四、JIT 编译有什么代价
JIT 的主要代价是首次运行时需要额外编译时间。
可能出现：
第一次启动 Kernel：
较慢
后续启动：
较快因为驱动通常可以缓存 JIT 生成的代码。
JIT 还存在一个重要条件：
当前 GPU Driver 必须能够理解程序中使用的 PTX ISA 版本。
如果使用非常新的 CUDA Toolkit 生成 PTX，但目标机器驱动过旧，驱动可能无法识别这份 PTX。
因此兼容性不仅与 GPU 架构有关，也与 Toolkit 和 Driver 版本有关。
二十五、Cubin 的兼容性
对于普通基线 Cubin，可以简化理解为：
相同 Major Compute Capability 内
较低 Minor 版本编译的 Cubin
可能在较高 Minor 版本 GPU 上运行例如官方兼容性指南给出的示例是：
图示：二十五、Cubin 的兼容性
而 PTX 的前向兼容范围更广，例如面向 9.x 基线生成的 PTX，可以在满足条件的更高 Compute Capability GPU 上由驱动重新编译。
二十六、为什么不只保留 PTX
如果 PTX 有前向兼容性，似乎可以完全不生成 Cubin。
但只使用 PTX 有几个问题：
- 首次运行有 JIT 开销；
- 依赖目标机器的 Driver JIT；
- 不同驱动生成的代码可能略有差异；
- 某些架构专用功能不能通过通用 PTX 使用；
- 部署行为不如预编译 Cubin 可预测。
因此常见做法是：
为当前主要 GPU 生成 Cubin
+
保留一个 PTX 版本用于未来兼容二十七、为什么不为所有 GPU 都生成 Cubin
每增加一个架构版本：
- 可执行文件体积增加；
- 编译时间增加；
- Device Link 时间增加；
- 发布包更大；
- 测试组合增多。
如果一个大型库包含上千个 Kernel，同时为很多架构生成原生代码，二进制体积可能非常大。
所以软件通常会在：
启动速度
文件大小
支持范围
架构优化程度之间权衡。
二十八、什么是 sm_90a 和 sm_100a
从 Hopper 开始，一些非常专用的指令不保证在未来架构中继续以相同形式存在。
例如某些：
- Tensor Core 指令；
- WGMMA；
- 特定 TMA 功能；
- 架构专用数据移动机制。
因此 CUDA 引入带 a 后缀的架构专用目标：
sm_90a
sm_100a
sm_120aa 可以理解为：
Architecture-Specific。
使用 sm_90a 编译时，可以使用 Hopper 完整的架构专用功能，但生成代码只保证在准确对应的 Compute Capability 上运行。
官方文档规定，架构专用目标只能在与其编译目标完全匹配的 Compute Capability 上执行。
二十九、为什么 sm_90 和 sm_90a 不相同
sm_90 使用基线功能：
尽量保持向后端未来架构兼容的公共功能sm_90a 可以额外使用 Hopper 专用能力：
更强功能
但兼容范围更窄可以表示为：
图示：二十九、为什么 sm_90 和 sm_90a 不相同
例如 CUTLASS 中某些 Hopper Tensor Core Kernel 需要以 90a 为目标编译；只指定普通 90 无法合法使用这些专用特性。
三十、什么是 f 后缀
从 Compute Capability 10.0 开始，还出现了：
compute_100f
sm_100ff 表示：
Family-Specific，架构家族专用功能。
它位于普通基线和完全架构专用之间：
普通目标：
compute_100
兼容范围最宽，功能最少
Family 目标：
compute_100f
可使用家族共有专用功能
Architecture 目标：
compute_100a
可使用具体架构完整专用功能官方文档以 Compute Capability 10.x 为例说明：
图示：三十、什么是 f 后缀
三十一、如何查看 GPU 的 Compute Capability
可以直接使用：
nvidia-smi --query-gpu=name,compute_cap --format=csv输出可能类似：
name, compute_cap
NVIDIA A100-SXM4-80GB, 8.0也可以在 CUDA 程序中读取：
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
printf("Compute Capability: %d.%d\n",
prop.major,
prop.minor);官方 CUDA Programming Guide 也提供了通过 nvidia-smi、Runtime API、Driver API 和 NVML 查询 Compute Capability 的方式。
三十二、__CUDA_ARCH__ 是什么
编译 Device Code 时，编译器会定义：
__CUDA_ARCH__例如：
图示：三十二、__CUDA_ARCH__ 是什么
可以进行条件编译：
__device__ float operation(float x)
{
#if __CUDA_ARCH__ >= 900
// 使用较新架构路径
return new_operation(x);
#else
return old_operation(x);
#endif
}需要注意：
__CUDA_ARCH__主要用于 Device Code，不应让普通 Host Code 依赖它。
当前 NVCC 文档定义该宏为虚拟架构编号，并说明它在不同 -gencode 前端编译阶段可能取不同值。
三十三、一个 PTX 指令是否对应一个 SASS 指令
不一定。
可能出现：
一对一
PTX add.f32
↓
SASS FADD一对多
一条复杂 PTX 指令可能被拆成多条 SASS。
多对一
例如：
乘法
+
加法可能被融合为：
FFMA完全消失
如果编译器发现结果没有使用，相关指令可能被优化删除。
被替换
例如除法、取模或常量乘法，可能被替换为：
- 乘倒数；
- 移位；
- 乘加；
- 位运算组合。
所以不能通过 PTX 指令数量直接精确推断最终周期数。
三十四、为什么 SASS 才能真正反映寄存器和指令
例如 CUDA 源代码：
float y = x * 2.0f;PTX 可能显示：
mul.f32但最终 SASS 可能：
- 使用
FADD x, x； - 使用
FMUL； - 与其他操作融合成
FFMA； - 在编译期直接常量折叠。
真正进入 Warp Scheduler 和执行管线的是 SASS。
因此性能分析的层次是：
CUDA C++：
理解算法
PTX：
理解编译器中间语义
SASS：
理解真实硬件指令三十五、如何查看可执行文件中的 PTX
假设已经编译：
nvcc -O3 -arch=sm_80 kernel.cu -o app可以执行：
cuobjdump --dump-ptx app或者缩写：
cuobjdump -ptx appcuobjdump 可以从可执行文件、目标文件、静态库或 Fat Binary 中提取和显示 PTX。
三十六、如何查看 SASS
执行：
cuobjdump --dump-sass app或者：
cuobjdump -sass app它会反汇编可执行文件中嵌入的 Cubin。
如果已经有单独 Cubin，还可以使用：
nvdisasm kernel_sm80.cubin两者区别可以简化为：
cuobjdump：
可以直接处理 Host 可执行文件、库和 Cubin
nvdisasm：
主要处理 Cubin，提供更丰富的反汇编分析功能这是 CUDA Binary Utilities 官方文档给出的主要差异。
三十七、如何同时查看源码行和 SASS
编译时加入行信息：
nvcc -O3 \
-lineinfo \
-arch=sm_80 \
-cubin kernel.cu \
-o kernel.cubin然后：
nvdisasm -g kernel.cubin或者：
nvdisasm --print-line-info kernel.cubin这样可以看到某些 SASS 指令对应哪一行 CUDA 源码。
nvdisasm -g 会利用 Cubin 中的调试行信息对反汇编结果进行标注。
三十八、如何查看寄存器使用量
编译时使用：
nvcc -O3 \
-Xptxas=-v \
-arch=sm_80 \
kernel.cu \
-o app可能输出：
ptxas info:
Used 40 registers,
0 bytes spill stores,
0 bytes spill loads,
8192 bytes smem重点关注：
Used registers
spill stores
spill loads
shared memory
constant memory也可以使用：
cuobjdump --dump-resource-usage app或：
cuobjdump -res-usage app该命令可以显示每个函数的寄存器、栈、Shared Memory 和 Local Memory 使用量。
三十九、什么是 Register Spill
如果一个线程需要的活跃变量太多，物理寄存器不足，编译器可能把部分数据放入 Local Memory。
变量太多
↓
寄存器无法容纳
↓
Spill 到 Local MemorySASS 中可能出现额外的 Local Memory Load/Store。
需要注意：
Local Memory虽然名字中有“Local”，但物理上通常位于片外显存地址空间，并经过缓存。
因此 Spill 可能显著增加访问延迟和内存流量。
查看 ptxas 报告和 SASS，是判断 Spill 的重要方法。
四十、为什么同一源代码在不同架构上寄存器数量不同
因为不同架构的：
- 指令集不同；
- 地址指令不同；
- Tensor Core 接口不同；
- 调度规则不同；
- 编译后端不同；
- 寄存器分配策略不同。
例如同一 Kernel：
sm_80：使用 48 个寄存器
sm_90：使用 44 个寄存器
sm_100：使用 52 个寄存器是完全可能的。
这也会进一步影响：
- Occupancy；
- Spill；
- 每个 SM 可驻留 Block 数；
- 最终性能。
因此每个目标架构都需要独立测试。
四十一、为什么手写 PTX 不一定更快
PTX 看起来已经接近汇编，因此有人可能认为手写 PTX 一定比 CUDA C++ 快。
实际上不一定。
原因包括：
- PTX 仍会被后端重新优化；
- PTX 不直接控制最终寄存器分配；
- PTX 不直接控制全部 SASS 调度；
- 编译器可能比人工更擅长常规优化；
- PTX ISA 比真实硬件抽象；
- 不同架构需要不同优化。
手写 PTX 更适合：
- 使用 CUDA C++ 暂未暴露的指令；
- 精确表达特殊语义；
- 编写底层库；
- 架构研究；
- 微基准测试。
对于普通 Kernel，应先写清晰 CUDA C++，再用 PTX 和 SASS 验证编译结果。
四十二、内联 PTX
CUDA C++ 中可以使用内联 PTX：
__device__ int lane_id()
{
int id;
asm("mov.u32 %0, %%laneid;"
: "=r"(id));
return id;
}其中：
mov.u32是 PTX 指令。
内联 PTX 可以访问某些特殊寄存器或指令，但需要注意：
- 约束类型必须正确；
- 编译器仍可能重排周围代码；
- 不同 Compute Capability 支持不同；
- 可移植性降低；
- 错误使用可能破坏优化。
四十三、SASS 能否直接手写
理论上可以研究二进制编码，但 NVIDIA 的常规开发流程并不以手写 SASS 为主要编程接口。
通常流程是：
CUDA C++
或
PTX
↓
官方编译工具
↓
SASS原因包括：
- SASS 与架构高度绑定；
- 指令编码和控制字段复杂；
- 不同架构变化明显；
- 缺少稳定的跨代保证；
- 手工处理寄存器和调度成本很高。
所以 PTX 是官方提供的更稳定低级编程接口。
四十四、PTX 版本和 Compute Capability 不是同一概念
PTX 文件开头可能出现：
.version 8.0
.target sm_80
.address_size 64这里包含两个不同概念。
.version
表示使用的 PTX ISA 语言版本。
例如：
PTX ISA 8.x
PTX ISA 9.x.target
表示这份 PTX 所假设的目标 Compute Capability 或功能目标。
例如：
sm_80
sm_90
sm_100所以：
PTX ISA 版本
≠
GPU Compute Capability一个 PTX ISA 版本可以描述多个 Compute Capability 的目标功能。
四十五、为什么新驱动能够支持旧程序
假设旧程序包含：
sm_80 Cubin
+
compute_80 PTX在未来 GPU 上：
- 原来的
sm_80Cubin 可能不能直接执行； - 但新的 Driver 可以读取
compute_80PTX； - 再生成未来 GPU 的原生指令。
这就是 PTX 前向兼容机制的价值。
反过来，旧驱动通常不能理解未来工具链产生的新 PTX 或新 GPU。
所以兼容性方向通常是：
新 Driver 支持较旧 CUDA 程序而不是：
旧 Driver 自动支持未来 GPU 和未来 PTX四十六、为什么库经常很大
例如深度学习库可能需要包含：
多种 GPU 架构
×
多种数据类型
×
多种矩阵尺寸
×
多种布局
×
多种算法即：
sm_80 FP16 GEMM
sm_80 BF16 GEMM
sm_90 FP8 GEMM
sm_100 FP4 GEMM
……每个组合都可能有多个预编译 Kernel。
所以 CUDA、PyTorch、cuDNN 或 CUTLASS 相关安装包可能非常大。
其中相当一部分空间用于存放不同 GPU 架构的原生代码和模板实例。
四十七、从 GPU 架构设计角度理解 ISA 分层
为什么 NVIDIA 不直接把 CUDA C++ 编译成一个永久固定的 SASS？
因为 GPU 架构持续变化：
执行单元数量变化
指令格式变化
Tensor Core 变化
Load/Store 机制变化
同步机制变化如果应用只分发固定机器码，那么新 GPU 很难复用旧程序。
通过：
CUDA C++
↓
PTX 虚拟 ISA
↓
目标相关 SASS可以把：
软件语义与：
具体微架构部分解耦。
PTX 提供了跨代稳定层，而 SASS 可以针对每代 GPU 进行深度优化。
四十八、从 RTL 角度看 SASS 指令执行
当 Cubin 被加载后，Kernel 的 SASS 指令存入 GPU 指令存储或缓存。
执行过程大致是：
图示：四十八、从 RTL 角度看 SASS 指令执行
例如：
FFMA R8, R4, R5, R8;解码后可能产生：
图示：四十八、从 RTL 角度看 SASS 指令执行
这说明 SASS 是连接编译器和真实硬件流水线的最终接口。
四十九、实际分析一个 Kernel 的推荐步骤
可以按以下层次进行。
第一步：看 CUDA C++
判断算法结构：
- 并行维度；
- 数据依赖；
- 访存模式；
- 同步；
- 分支。
第二步：看编译报告
nvcc -Xptxas=-v ...检查：
- 寄存器数；
- Shared Memory；
- Spill；
- Stack。
第三步：看 PTX
nvcc -ptx ...检查：
- 是否使用目标指令；
- 分支结构；
- 内存空间；
- 原子操作；
- 向量化。
第四步：看 SASS
cuobjdump -sass app检查：
- 真实 Load/Store；
- FFMA 数量；
- 分支；
- Spill；
- Tensor Core 指令；
- 指令调度。
第五步：使用性能分析器
结合 Nsight Compute 查看：
- 指令吞吐率；
- Stall Reason；
- Cache Hit；
- Memory Transactions；
- Occupancy；
- Eligible Warps。
五十、本课核心结论
第一，CUDA .cu 文件中同时包含 CPU Host Code 和 GPU Device Code。
第二，nvcc 是编译驱动，负责协调 Host 编译器、CUDA 前端、ptxas 和链接工具。
第三，PTX 是 NVIDIA 的虚拟并行指令集，不是某块具体 GPU 的最终机器码。
第四，SASS 是针对特定 GPU 架构的原生机器指令表示。
第五，ptxas 将 PTX 优化并转换成特定架构的 SASS。
第六，Cubin 是包含 GPU 原生代码、符号和资源信息的二进制文件。
第七，Compute Capability 描述 GPU 的功能和资源版本，不是单纯的性能评分。
第八：
compute_80 = 虚拟 PTX 架构
sm_80 = 真实机器架构第九，Fat Binary 可以同时包含多个架构的 Cubin 和 PTX。
第十，运行时 Driver 优先选择兼容 Cubin；没有时可以对 PTX 执行 JIT 编译。
第十一，Cubin 启动更直接，PTX 提供更强的前向兼容能力。
第十二，sm_90a、sm_100a 等目标允许使用架构专用功能，但兼容范围较窄。
第十三，PTX 虚拟寄存器数量不等于最终物理寄存器数量。
第十四，真正决定执行管线、寄存器分配和指令调度的是最终 SASS。
第十五，可以使用：
cuobjdump -ptx
cuobjdump -sass
cuobjdump -res-usage
nvdisasm查看 GPU 二进制内容。
下一课将继续进入编译器与性能之间的关系：
第十九课：编译器如何优化 CUDA Kernel
主要包括：
常量传播
死代码消除
循环展开
指令融合
FMA
地址计算优化
寄存器分配
函数内联
分支转谓词
向量化 Load/Store
编译器优化为什么有时反而降低性能
volatile、restrict 和内联 PTX 的作用