The ISA contract

Compile four multiplies and four adds on an ordinary laptop, disassemble the object file, and every line comes back in two columns. On the right, addss %xmm0, %xmm2. On the left, the four bytes f3 0f 58 d0. They are not two things that correspond. They are one instruction written down twice, once for a person and once for the decoder, and an assembler turns either into the other without losing anything.

The left column is machine code, the encoding the hardware actually eats. The right column is assembly, a notation for humans with mnemonics instead of opcodes and names instead of register numbers. Because the two map one to one, people say "writing assembly" and "the machine's instructions" as if they were the same claim, and mostly nothing goes wrong. The distinction that does matter sits one level up from this pair.

0000000000000000 <_dot4>:
       0: 55                           	pushq	%rbp
       1: 48 89 e5                     	movq	%rsp, %rbp
       4: f3 0f 10 07                  	movss	(%rdi), %xmm0
       8: f3 0f 10 4f 04               	movss	0x4(%rdi), %xmm1
       d: f3 0f 59 06                  	mulss	(%rsi), %xmm0
      11: 0f 57 d2                     	xorps	%xmm2, %xmm2
      14: f3 0f 58 d0                  	addss	%xmm0, %xmm2

Those four bytes have meant the same thing since 1999, and they mean it on parts from two different companies. The agreement covers which instructions exist, how they are encoded, which registers they may name, what the memory model guarantees about the order other cores see writes in. All of that together is the instruction set architecture, and it is a contract: code compiled against it keeps running on machines nobody had designed when it was compiled.

What the contract deliberately leaves out is everything about how the work gets done. How many multiplies issue per cycle, how deep the pipeline runs, whether there is a reorder buffer at all, how large the caches are, how the branch predictor is organized. That is the microarchitecture, one implementation of the contract, and it is redesigned every generation or two. The dump above was produced on a Core i9-9880H; the same bytes run unchanged on an AMD part that shares none of its internals.

The ISA is the part that may not change. Everything under it is free to.

On a CPU the contract and the machine sit at the same address, so the two words collapse in casual speech and no harm comes of it. NVIDIA moved the line. There are two instruction sets in a CUDA binary, they are not versions of each other, and only one of them is the contract.

The first instruction set is PTX, and the specification opens by saying exactly what it is: "a low-level parallel thread execution virtual machine and instruction set architecture (ISA)". Virtual machine, meaning no chip decodes it. PTX has as many registers as a kernel wants, one type per instruction, and special registers with names like %tid.x rather than an address. It is an ISA in the sense of being a complete, documented, versioned contract, and in no other sense.

What the contract is for shows up in the goals list, one line of which reads "Provide a stable ISA that spans multiple GPU generations." A virtual ISA is a contract with the implementation deliberately missing: something for a compiler to target and something for a translator to consume, with the translation postponed. The specification says when the postponement ends, too: "PTX programs are translated at install time to the target hardware instruction set."

Read the dump below and the virtual part is visible in the register names. Every value gets a fresh one, %f1 through %f4 for floats and %rd1 through %rd7 for addresses, because nothing here is competing for a physical register file yet. The thread index arrives through %ctaid.x, %ntid.x, and %tid.x; the bounds check becomes a setp.ge.s32 writing a predicate and a @%p1 bra reading it; the multiply-add is one fma.rn.f32 with its rounding mode spelled out.

.visible .entry saxpy(
	.param .f32 saxpy_param_0,
	.param .u64 saxpy_param_1,
	.param .u64 saxpy_param_2,
	.param .u32 saxpy_param_3
)
{
	ld.param.f32 	%f1, [saxpy_param_0];
	ld.param.u64 	%rd1, [saxpy_param_1];
	ld.param.u64 	%rd2, [saxpy_param_2];
	ld.param.u32 	%r2, [saxpy_param_3];
	mov.u32 	%r3, %ctaid.x;
	mov.u32 	%r4, %ntid.x;
	mov.u32 	%r5, %tid.x;
	mad.lo.s32 	%r1, %r3, %r4, %r5;
	setp.ge.s32 	%p1, %r1, %r2;
	@%p1 bra 	$L__BB0_2;
	cvta.to.global.u64 	%rd3, %rd2;
	cvta.to.global.u64 	%rd4, %rd1;
	mul.wide.s32 	%rd5, %r1, 4;
	add.s64 	%rd6, %rd4, %rd5;
	ld.global.f32 	%f2, [%rd6];
	add.s64 	%rd7, %rd3, %rd5;
	ld.global.f32 	%f3, [%rd7];
	fma.rn.f32 	%f4, %f2, %f1, %f3;
	st.global.f32 	[%rd7], %f4;
$L__BB0_2:
	ret;
}

