PTX and SASS: Intermediate Representations and Tuning Methods for GPU Kernels

Note: This is a review article (no first-hand experimental data), based on public compiler-chain documentation and architecture analysis.

Experiment Objective

Ch08's CUDA compilation pipeline article covered how nvcc compiles .cu into executables; this article goes one layer deeper: what the compilation intermediates actually look like, and when an engineer needs to "read the assembly" to tune. Goal: understand the division of labor and generation paths of PTX and SASS, how to disassemble kernels with cuobjdump/nvdisasm, and the three classes of "only visible in assembly" optimizations (instruction selection, register pressure, tail effects) — plus when to look and when not to waste time.

Code Walkthrough

Part 1: The compile chain from .cu to SASS

# Normal compile: nvcc generates a fatbin (PTX plus per-arch SASS)
nvcc -arch=sm_90 -O3 -o kernel kernel.cu

# PTX only (virtual IR, cross-arch)
nvcc -arch=compute_90 -ptx -o kernel.ptx kernel.cu

# Disassemble SASS (real machine code, arch-bound)
cuobjdump -sass kernel > kernel.sass

# Or: nvdisasm directly on the cubin
nvdisasm -c kernel.cubin > kernel.sass

Key points:

1. PTX is a "virtual ISA," SASS is a "real-machine ISA": PTX targets compute_XX (virtual arch), is cross-generation, and is the foundation of JIT and driver compatibility; SASS targets sm_XX (real arch) and is the machine code that actually executes. With -arch=sm_90, nvcc embeds both PTX and SASS.
2. cuobjdump -sass is the entry to real execution code: register allocation, instruction scheduling, and bank-conflict elimination done by the optimizer are visible at the SASS layer — PTX cannot show them.
3. Disassembly is "reading," not "editing": engineers do not hand-write SASS (except extreme FFMA-merging for specific archs); the value is confirming whether the compiler did the expected optimization and finding what it failed to do.

Part 2: Four signal classes worth reading in SASS (simplified example)

/* Signal 1: is FFMA fused? — "FFMA R6, R4, R5, R6" does multiply-add in one instruction */
/* Signal 2: register pressure — the more registers a warp uses, the lower the occupancy */
/* Signal 3: LDL/STL memory instructions — vectorized? (LDG.E.128 = 16B at once) */
/* Signal 4: bank conflicts — shared-memory bank distribution at the SASS level */

Reading method:

• Fused multiply-add (FFMA): a*b+c should be a single FFMA in SASS. If the compiler emitted FMUL + FADD, fusion was missed and the code style may block it (e.g., intermediate-variable phrasing).
• Memory vectorization: contiguous float4 loads should appear as LDG.E.128; if they degrade to four LDG.E.32, memory efficiency drops 4× — fixable by aligning the code.
• Register pressure: the register-number spans in SASS show how many registers each thread uses; above a threshold (e.g., 128), occupancy drops — matching the low Occupancy seen in Nsight.

Principles

Why two layers of intermediate representation

• PTX (virtual ISA) exists for compatibility: the same PTX can JIT into new-arch SASS, so old programs run on new cards without recompiling — the carrier of CUDA's "forward compatibility" promise.
• SASS (real ISA) exists for performance: each sm_XX architecture differs in instruction set, register file, and memory hierarchy; SASS is the arch-specialized optimal code.
• The division of labor: the front end generates PTX (language-level optimization); the back end turns PTX into SASS (arch-level optimization: register allocation, instruction scheduling). For performance questions, read the back-end artifact SASS.

Three "only visible in assembly" optimization classes

1. Instruction-selection problems: the compiler emitted inefficient sequences (unfused multiply-add, redundant mov, unvectorized memory access). The fix is usually in source — changing phrasing (float4 structs, #pragma unroll, removing intermediate variables) lets the compiler "see" the fusion opportunity.
2. Register-pressure problems: too many registers per thread → low occupancy → latency unhidden. Confirm in SASS, then use __launch_bounds__ to cap registers or refactor the kernel.
3. Tail effects and branch divergence: BRA and predicated instructions in SASS reveal intra-warp divergence. Divergence serializes warp branches — a typical source of "full utilization but low efficiency."

When to read assembly (judgment rules)

Scenario	Worth reading SASS?
kernel already memory/compute-bound (Nsight verdict)	yes: confirm vectorization and fusion
low occupancy but code looks normal	yes: confirm register pressure
ordinary business code, non-hot kernels	no: compiler defaults are good enough
the last 20% of performance tuning	yes: this is a finishing tool, not a starting point

Key insight: SASS tuning is "evidence confirmation," not "guesswork optimization" — first use Nsight (previous article) to classify the bottleneck, then read SASS to confirm the mechanism, then fix the source. Skipping classification to dig into assembly is needle-in-a-haystack.

Results

(Estimated, not measured)

Signal	Problem	Fix	Typical gain
unfused FMUL+FADD	multiply-add not fused	rephrase so compiler fuses	1.1-1.3×
LDG.E.32 ×4	memory not vectorized	float4 alignment	1.5-3×
>128 registers/thread	occupancy crushed	__launch_bounds__	1.2-1.5×
obvious branch divergence	warp serialization	reorder data to remove divergence	1.3-2×

Engineering Takeaways

1. SASS is a confirmation tool, not an optimization starting point: Nsight to classify (previous article), SASS to confirm (this one), source to fix — the reverse order is twice the work for half the result.
2. Fix code, not assembly: 99% of the time, change the source so the compiler emits better SASS; hand-editing SASS only matters in extreme finishing scenarios.
3. Archive disassembly results: keep SASS and ncu reports of performance-sensitive kernels as baselines; diff after compiler/driver upgrades to catch silent regressions.
4. Compiler versions change SASS: the same source can differ 10%-20% in SASS across CUDA versions — regression comparison after CUDA upgrades is a discipline.

Next article preview: with the compiler and kernel chain covered, back to the framework layer — "DeepSpeed ZeRO: Three-Stage Training-Memory Optimization, Principles and Practice."