Designing a Tiny GPU ISA

Eleven modules ago this course promised that text becomes silicon. The capstone collects on that promise: over the next five lessons we will design and simulate a complete, working GPU — PrimerGPU — with SIMT cores, a block dispatcher and an honest, bottlenecked memory system. And every machine ever built starts in the same place: not with gates, but with a contract. Before a single flip-flop exists, someone writes down exactly what software may ask the hardware to do — the instruction set architecture. We did this once before, when we planned a CPU around an RV32I subset; now we do it again for a machine whose whole reason to exist is running one program as ten thousand threads. After touring real GPUs all the way out to chiplet monsters, it is time to build our own — small enough to fit in your head, real enough to breathe.

Design goals: minimal, but sufficient

A good tiny ISA is ruthless about what it leaves out. Our contract has three clauses:

Every choice from here on — 4-bit register fields, 8-bit immediates, a single condition-flag set — follows from squeezing those clauses into sixteen bits.

The PrimerGPU ISA: eleven instructions

Each instruction is one 16-bit word: a 4-bit opcode in bits 15–12, then register or immediate fields. Rd is the destination (bits 11–8), Rs and Rt are sources (bits 7–4 and 3–0), imm8 is an 8-bit unsigned immediate (bits 7–0):

Instruction15–1211–87–43–0Meaning
NOP0000do nothing for one instruction slot
BRnzp0001nzp·0imm8if (flags ∧ nzp mask) ≠ 0, PC ← imm8
CMP0010RsRtset N/Z/P flags from Rs − Rt
ADD0011RdRsRtRd ← Rs + Rt
SUB0100RdRsRtRd ← Rs − Rt
MUL0101RdRsRtRd ← Rs × Rt
DIV0110RdRsRtRd ← Rs ÷ Rt (integer)
LDR0111RdRsRd ← mem[Rs]
STR1000RsRtmem[Rs] ← Rt
CONST1001Rdimm8Rd ← imm8
RET1111this thread is done

Two entries repay a second look. BRnzp is a predicated branch borrowed in spirit from the LC-3: CMP compares two registers and records whether the result was Negative, Zero or Positive; the branch carries a 3-bit mask saying which of those outcomes it fires on. BRnzp n, #4 branches to address 4 if the last compare came out negative — a loop condition in two instructions. And RET is not "return a value"; it is a thread announcing my work here is done — the signal the machinery of the next three lessons will collect, count, and turn into "kernel complete".

Reading a word

Sixteen bits, four fields, no exceptions to memorise — here is instruction 9 of the kernel we are about to write, taken apart:

Fixed fields are a gift to the decoder: bits 11–8 are always where a destination would live, bits 7–4 always the first source. In the next lesson the decode logic will be a page of combinational SystemVerilog precisely because we kept this table boring.

Sixteen registers — and the three that make it a GPU

Every thread owns sixteen 16-bit registers. Thirteen are ordinary scratch space. The last three are the whole trick:

RegistersNameAccessContents
R0R12general purposeread / writewhatever the kernel computes
R13%blockIdxread-onlywhich block this thread belongs to
R14%blockDimread-onlyhow many threads per block
R15%threadIdxread-onlythis thread's position within its block

This is the SIMT index magic — as architecture. Every thread runs the same instruction words, but each thread's R13R15 are wired to its own coordinates by the hardware. The one expression every kernel opens with,

i \;=\; \texttt{\%blockIdx} \times \texttt{\%blockDim} + \texttt{\%threadIdx},

therefore evaluates to a different number in every thread: thread 2 of block 1 with 4-thread blocks gets i = 1 \times 4 + 2 = 6, its neighbour gets 7. One program text, ten thousand different executions — the entire SIMT idea, delivered by three read-only registers. Addressing follows for free: with arrays laid out flat in one data memory (say a at 0, b at 8, c at 16 for 8-element vectors), thread i touches a[i] at address 0+i, b[i] at 8+i — plain register arithmetic, no special addressing modes needed.