Push the same kernel one step further and the second instruction set appears. SASS is what the streaming multiprocessorsOne of a GPU’s many cores; each keeps many warps resident and switches among them to hide memory latency.taught in /l/tpu → decode, and everything postponed in PTX has been settled in it. Registers are physical and numbered, R0 through R7 here. The kernel parameters are no longer named; they are offsets into a constant bank, c[0x0][0x210] and c[0x0][0x218]. The bounds check is now ISETP.GE.AND writing the predicate register P0, and the threads that fail it leave immediately through a predicated @P0 EXIT.

One line is worth pausing on. The whole body of the PTX, four separate address computations and two loads, has collapsed into IMAD.WIDE pairs and LDG.E with a descriptor operand. Nothing in the source asked for that. It is the assembler choosing instructions for one particular generation of hardware, which is what an assembler for a real ISA does.

saxpy:
 LDC R1, c[0x0][0x28]
 S2R R0, SR_TID.X
 S2UR UR4, SR_CTAID.X
 LDC R7, c[0x0][RZ]
 IMAD R7, R7, UR4, R0
 ULDC UR4, c[0x0][0x228]
 ISETP.GE.AND P0, PT, R7, UR4, PT
 @P0 EXIT
 LDC.64 R2, c[0x0][0x218]
 ULDC.64 UR4, c[0x0][0x208]
 ULDC UR6, c[0x0][0x210]
 LDC.64 R4, c[0x0][0x220]
 IMAD.WIDE R2, R7, 0x4, R2
 LDG.E R2, desc[UR4][R2.64]
 IMAD.WIDE R4, R7, 0x4, R4
 LDG.E R7, desc[UR4][R4.64]
 FFMA R7, R2, UR6, R7
 STG.E desc[UR4][R4.64], R7
 EXIT

Now run the experiment that makes the split concrete. Take that source, one nvcc, and ask for two architectures. The PTX comes back byte for byte identical, the same thirty lines both times. The SASS does not, and the differences are not cosmetic: MOV R1 where Hopper used LDC R1, a bare LDG.E.SYS R2, [R2] where Hopper carried a descriptor, an IMAD.WIDE that folds the base address straight out of the constant bank, and every parameter at a different offset, 0x160 here against 0x210 there.

Nothing about the program changed between those two dumps. The contract held and the floor moved, which is the arrangement the goals list was describing. It also explains why NVIDIA behaves so differently about the two levels. The PTX specification is a document with a version number you can read cover to cover. For SASS the binary utilities manual gives you nvdisasm, which "extracts information from standalone cubin files and presents them in human readable format", an opcode table per architecture, and no assembler. SASS is something you read.

saxpy:
 MOV R1, c[0x0][0x28]
 @!PT SHFL.IDX PT, RZ, RZ, RZ, RZ
 S2R R4, SR_CTAID.X
 S2R R3, SR_TID.X
 IMAD R4, R4, c[0x0][0x0], R3
 ISETP.GE.AND P0, PT, R4, c[0x0][0x178], PT
 @P0 EXIT
 MOV R5, 0x4
 IMAD.WIDE R2, R4, R5, c[0x0][0x168]
 IMAD.WIDE R4, R4, R5, c[0x0][0x170]
 LDG.E.SYS R2, [R2]
 LDG.E.SYS R7, [R4]
 FFMA R7, R2, c[0x0][0x160], R7
 STG.E.SYS [R4], R7
 EXIT

Two tools cross the gap, and knowing which one ran explains most surprising launch times. At build time it is ptxas, which the nvcc manual calls "the PTX optimizing assembler". You tell nvcc which architectures to target, and it names them in two vocabularies: it produces "a true binary load image for each real architecture (such as sm_100), and PTX code for the virtual architecture (such as compute_100)". Both go into the same object file, so a shipped binary usually carries several cubins and one copy of the PTX they all came from.

The second tool is the driver, and it runs when the first list comes up short: "During runtime, such embedded PTX code is dynamically compiled by the CUDA runtime system if no binary load image is found for the current GPU." That is the mechanism behind a binary from two years ago starting on a card that did not exist when it was built. It is also the mechanism behind that binary taking an unreasonable amount of time on its first launch, because a compiler you never invoked is running inside the driver, on your critical path, over code you cannot inspect.

The wall lesson in the compiler unit already puts ptxas and SASS where the toolchain stops (/l/xla/dumps-on-demand, which names libtpu and LLOThe TPU’s near-assembly, closed inside libtpu. The readable world ends one layer above, at Mosaic.taught in /l/vliw-bundles-and-llo → in the same breath). This lesson says what is behind that door and why the vendor wants it shut: the floor is redesigned every generation, and a contract one level up is what lets them do it. The next question is not which instructions are down there. It is who decided which of them run at the same time.