<<<CUDA C++ Coursegrid · block · warp · lane
EN中文
Stage 4 · Going lower12 / 15about 22 min

PTX and SASS: see what the compiler actually emitted

The nvcc pipeline, virtual vs real architecture, and how to read the assembly

After this lesson you can
  • Read nvcc's two-stage compile model and the structure of a fatbin
  • Tell compute_XX from sm_XX, and know when JIT happens
  • Dump and read SASS with cuobjdump / nvdisasm
  • Know when inline PTX is useful and where it stops helping

When you want to know why a kernel is slower than you expected, source often cannot answer. The answer is in the compile artifact: whether the compiler unrolled the loop, whether it emitted FMA, how it allocated registers, whether it inserted spill instructions. To see that, you drop to PTX and SASS.

      kernel.cu
          │
    ┌─────┴─────┐  nvcc frontend splits here
    │           │
 host code    device code
    │           │
  g++/cl      cicc (NVVM / LLVM backend)
    │           │
    │        kernel.ptx      ← virtual ISA, GPU-agnostic, forward compatible
    │           │
    │        ptxas           ← assembler: register allocation, instruction scheduling
    │           │
    │        kernel.cubin    ← SASS, real machine code for one architecture
    │           │
    │      fatbinary         ← can pack PTX + cubin for several architectures
    └─────┬─────┘
          │
      executable

  At runtime the driver looks for a cubin matching this GPU;
  if none, it JITs from PTX (first launch is noticeably slower)
nvcc compile pipeline

compute_XX vs sm_XX

FlagMeaningOutput
-arch=compute_80Virtual architecture: PTX targetPTX (can JIT to newer GPUs)
-code=sm_80Real architecture: SASS targetcubin (runs only on sm_80)
-arch=sm_80Shorthand for both of the abovePTX + cubin, both 80
-gencode arch=compute_80,code=sm_80Precise control of one targetsm_80 cubin
Build one fat binary for several GPUs
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 \  # keep PTX so future GPUs can JIT7  -o app main.cu8 9# The last line matters: keep a PTX copy as a forward-compatible fallback,10# otherwise a newer GPU such as sm_100 reports no kernel image is available.

Reading PTX

Dump PTX
1nvcc -ptx -arch=sm_80 -o kernel.ptx kernel.cu2# or extract it from an already-built binary3cuobjdump -ptx ./app
PTX for vecAdd (excerpt)
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;                  // predicated branch: out of range, exit immediately17 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}

Reading SASS: the instructions that actually run

PTX is still an IR. ptxas applies a lot of optimization before it emits SASS. SASS is what the hardware actually runs. Register count, instruction scheduling, and dual issue are all decided at this layer. To know what the compiler finally did, you have to read SASS.

Dump and read SASS
1# disassemble the executable2cuobjdump -sass ./app | less3 4# one kernel only5cuobjdump -sass ./app | awk '/vecAdd/,/^$/'6 7# disassemble a cubin, with control info8nvcc -cubin -arch=sm_80 -o kernel.cubin kernel.cu9nvdisasm -c kernel.cubin10 11# also print register usage and spills (the command you will use most)12nvcc -O3 -arch=sm_80 -Xptxas -v -c kernel.cu
SASS looks like this (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 ;                        // out-of-range threads exit immediately7        /*0060*/  LDG.E R4, [R2.64] ;               // global load8        /*0070*/  LDG.E R5, [R6.64] ;9        /*0080*/  FADD R7, R4, R5 ;10        /*0090*/  STG.E [R8.64], R7 ;               // global store11        /*00a0*/  EXIT ;
SASS instructionMeaningWhat it tells you
LDG.E.128128-bit vectorized global loadGood: vectorization took effect
LDG.E32-bit global loadConsider switching to float4 vectorization
LDS / STSShared memory read/writeExpected
LDL / STLLocal memory read/writeAlarm: registers spilled
HMMA / IMMATensor Core matrix multiply-accumulateGood: Tensor Core is in use
FFMAFused multiply-addGood: one instruction for a*b+c
BSSY / BSYNCBranch synchronizationWarp divergence is present

Inline PTX

In a few cases you need an instruction C++ cannot express. CUDA lets you inline PTX assembly inside a kernel. The syntax matches GCC extended asm.

Typical uses of inline PTX
1// 1) Non-temporal load: skip L1, avoid polluting the cache (streaming data)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) Read the hardware clock counter for fine-grained timing9__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) Hint L2 cache residency policy (Ampere+)16__device__ __forceinline__ void prefetchL2(const void* p) {17    asm volatile("prefetch.global.L2 [%0];" :: "l"(p));18}

Self-check

Answer in your head first, then open the solution. If you cannot, reread that section.