gpu compiler
**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).
- Fast math (`-use_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.