PTX, SASS, and What a Fat Binary Contains

A device binary is not one artefact. The compilation path from a kernel source file to instructions the machine executes has two stages, and they can happen at different times on different machines. The stage boundary is a virtual instruction set, and whether a distributed binary contains the intermediate form, the final form, or both decides whether it still runs on hardware released after it was built.

Two compilers, one of them on the user’s machine

The offline stage runs the vendor compiler at build time. It parses the CUDA C++, performs the device-side optimisations, and emits PTX: a typed, register-based virtual assembly with an unlimited register namespace, explicit address spaces, and instructions such as fma.rn.f32, ld.global.nc.f32, and bar.sync 0. PTX is documented, versioned, and forward-compatible in the sense that a driver for newer hardware can consume PTX produced for an older ISA version.

The second stage is a compiler embedded in the driver. It translates PTX into SASS, the machine code of the specific architecture: the encoding that the instruction fetch unit decodes, with the register file size, the fixed number of warp schedulers, and the scoreboard hazards of that chip. This compilation happens at module load time, on the machine where the program runs, and its output is cached by the driver so the cost is paid once per binary per driver version.

The two stages explain the compatibility rules. SASS is tightly coupled to one architecture: an embedded opcode field, an encoding for the register bank, and a fixed pipeline schedule. It is not portable across major generations in every case, which is why a cubin compiled for one architecture may fail to load on another. PTX is portable upward, which is why shipping PTX alongside the SASS is the way to reach hardware that did not exist when the program was built.

What the fat binary is

nvcc produces a fat binary: a container holding one or more cubins, one or more PTX versions, or both, along with metadata. The container is embedded in the host object file, and the runtime linker reads it before the first kernel launch.

The contents are chosen by a pair of options repeated once per target: an architecture, which selects the instruction set the code is generated against, and a code target, which selects what is emitted. code=sm_XY embeds machine code for that architecture; code=compute_XY embeds PTX for that virtual architecture. Passing the same architecture with both code forms, once with sm and once with compute, produces the common pair: one native image and one intermediate fallback.

The loader’s selection order at module load is fixed. It prefers a cubin built for the device’s architecture or, absent an exact match, one built for an earlier minor revision of the same major architecture, which is binary compatible. Failing that, it takes the highest PTX version in the container that the driver supports and compiles it for the running device with the driver’s compiler. Failing both, the load fails.

// Query the architecture the runtime reports for the current device.
cudaDeviceProp prop;
cudaGetDeviceProperties(&prop, 0);
// prop.major and prop.minor are the compute capability,
// the same pair encoded in the sm_XY suffix of a cubin.
printf("%d.%d\n", prop.major, prop.minor);

The pair is a query, not a compile-time constant, and the string form is what a build system needs when it generates architecture flags.

Why the arch list is a build-time decision with runtime consequences

An application that ships SASS for one architecture runs at full speed there and pays PTX compilation everywhere else, if it shipped PTX at all. An application that ships only SASS for an architecture the user does not have fails to launch its kernels entirely, typically with an error naming the mismatch between the image and the device.

Ship PTX only, and every install compiles everything on first use, which shifts a cost onto the user that the build machine could have paid once, and makes the resulting code quality depend on the driver’s compiler rather than on the toolchain the application was built with. The driver’s PTX compiler is derived from the same code base as the offline compiler, but its version tracks the driver, so a new driver may produce different machine code from identical PTX.

The practical compromise, applied by most distributed binaries, is a small set of architecture-specific cubins covering the architectures in wide use plus one PTX image for the newest supported virtual architecture. The set is a maintenance burden that grows with each hardware generation, and it is the reason the PTX image stays in the container: PTX is the only artefact that covers hardware the build machine never saw.

SPIR-V is the same split with a different boundary

Vulkan and OpenCL 2.x use SPIR-V for the same purpose PTX serves: a documented intermediate representation that a driver lowers to its own machine code. The boundary is in a different place. In the Vulkan model, the intermediate form is the only artefact most applications ship, and the driver’s lowering is the normal path rather than a fallback. There is no equivalent of the embedded cubin because no vendor publishes native instructions to target.

The design decision is that the pipeline compilation is explicit: at pipeline creation, the driver compiles the SPIR-V module, and an implementation is permitted to cache the result on disk keyed by a hash of the module and the pipeline state. The cache is not transparent the way the driver’s PTX cache is; the application supplies the key and the storage.

The trade is visible in both directions. Vendors can recompile the same SPIR-V to exploit a new instruction without the application reissuing a binary, which is what makes a SPIR-V binary portable to hardware that did not exist when it was produced. The application loses any guarantee about how its code maps to instructions, and a shader that compiled well on one vendor may lower badly on another.

Reading the artefact as a portability statement

For any binary in this family, the useful question is what it contains and what the loader may do with it:

• A cubin names one architecture and encodes its instruction set exactly. Portable to that architecture and to newer ones only through a compatibility layer.
• A PTX image names a virtual architecture and a language version. Portable upward, compiled at load, with performance dependent on the driver’s compiler.
• A SPIR-V module names a version and a set of capability declarations that the device must satisfy. Portable across vendors subject to those capabilities, lowered per driver.

The version number in each case is a contract about the feature set, not a quality marker. PTX 7.x added features that older drivers cannot assemble, so a binary built with a newer toolkit requires a driver at or above the matching release. The failure mode is a load error or an unresolved-symbol style diagnostic at module load, well before any kernel runs, which is at least a clean failure.

What the layers above do with this

Portable single-source models sit on top of exactly this chain. A SYCL or OpenMP offload compilation produces SPIR-V for a generic target or native code for a specific one, selected by a flag. The mechanism is inherited wholesale, along with the trade: choose the generic target and pay per-device lowering; choose the native target and pay with a binary per architecture. The layering is described in Portable Single-Source Layers and Where They Leak, and the runtime side of the same choice, compiling or loading at startup and caching the result, is treated in Driver-Level Runtime Compilation and Module Loading.