GPU Compiler Pipeline and PTX is the compilation infrastructure that transforms CUDA C++ source code through multiple intermediate representations into machine code optimized for a specific GPU microarchitecture — a multi-stage process that performs aggressive optimization (instruction selection, register allocation, instruction scheduling, memory access optimization) to achieve near-peak hardware performance. Understanding the GPU compiler pipeline helps performance engineers write kernels that the compiler can optimize effectively and debug performance issues when automatic optimization falls short.
CUDA Compilation Pipeline
``
CUDA C++ Source (.cu)
↓
[NVCC Frontend]
↓ (splits host and device code)
Host C++ → [GCC/Clang] → Host binary
Device code → [NVVM IR] (LLVM-based)
↓
[PTX Code Generator] → PTX (Parallel Thread Execution) assembly
↓
[PTX Assembler (ptxas)] → SASS (native GPU machine code)
↓
[Linked] → Executable with embedded GPU binary
PTX (Parallel Thread Execution) — The GPU IR
- PTX is NVIDIA's virtual ISA — architecture-independent intermediate assembly.
- Like Java bytecode for GPUs: PTX compiled once → can be JIT-compiled to any SM architecture at runtime.
- PTX advantages:
- Forward compatibility: PTX from CUDA 9 still runs on new GPUs (JIT-recompiled).
- Portable: Target different GPU generations without recompiling source.
- PTX registers: Virtual (unlimited) → ptxas allocates physical registers.
PTX Example
`ptx`
.kernel vector_add (.param .u64 A, .param .u64 B, .param .u64 C)
{
.reg .u32 %r<4>;
.reg .f32 %f<3>;
.reg .u64 %rd<4>;
ld.param.u64 %rd0, [A];
cvta.to.global.u64 %rd0, %rd0;
mov.u32 %r0, %tid.x; // thread index
ld.global.f32 %f0, [%rd0+%r0*4]; // load A[i]
// ...
st.global.f32 [%rd2+%r0*4], %f2; // store C[i]
}
SASS (Streaming Assembler) — Native GPU ISA
- Architecture-specific machine code (SM80 for A100, SM90 for H100).
- Not publicly documented by NVIDIA (reverse-engineered by community).
- cuobjdump -sass kernel.cubin: Disassemble SASS from compiled kernel.
- SASS reveals: Actual instructions, register usage, memory access patterns, predication.
Key Compiler Optimizations
1. Instruction Selection
- Map CUDA math to optimal GPU instructions.
- __fmaf_rn(a,b,c) → FMAD instruction (fused multiply-add in one instruction → no rounding between multiply and add).-use_fast_math
- Fast math (): Replace division/sqrt with approximate hardware instructions → 2–5× faster, slightly less accurate.
2. Register Allocation
- Minimize register spills (to local memory) → high register pressure → expensive.
- ptxas: Limits max registers per thread (--maxrregcount=64) → trade register pressure for higher occupancy.
- Tradeoff: Fewer registers → more threads can run → better latency hiding vs. more registers → faster per-thread computation.
3. Instruction Scheduling
- Reorder instructions to hide memory latency → issue independent instructions while waiting for load.
- Dual-issue: H100 can issue 2 independent instructions simultaneously if no data dependency.
4. Memory Access Coalescing
- Compiler analyzes access patterns → generates coalesced ld.global instructions where possible.
- Shared memory bank conflict detection: Some compilers warn about bank conflicts.
5. Loop Unrolling
- #pragma unroll N`: Unroll inner loop N times → reduce loop overhead, enable instruction-level parallelism.
- Caveat: Too much unrolling → register pressure → spills → performance regression.
Compilation Flags
| Flag | Effect |
|------|--------|
| -O3 | Maximum optimization |
| --use_fast_math | Approximate math (FMAD, fast sqrt) |
| -arch=sm_90 | Target H100 architecture |
| --maxrregcount=64 | Limit registers (increase occupancy) |
| -lineinfo | Keep source line info for profiling |
| -Xptxas -v | Verbose register/shared memory usage report |
The GPU compiler pipeline is the invisible performance engineer inside every CUDA program — by transforming high-level C++ tensor operations into optimally scheduled, register-allocated, memory-coalesced machine instructions through a multi-stage compilation process, NVCC and ptxas routinely achieve 70–90% of theoretical GPU peak performance for well-structured kernels, making the compiler as important as the hardware architecture in determining whether a GPU workload achieves its potential throughput.