# Compilation and machine code · Foundations A CUDA source file mixes host and device code, and nvcc, the CUDA compiler driver, splits it: the host code is forwarded to an ordinary C++ compiler, while the device functions go through NVIDIA’s own compilers. The device side compiles in two stages. First the CUDA C++ is compiled to **PTX** (Parallel Thread Execution, a virtual machine instruction set architecture that acts as the assembly language of the CUDA platform.), assembly for a virtual GPU defined only by its capabilities, not by any physical chip. Then an assembler called ptxas compiles the PTX into a *cubin*, an ELF binary holding the actual machine instructions for one real architecture. The compiled GPU code is embedded back into the host executable, which can carry both cubins and the PTX text itself. The split exists because NVIDIA does not guarantee binary compatibility across GPU generations: instruction sets and encodings change between major compute capabilities, so a cubin built for one generation will not run on the next. Every nvcc compilation therefore names two architectures, a virtual one (compute\_80) stating which features the code needs, and a real one (sm\_80) naming the processor to build for. Because the executable can also embed the PTX, the GPU driver can **JIT compile** (Compiling the embedded PTX to machine code at application runtime, done by a compiler built into the GPU driver, for the exact GPU the program is about to run on.) it at runtime for whatever GPU is actually present, including architectures that did not exist when the application shipped. > Figure. The two-stage device compile. nvcc compiles the device side of a .cu file to PTX, assembly for a virtual GPU, and ptxas compiles that into the machine instructions (SASS) of one real architecture. Because the executable can also carry the PTX, the driver can JIT compile it at runtime for the GPU actually present. Everything above the cubin is an intermediate form, which is why kernel authors read the machine code when it matters. PTX is what the compiler promised; the disassembled cubin, called **SASS** (The native assembly of a real GPU architecture, disassembled from a cubin by cuobjdump or nvdisasm.), is what the hardware will execute, and the two can differ because ptxas optimizes, allocates registers, and schedules instructions on the way down. The toolkit ships two disassemblers: cuobjdump accepts both cubins and host executables and can extract the embedded PTX, while nvdisasm accepts only cubins but adds control flow analysis and richer display options. When this book claims a kernel change did something, disassembling is how that claim gets checked: against the instructions the compiler actually emitted, not the ones the source seems to imply. What PTX looks like is worth seeing once. NVIDIA’s Understanding PTX post compiles the same bounds-checked vecAdd kernel from section 1.1 and prints the result. Most of it reads like assembly for any machine, loads, an add, a store, with the CUDA thread model showing through in special registers: `%tid.x`, `%ntid.x`, and `%ctaid.x` are `threadIdx.x`, `blockDim.x`, and `blockIdx.x`. > Figure. vecAdd compiled to PTX, from Understanding PTX (declarations and parameter loads elided) The kernel survives translation recognizably. A single `mad.lo.s32`, a multiply-add, computes the entire `workIndex` expression; `setp.ge.u32` and the predicated branch `@%p1 bra` are the bounds check, comparing the index against the length parameter (loaded into `%r2` in the elided lines) and jumping past the body when it is out of range; and the load, add, store triple in the middle is the arithmetic itself. Because PTX targets a virtual machine, the post notes, ptxas behaves more like a compiler than a traditional assembler when it lowers this into a cubin, which is where the register allocation and scheduling that separate PTX from SASS happen. The post draws the comparison this whole pipeline rests on: PTX plays the role in CUDA that LLVM IR plays for clang, a portable middle stage that many front ends can target and any supported GPU can be compiled from. Source Compilation trajectory and virtual versus real architectures from the [CUDA Compiler Driver NVCC](https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/) documentation; PTX’s role, the vecAdd PTX listing, and JIT compilation from [Understanding PTX](https://developer.nvidia.com/blog/understanding-ptx-the-assembly-language-of-cuda-gpu-computing/); cuobjdump and nvdisasm from [CUDA Binary Utilities](https://docs.nvidia.com/cuda/cuda-binary-utilities/).