GPU Glossary
/
Device Software·3 min read

Parallel Thread ExecutionPTX

A forward-compatible intermediate representation and virtual instruction set for NVIDIA GPUs, JIT-compiled by the driver into device-specific SASS.

Parallel Thread eXecution (PTX) is an intermediate representation (IR) for code that runs on a parallel processor, almost always an NVIDIA GPU. It is one of the formats output by nvcc, the NVIDIA CUDA Compiler Driver. Many NVIDIA engineers pronounce it "pee-tecks"; everyone else says "pee-tee-ecks".

Virtual machine, ISA and IR#

NVIDIA's documentation calls PTX both a "virtual machine" and an "instruction set architecture".

To a programmer, PTX is an instruction set for a virtual machine model. Programmers or compilers that emit PTX can be confident their program will have the same semantics on many different physical machines, including ones that don't exist yet. In that sense it resembles CPU ISAs such as x86_64, aarch64 or SPARC.

Unlike those ISAs, PTX is very much an intermediate representation, like LLVM-IR. The PTX parts of a CUDA binary are just-in-time (JIT) compiled by the host CUDA Drivers into device-specific SASS for execution.

Thanks to this JIT step, PTX is forward-compatible: GPUs with a matching or higher compute capability can run the program.

Key idea PTX is a "narrow waist" between hardware and software: compilers target PTX, and the driver translates it for whatever GPU is present.

Example PTX#

.reg .f32 %f<7>;

A directive to the PTX-to-SASS compiler saying this kernel uses seven 32-bit floating point registers. Registers are allocated dynamically to groups of threads (warps) from the SM's register file.

fma.rn.f32 %f5, %f4, %f3, 0f3FC00000;

A fused multiply-add (fma): multiply registers f3 and f4, add the constant 0f3FC00000, and store the result in f5. All values are 32-bit floating point. The rn suffix sets the rounding mode to IEEE 754 "round even", which is the default.

mov.u32 %r1, %ctaid.x;
mov.u32 %r2, %ntid.x;
mov.u32 %r3, %tid.x;

move the x-axis values of the cooperative thread array index, the cooperative thread array dimension index (ntid), and the thread index into three u32 registers, r1 to r3.

The PTX machine model#

PTX exposes several levels of parallelism to the programmer. These map directly onto hardware through the PTX machine model.

The PTX machine model. Modified from the PTX documentation.
The PTX machine model. Modified from the PTX documentation.

In this model, a single instruction unit serves multiple processors. Each processor runs one thread, but all of them must execute the same instructions: hence parallel thread execution. Threads coordinate through shared memory and produce different results through their private registers.

Versions and documentation#

The latest PTX documentation is available from NVIDIA. PTX instruction sets are versioned by "compute capability", which is synonymous with "minimum supported Streaming Multiprocessor architecture version".

Writing PTX by hand#

Hand-written in-line PTX is uncommon outside the cutting edge of performance, much like in-line x86_64 assembly (used, for example, in high-performance vectorized query operators in analytical databases and in performance-sensitive parts of operating system kernels).

As of September 2025, in-line PTX is the only way to use some Hopper-specific hardware features, such as the wgmma and tma instructions. Examples include Flash Attention 3 and the Machete w4a16 kernels.