The first kernel: vector add, line by line

Here is c[i] = a[i] + b[i] over 8 elements, as every thread of every block will run it — 13 words, no loops, because the parallelism is the loop:

MUL R0, %blockIdx, %blockDim ; R0 = blockIdx * blockDim ADD R0, R0, %threadIdx ; R0 = i — which element am I? CONST R1, #0 ; base address of a CONST R2, #8 ; base address of b CONST R3, #16 ; base address of c ADD R4, R1, R0 ; R4 = &a[i] LDR R4, R4 ; R4 = a[i] ADD R5, R2, R0 ; R5 = &b[i] LDR R5, R5 ; R5 = b[i] ADD R6, R4, R5 ; R6 = a[i] + b[i] ADD R7, R3, R0 ; R7 = &c[i] STR R7, R6 ; c[i] = R6 RET ; this thread is done

Walk it as thread 3 of block 1 (blocks of 4): lines 1–2 compute i = 1\times4+3 = 7; lines 3–5 pin down the memory layout; lines 6–9 fetch a[7] and b[7] (note the idiom LDR R4, R4 — the address goes in, the data comes out, reusing the register); line 10 adds; lines 11–12 store to c[7]; and RET reports the thread finished. Launch this with 2 blocks of 4 threads and the eight copies each pick a distinct i \in \{0,\dots,7\} — the whole vector is covered exactly once, with no thread ever mentioning any other.

An assembler in a screenful

A 16-bit fixed-format ISA is so regular that the assembler is barely a program: shift the opcode up twelve, OR the fields in. Here it is, encoding the kernel and then disassembling it back — the round trip is the proof the encoding loses nothing:

