<<<CUDA C++ 学习路线grid · block · warp · lane
EN中文
阶段 4 · 深入底层12 / 15约 22 分钟

PTX 与 SASS：看到编译器到底生成了什么

nvcc 编译流水线、虚拟架构与真实架构、以及怎么读汇编

学完这一课你会
  • 看懂 nvcc 的两阶段编译模型和 fatbin 的结构
  • 分清 compute_XX 与 sm_XX，知道 JIT 什么时候发生
  • 会用 cuobjdump / nvdisasm 导出并阅读 SASS
  • 了解内联 PTX 的用法及其边界

当你想知道「为什么这个 kernel 比预期慢」时，源码经常已经给不出答案。真正的答案在编译产物里：编译器有没有展开循环、有没有用上 FMA、寄存器分配得怎么样、有没有偷偷插入溢出指令。要看到这些，你得下到 PTX 和 SASS 这两层。

      kernel.cu
          │
    ┌─────┴─────┐  nvcc 前端拆分
    │           │
 host 代码    device 代码
    │           │
  g++/cl      cicc（NVVM / LLVM 后端）
    │           │
    │        kernel.ptx      ← 虚拟 ISA，与具体 GPU 无关，向前兼容
    │           │
    │        ptxas           ← 汇编器：寄存器分配、指令调度
    │           │
    │        kernel.cubin    ← SASS，特定架构的真实机器码
    │           │
    │      fatbinary         ← 可同时打包多个架构的 PTX + cubin
    └─────┬─────┘
          │
      可执行文件

  运行时：驱动先找匹配当前 GPU 的 cubin；
          找不到就拿 PTX 现场 JIT（首次启动会明显变慢）
nvcc 的编译流水线

compute_XX 与 sm_XX 的区别

写法含义产物
-arch=compute_80虚拟架构：PTX 的目标PTX（可 JIT 到更新的卡）
-code=sm_80真实架构：SASS 的目标cubin（只能在 sm_80 上跑）
-arch=sm_80上面两者的简写PTX + cubin 都是 80
-gencode arch=compute_80,code=sm_80精确控制单个目标sm_80 的 cubin
为多张卡构建一个通用二进制
1nvcc -O3 \2  -gencode arch=compute_75,code=sm_75 \   # Turing  T4 / RTX 20xx3  -gencode arch=compute_80,code=sm_80 \   # Ampere  A1004  -gencode arch=compute_86,code=sm_86 \   # Ampere  RTX 30xx5  -gencode arch=compute_90,code=sm_90 \   # Hopper  H1006  -gencode arch=compute_90,code=compute_90 \  # 保留 PTX，未来新卡可 JIT7  -o app main.cu8 9# 最后一行很关键：留一份 PTX 作为「向前兼容的后备」，10# 否则在 sm_100 之类的新卡上会直接报 no kernel image is available。

读 PTX

导出 PTX
1nvcc -ptx -arch=sm_80 -o kernel.ptx kernel.cu2# 或者从已编译的二进制里抽3cuobjdump -ptx ./app
vecAdd 的 PTX（节选）
1.visible .entry _Z6vecAddPKfS0_Pfi(2    .param .u64 a, .param .u64 b, .param .u64 c, .param .u32 n3)4{5    .reg .pred  %p<2>;6    .reg .f32   %f<4>;7    .reg .b32   %r<6>;8    .reg .b64   %rd<11>;9 10    ld.param.u32 %r2, [n];11    mov.u32      %r3, %ctaid.x;        // blockIdx.x12    mov.u32      %r4, %ntid.x;         // blockDim.x13    mov.u32      %r5, %tid.x;          // threadIdx.x14    mad.lo.s32   %r1, %r3, %r4, %r5;   // i = blockIdx*blockDim + threadIdx15    setp.ge.s32  %p1, %r1, %r2;        // p1 = (i >= n)16    @%p1 bra     END;                  // 谓词跳转：越界就直接退出17 18    ld.global.f32 %f1, [%rd5];19    ld.global.f32 %f2, [%rd8];20    add.f32       %f3, %f2, %f1;21    st.global.f32 [%rd10], %f3;22END:23    ret;24}

读 SASS：真实执行的指令

PTX 还是中间表示，ptxas 会对它做大量优化后才生成 SASS。SASS 才是硬件真正执行的东西，寄存器数、指令调度、双发射都在这一层确定。想知道编译器最终做了什么，只能看 SASS。

导出并阅读 SASS
1# 从可执行文件反汇编2cuobjdump -sass ./app | less3 4# 只看某个 kernel5cuobjdump -sass ./app | awk '/vecAdd/,/^$/'6 7# 从 cubin 反汇编，带控制信息8nvcc -cubin -arch=sm_80 -o kernel.cubin kernel.cu9nvdisasm -c kernel.cubin10 11# 同时打印寄存器用量和溢出情况（最常用的一条）12nvcc -O3 -arch=sm_80 -Xptxas -v -c kernel.cu
SASS 长这样（Ampere）
1        /*0000*/  MOV R1, c[0x0][0x28] ;2        /*0010*/  S2R R0, SR_CTAID.X ;              // blockIdx.x3        /*0020*/  S2R R3, SR_TID.X ;                // threadIdx.x4        /*0030*/  IMAD R0, R0, c[0x0][0x0], R3 ;    // i = bid*bdim + tid5        /*0040*/  ISETP.GE.AND P0, PT, R0, c[0x0][0x178], PT ;6        /*0050*/  @P0 EXIT ;                        // 越界线程直接退出7        /*0060*/  LDG.E R4, [R2.64] ;               // 全局加载8        /*0070*/  LDG.E R5, [R6.64] ;9        /*0080*/  FADD R7, R4, R5 ;10        /*0090*/  STG.E [R8.64], R7 ;               // 全局存储11        /*00a0*/  EXIT ;
SASS 指令含义看到它意味着
LDG.E.128128 位向量化全局加载好：向量化生效了
LDG.E32 位全局加载可以考虑改用 float4 向量化
LDS / STS共享内存读写正常
LDL / STL本地内存读写警报：寄存器溢出了
HMMA / IMMATensor Core 矩阵乘累加好：Tensor Core 用上了
FFMA融合乘加好：一条指令完成 a*b+c
BSSY / BSYNC分支同步存在 warp divergence

内联 PTX

极少数情况下，你需要一条 C++ 层面表达不出来的指令。CUDA 允许在 kernel 里内联 PTX 汇编，语法和 GCC 的扩展汇编一致。

内联 PTX 的典型用法
1// 1) 非临时加载：绕过 L1，避免污染缓存（流式数据）2__device__ __forceinline__ float ldNonTemporal(const float* p) {3    float v;4    asm volatile("ld.global.nc.f32 %0, [%1];" : "=f"(v) : "l"(p));5    return v;6}7 8// 2) 读取硬件时钟计数器，做细粒度计时9__device__ __forceinline__ unsigned long long clock64_() {10    unsigned long long t;11    asm volatile("mov.u64 %0, %%clock64;" : "=l"(t));12    return t;13}14 15// 3) 提示 L2 缓存驻留策略（Ampere+）16__device__ __forceinline__ void prefetchL2(const void* p) {17    asm volatile("prefetch.global.L2 [%0];" :: "l"(p));18}

自测

先自己在心里回答一遍，再展开对照。答不上来的说明这一段值得重读。