Compilation and machine code

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, 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 it at runtime for whatever GPU is actually present, including architectures that did not exist when the application shipped.

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, 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.

.visible .entry _Z6vecAddPfS_S_j(
.param .u64 _Z6vecAddPfS_S_j_param_0,
// ...
.param .u32 _Z6vecAddPfS_S_j_param_3
)
{
// ...
mov.u32 %r3, %tid.x;
mov.u32 %r4, %ntid.x;
mov.u32 %r5, %ctaid.x;
mad.lo.s32 %r1, %r5, %r4, %r3;
setp.ge.u32 %p1, %r1, %r2;
@%p1 bra $L__BB0_2;

// ...
ld.global.f32 %f1, [%rd8];
ld.global.f32 %f2, [%rd6];
add.f32 %f3, %f2, %f1;
// ...
st.global.f32 [%rd10], %f3;

$L__BB0_2:
ret;
}

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.

nvcc x.cu --gpu-architecture=compute_80 --gpu-code=compute_80,sm_86,sm_89