2.2GPU Instruction Set Architectures

How does the CUDA kernel from §2.1 become instructions a GPU actually executes? The answer reveals something unusual about the GPU ecosystem: unlike x86 — backward compatible all the way to the Intel 8086 of 1976, so 40-year-old binaries can in principle still run — GPU vendors never promised hardware-level compatibility. What they promised instead shaped their entire compilation stack.

2.2.1 NVIDIA: a documented virtual ISA over an undocumented real one

With many vendors shipping many different GPU designs, early programmable GPUs already relied on ISA virtualization (OpenGL’s shading language, Microsoft’s HLSL). When NVIDIA introduced CUDA in 2007 they followed the same path: kernels compile first to PTXPTXNVIDIA's documented virtual RISC-like ISA (limitless virtual registers); the compiler target that gets lowered to hardware code.Glossary → — the Parallel Thread Execution virtual ISA — which is fully documented with every CUDA release (documented well enough that this book’s authors built the GPGPU-Sim simulator on it). PTX looks like a classic RISC ISA (ARM, MIPS, SPARC) crossed with a compiler’s intermediate representation: loads and stores, predicationpredicationExecuting an instruction conditionally under a predicate register (@%p1 in PTX, @P0 in SASS) instead of branching.Glossary → — and an essentially limitless supply of virtual registers, each definition typically getting a fresh one, much like static single assignment form.

Before running, PTX is lowered — by the driver or the standalone ptxas — to SASSSASSNVIDIA's hardware ISA ("Streaming ASSembler"); only partially documented, free to change every generation.Glossary → (“Streaming ASSembler”), the real hardware ISA. And SASS is only partially documented, quite deliberately: with no customer expectation of hardware compatibility, NVIDIA can redesign the ISA completely between generations. The community pushed back with reverse engineering — decuda arrived within about a year of CUDA-capable hardware (2007, for the GeForce 8), understood SASS well enough to build an assembler, and successors covered Fermi and Maxwell; NVIDIA eventually shipped cuobjdump and began listing opcode names (though, as of the book’s writing, not operand formats or semantics).

Step through the same SAXPY kernel at three levels — watch how the logical structure survives while everything else changes:

1.visible .entry _Z5saxpyifPfS_(2.param .u32 _Z5saxpyifPfS__param_0,3.param .f32 _Z5saxpyifPfS__param_1,4.param .u64 _Z5saxpyifPfS__param_2,5.param .u64 _Z5saxpyifPfS__param_36)7{8.reg .pred %p<2>;9.reg .f32 %f<5>;10.reg .b32 %r<6>;11.reg .b64 %rd<8>;12 13ld.param.u32 %r2, [_Z5saxpyifPfS__param_0];14ld.param.f32 %f1, [_Z5saxpyifPfS__param_1];15ld.param.u64 %rd1, [_Z5saxpyifPfS__param_2];16ld.param.u64 %rd2, [_Z5saxpyifPfS__param_3];17mov.u32 %r3, %ctaid.x;18mov.u32 %r4, %ntid.x;19mov.u32 %r5, %tid.x;20mad.lo.s32 %r1, %r4, %r3, %r5;21setp.ge.s32 %p1, %r1, %r2;22@%p1 bra BB0_2;23 24cvta.to.global.u64 %rd3, %rd2;25cvta.to.global.u64 %rd4, %rd1;26mul.wide.s32 %rd5, %r1, 4;27add.s64 %rd6, %rd4, %rd5;28ld.global.f32 %f2, [%rd6];29add.s64 %rd7, %rd3, %rd5;30ld.global.f32 %f3, [%rd7];31fma.rn.f32 %f4, %f2, %f1, %f3;32st.global.f32 [%rd7], %f4;33 34BB0_2:35ret;36}

1/*0000*/     MOV R1, c[0x1][0x100];2/*0008*/     S2R R0, SR_CTAID.X;3/*0010*/     S2R R2, SR_TID.X;4/*0018*/     IMAD R0, R0, c[0x0][0x8], R2;5/*0020*/     ISETP.GE.AND P0, PT, R0, c[0x0][0x20], PT;6/*0028*/ @P0 BRA.U 0x78;7/*0030*/ @!P0 MOV32I R5, 0x4;8/*0038*/ @!P0 IMAD R2.CC, R0, R5, c[0x0][0x28];9/*0040*/ @!P0 IMAD.HI.X R3, R0, R5, c[0x0][0x2c];10/*0048*/ @!P0 IMAD R4.CC, R0, R5, c[0x0][0x30];11/*0050*/ @!P0 LD.E R2, [R2];12/*0058*/ @!P0 IMAD.HI.X R5, R0, R5, c[0x0][0x34];13/*0060*/ @!P0 LD.E R0, [R4];14/*0068*/ @!P0 FFMA R0, R2, c[0x0][0x24], R0;15/*0070*/ @!P0 ST.E [R4], R0;16/*0078*/     EXIT;

