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
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.
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):
| Instruction | 15–12 | 11–8 | 7–4 | 3–0 | Meaning |
|---|---|---|---|---|---|
NOP | 0000 | — | — | — | do nothing for one instruction slot |
BRnzp | 0001 | nzp·0 | imm8 | if (flags ∧ nzp mask) ≠ 0, PC ← imm8 | |
CMP | 0010 | — | Rs | Rt | set N/Z/P flags from Rs − Rt |
ADD | 0011 | Rd | Rs | Rt | Rd ← Rs + Rt |
SUB | 0100 | Rd | Rs | Rt | Rd ← Rs − Rt |
MUL | 0101 | Rd | Rs | Rt | Rd ← Rs × Rt |
DIV | 0110 | Rd | Rs | Rt | Rd ← Rs ÷ Rt (integer) |
LDR | 0111 | Rd | Rs | — | Rd ← mem[Rs] |
STR | 1000 | — | Rs | Rt | mem[Rs] ← Rt |
CONST | 1001 | Rd | imm8 | Rd ← imm8 | |
RET | 1111 | — | — | — | this 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".
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.
Every thread owns sixteen 16-bit registers. Thirteen are ordinary scratch space. The last three are the whole trick:
| Registers | Name | Access | Contents |
|---|---|---|---|
R0–R12 | general purpose | read / write | whatever the kernel computes |
R13 | %blockIdx | read-only | which block this thread belongs to |
R14 | %blockDim | read-only | how many threads per block |
R15 | %threadIdx | read-only | this 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 R13–R15 are wired to its own
coordinates by the hardware. The one expression every kernel opens with,
therefore evaluates to a different number in every thread: thread 2 of block 1 with
4-thread blocks gets a at 0, b at 8, c at 16 for
8-element vectors), thread
R13–R15 are read-only
and hold %blockIdx, %blockDim, %threadIdx;CMP setting N/Z/P flags + BRnzp predicated on
them; RET ends the thread.
Here is
Walk it as thread 3 of block 1 (blocks of 4): lines 1–2 compute
LDR R4, R4 — the address goes in, the data comes out, reusing the register); line 10
adds; lines 11–12 store to RET reports the
thread finished. Launch this with 2 blocks of 4 threads and the eight copies each pick a distinct
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:
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
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.