// ── The PrimerGPU assembler: mnemonics in, 16-bit words out — and back again. ── const OPC: Record<string, number> = { NOP: 0b0000, BRnzp: 0b0001, CMP: 0b0010, ADD: 0b0011, SUB: 0b0100, MUL: 0b0101, DIV: 0b0110, LDR: 0b0111, STR: 0b1000, CONST: 0b1001, RET: 0b1111, }; const SPECIALS: Record<string, number> = { "%blockIdx": 13, "%blockDim": 14, "%threadIdx": 15 }; function reg(tok: string): number { if (tok in SPECIALS) return SPECIALS[tok]; return parseInt(tok.slice(1), 10); // "R7" -> 7 } function assemble(line: string): number { const [mn, ...ops] = line.replace(/,/g, " ").split(/\s+/).filter(Boolean); const op = OPC[mn]; if (mn === "NOP" || mn === "RET") return op << 12; if (mn === "CONST") return (op << 12) | (reg(ops[0]) << 8) | parseInt(ops[1].slice(1), 10); if (mn === "BRnzp") { // ops[0] = mask letters "nzp", ops[1] = "#imm" const m = ops[0]; const mask = (m.includes("n") ? 4 : 0) | (m.includes("z") ? 2 : 0) | (m.includes("p") ? 1 : 0); return (op << 12) | (mask << 9) | parseInt(ops[1].slice(1), 10); } if (mn === "CMP") return (op << 12) | (reg(ops[0]) << 4) | reg(ops[1]); if (mn === "LDR") return (op << 12) | (reg(ops[0]) << 8) | (reg(ops[1]) << 4); if (mn === "STR") return (op << 12) | (reg(ops[0]) << 4) | reg(ops[1]); // ADD, SUB, MUL, DIV: Rd, Rs, Rt return (op << 12) | (reg(ops[0]) << 8) | (reg(ops[1]) << 4) | reg(ops[2]); } function rname(r: number): string { return r === 13 ? "%blockIdx" : r === 14 ? "%blockDim" : r === 15 ? "%threadIdx" : "R" + r; } function disassemble(w: number): string { const op = w >> 12, rd = (w >> 8) & 15, rs = (w >> 4) & 15, rt = w & 15, imm = w & 255; const name = Object.keys(OPC).find((k) => OPC[k] === op)!; if (name === "NOP" || name === "RET") return name; if (name === "CONST") return `CONST ${rname(rd)}, #${imm}`; if (name === "BRnzp") { const m = (rd >> 1); // bits 11..9 const s = (m & 4 ? "n" : "") + (m & 2 ? "z" : "") + (m & 1 ? "p" : ""); return `BRnzp ${s}, #${imm}`; } if (name === "CMP") return `CMP ${rname(rs)}, ${rname(rt)}`; if (name === "LDR") return `LDR ${rname(rd)}, ${rname(rs)}`; if (name === "STR") return `STR ${rname(rs)}, ${rname(rt)}`; return `${name} ${rname(rd)}, ${rname(rs)}, ${rname(rt)}`; } // ── The vector-add kernel: c[i] = a[i] + b[i] for 8 elements (a@0, b@8, c@16). ── const kernel = [ "MUL R0, %blockIdx, %blockDim", // R0 = blockIdx * blockDim "ADD R0, R0, %threadIdx", // R0 = i, this thread's element index "CONST R1, #0", // base address of a "CONST R2, #8", // base address of b "CONST R3, #16", // base address of c "ADD R4, R1, R0", // R4 = &a[i] "LDR R4, R4", // R4 = a[i] "ADD R5, R2, R0", // R5 = &b[i] "LDR R5, R5", // R5 = b[i] "ADD R6, R4, R5", // R6 = a[i] + b[i] "ADD R7, R3, R0", // R7 = &c[i] "STR R7, R6", // c[i] = R6 "RET", // this thread is done ]; console.log("addr | binary | hex | disassembled"); console.log("-----+-------------------+------+---------------------------"); kernel.forEach((line, a) => { const w = assemble(line); const bin = w.toString(2).padStart(16, "0").replace(/(....)(?=.)/g, "$1 "); const hex = w.toString(16).padStart(4, "0"); console.log(` ${String(a).padStart(2)} | ${bin} | ${hex} | ${disassemble(w)}`); }); // Round trip: re-assembling every disassembled line must give the same word. const ok = kernel.every((line) => assemble(disassemble(assemble(line))) === assemble(line)); console.log(`\nround trip assemble(disassemble(w)) === w for all 13 words: ${ok}`);

Twenty-six bytes of machine code. That is the entire program our GPU will spend the next four lessons learning to execute — first in one core, then dispatched across several, then through a realistic memory system, and finally end to end with every cycle accounted for.

Because the budget fits. RV32I needs 32-bit instructions largely to afford three 5-bit register fields (32 registers) plus generous immediates. We chose 16 registers, so a register field costs 4 bits, and three of them plus a 4-bit opcode land on exactly sixteen — with the immediate formats (CONST, BRnzp) fitting an 8-bit payload in the same width. Real GPU ISAs make the opposite trade: NVIDIA's SASS instructions are 128 bits, spending lavish encoding space on predication, operand caches and scheduling hints, because instruction fetch is shared by 32 threads at once and amortised into irrelevance. Our design keeps the fetch machinery small instead — the right call for a GPU you must hold in your head. This "minimal but real" style is borrowed from the open-source tiny-gpu project, which proved a working SIMT machine fits in a few files of readable SystemVerilog.

Delete %blockIdx, %blockDim and %threadIdx from the table and everything still assembles — but you no longer have a GPU ISA; you have a slow CPU ISA that happens to be replicated. Every copy of the kernel would compute the same i, load the same two words, and store the same sum to the same address eight times over. The index registers are the sole mechanism by which one program text becomes ten thousand different executions: they are the only per-thread input, so they are the only source of divergence in what threads do. When CUDA writes threadIdx.x, this is what it compiles down to — a read of a register the hardware filled in. Miss this and you miss what a GPU fundamentally is: not a fast chip, but a machine for turning indices into work.