1         /* control instruction */2/*0008*/     MOV R1, c[0x0][0x20];3/*0010*/     S2R R0, SR_CTAID.X;4/*0018*/     S2R R2, SR_TID.X;5         /* control instruction */6/*0028*/     XMAD.MRG R3, R0.reuse, c[0x0][0x8].H1, RZ;7/*0030*/     XMAD R2, R0.reuse, c[0x0][0x8], R2;8/*0038*/     XMAD.PSL.CBCC R0, R0.H1, R3.H1, R2;9         /* control instruction */10/*0048*/     ISETP.GE.AND P0, PT, R0, c[0x0][0x140], PT;11/*0050*/ @P0 EXIT;12/*0058*/     SHL R2, R0.reuse, 0x2;13         /* control instruction */14/*0068*/     SHR R0, R0, 0x1e;15/*0070*/     IADD R4.CC, R2.reuse, c[0x0][0x148];16/*0078*/     IADD.X R5, R0.reuse, c[0x0][0x14c];17         /* control instruction */18/*0088*/     IADD R2.CC, R2, c[0x0][0x150];19/*0090*/     IADD.X R3, R0, c[0x0][0x154];20/*0098*/     LDG.E R0, [R4];21         /* control instruction */22/*00a8*/     LDG.E R6, [R2];23/*00b0*/     FFMA R0, R0, c[0x0][0x144], R6;24/*00b8*/     STG.E [R2], R0;25         /* control instruction */26/*00c8*/     EXIT;27/*00d0*/     BRA 0xd0;28/*00d8*/     NOP;

Three observations worth pausing on:

1. Parameters live in a dedicated address space in PTX but in banked constant memoryconstant memoryBanked read-only memory (SASS c[bank][offset]) holding e.g. kernel parameters; readable by non-load/store instructions.Glossary → (c[bank][offset]) in SASS — readable directly by ALU instructions, no load needed.
2. Pascal’s undisassembled rows are control instructionscontrol instructionsKepler+ scheduling words (1 per 3 instructions on Maxwell/Pascal) carrying stall counts, yield hints and dependency barriers — replacing hardware scoreboard checks.Glossary → (introduced with Kepler): one 64-bit word per three instructions carrying stall counts, yield hints, and dependency barriers per instruction — the compiler doing the dependency bookkeeping a hardware scoreboard would otherwise do, reminiscent of the Tera Computer System’s explicit-dependence lookahead.
3. .reuse flags (Maxwell onward) mark operands servable from an operand reuse cacheoperand reuse cacheMaxwell+ structure implied by .reuse flags: lets an operand be read multiple times per register-file access, saving energy.Glossary → without a fresh register-file read — an energy optimization foreshadowing the register-file research of §3.6.

2.2.2 AMD: a fully documented hardware ISA — with a scalar unit

AMD went the opposite way: with Southern Islands (the first Graphics Core Next generation), they published the complete hardware-level ISA specification, which is why lower-level academic simulators favor AMD. (AMD still has a virtual ISA too — HSAILHSAILAMD's virtual ISA in the Heterogeneous System Architecture stack.Glossary →, part of the Heterogeneous System Architecture stack.)

The architectural headline is GCN’s split into separate scalar and vector instructionsscalar vs. vector instructionsGCN split: v_ ops compute per-thread values on 4 vector units; s_ ops compute one value shared by the wavefront on a scalar unit.Glossary → — something NVIDIA GPUs (through Volta) don’t have. Each compute unit couples one scalar unit with four vector units: v_ instructions compute a different 32-bit value per thread; s_ instructions compute a single value shared by the whole wavefront. Where does that help? Control flow — and any computation that’s the same regardless of thread ID (the scalarization opportunity §3.5 explores).

The book’s if/else example makes the machinery concrete. The exec maskexec maskGCN special register predicating which vector lanes execute; the mechanism behind if/else on SIMT hardware.Glossary → predicates the vector lanes; scalar instructions save, invert, and restore it around the two branch bodies. Step through with the lane strip:

1float fn0(float a, float b)2{3    if (a > b)4        return (a * a - b);5    else6        return (b * b - a);7}

1// r0 = "a", r1 = "b"; result in r22 3v_cmp_gt_f32 r0, r1        // a > b -> vcc4s_mov_b64 s0, exec         // save current exec mask5s_and_b64 exec, vcc, exec  // do "if"6s_cbranch_vccz label0      // branch if all lanes fail7v_mul_f32 r2, r0, r0       // result = a * a8v_sub_f32 r2, r2, r1       // result = result - b9label0:10s_not_b64 exec, exec       // do "else"11s_and_b64 exec, s0, exec   // do "else"12s_cbranch_execz label1     // branch if all lanes fail13v_mul_f32 r2, r1, r1       // result = b * b14v_sub_f32 r2, r2, r0       // result = result - a15label1:16s_mov_b64 exec, s0         // restore exec mask

Both sides of the branch execute serially, each under a partial mask — divergence costs throughput even here, at the ISA level. NVIDIA hardware automates this same bookkeeping with a reconvergence stack, which is exactly where Chapter 3 picks up (§3.1.1).

One more GCN insight for later: instead of hardware interlocks for long-latency dependences, GCN uses S_WAITCNTS_WAITCNTGCN instruction waiting until outstanding-operation counters (vector memory, LDS/GDS, export) drop below a threshold; software-managed dependency resolution.Glossary →. Each wavefront has three counters — outstanding vector memory operations, LDS/GDS operations, and register exports — and the compiler inserts S_WAITCNT to hold the wavefront until a counter drops below a threshold. Compare that with NVIDIA’s control instructions above: two vendors, the same conclusion — let the compiler manage dependences; spend the silicon on lanes.

The two philosophies, side by side