The GPU Register File and Operand Collector

Quick quiz: what is the biggest memory on a modern GPU die? Not the L2 cache. Not the L1s. It is the register file. Each SM carries about 256 KB of registers — more than its L1 and shared memory combined — and a hundred-SM GPU therefore holds tens of megabytes of registers, dwarfing every cache on the chip. On a CPU, registers are a few hundred bytes of precious flip-flops; on a GPU they are the single largest SRAM budget. This lesson explains why the sizes flipped, and meets the clever machinery — banks and the operand collector — that makes such a monster readable at all.

The "why" you already know. The zero-cost warp switch works because every resident warp's registers live on-chip permanently — nothing saved, nothing restored. Up to 48–64 warps × 32 threads × dozens of registers each: the multitude has to live somewhere, and the register file is that somewhere. It is the warp scheduler's trick, made of SRAM.

A register 128 bytes wide

First, recalibrate what "a register" means. When a warp's instruction names register r5, all 32 lanes mean their own r5 — so the hardware stores one architectural register as a row of 32 \times 32\text{ bits} = 128 bytes. A single GPU register is the size of a whole CPU cache line. Reading "one operand" means pulling 128 bytes out of SRAM; a fused multiply-add (a \cdot b + c) needs three such reads. Now count the traffic: 32 lanes × 3 operands per FMA, and the SM's four partitions each issuing every cycle. A register file built the CPU way — one flat array of flip-flops with a read port per consumer — would need a port count no silicon can route. Multi-ported SRAM pays roughly quadratic area for its ports; at this scale the flat design is simply impossible.

The solution is the standard one for impossible port counts: fake many ports with many banks. The register file is split into single-ported SRAM banks (say 4 per partition), registers striped across them — register r_n lives in bank n \bmod 4. Four banks can serve four different-bank requests in one cycle; two requests to the same bank must take turns.

one architectural register = 32 lanes × 32 bits = 128 bytes — a cache-line-sized vector;
all resident warps' registers live in the RF permanently — it is the physical substrate of the zero-cost warp switch;
impossible port counts are faked with banked single-ported SRAM plus an operand collector that gathers operands over cycles;
the RF is an occupancy budget: registers-per-thread × resident threads × 4 B must fit in the RF.

The operand collector: assembling an instruction's ingredients

Banks create a new problem: an instruction's three operands may not all be readable this cycle. Enter the operand collector — a small set of staging buffers, one per in-flight instruction. When the scheduler issues an instruction, it parks in a collector slot and its operand requests go to the banks; an arbiter grants each bank to one request per cycle, and operands trickle into the slot as their banks free up. When all three have arrived, the instruction fires into the ALU lanes. Spread across banks, collection takes one cycle; piled onto one bank, it serialises — quietly costing throughput:

Who prevents the pile-ups? Mostly the compiler: register allocation on a GPU is bank-aware, choosing register numbers for hot instructions so their operands stripe across different banks. It is a strange thought from CPU-land — that renaming r6 to r7 could speed up your inner loop — but on a banked register file, register numbers have physics.

Try it: a bank-conflict simulator

The arbiter's behaviour is easy to model: each bank serves one request per cycle, so an instruction's collection time is the depth of its worst bank queue. Watch a conflict-free FMA, a worst-case one, and a partial clash:

// Operand collection: register r lives in bank r % NUM_BANKS; one read per bank per cycle. const NUM_BANKS = 4; function collect(name: string, regs: number[]): void { // settle: sort each operand request into its bank's queue const queues: number[][] = Array.from({ length: NUM_BANKS }, () => []); for (const r of regs) queues[r % NUM_BANKS].push(r); // snap, cycle by cycle: every bank hands over the front of its queue const cycles = Math.max(...queues.map((q) => q.length)); console.log(`${name}: operands ${regs.map((r) => "r" + r).join(", ")}`); for (let c = 1; c <= cycles; c++) { const served: string[] = []; for (let b = 0; b < NUM_BANKS; b++) { const r = queues[b].shift(); if (r !== undefined) served.push(`bank${b}→r${r}`); } console.log(` cycle ${c}: ${served.join(" ")}`); } console.log(` collected in ${cycles} cycle(s)\n`); } collect("FMA A (spread)", [2, 7, 9]); // banks 2, 3, 1 — no conflict collect("FMA B (worst)", [2, 6, 10]); // banks 2, 2, 2 — 3-way conflict collect("FMA C (partial)", [4, 8, 3]); // banks 0, 0, 3 — 2-way conflict

Real collectors juggle several instructions at once, so a bank cycle wasted by one instruction can be used by another — the average cost of conflicts is softened, but never zero. The lesson generalises: whenever you meet banked SRAM, the question is always "who collides?" — you will ask it again, almost verbatim, for shared memory two lessons from now.

The occupancy currency

One more consequence, and it is the big one. The register file is a fixed pot that every resident thread draws from:

\text{regs/thread} \times \text{resident threads} \times 4\,\text{B} \;\le\; \text{RF size}

With a 256 KB register file, a kernel using 32 registers per thread can keep 256\,\text{KB} / (32 \times 4\,\text{B}) = 2048 threads (64 warps) resident — more than the SM's warp slots, so registers don't bind. But a register-hungry kernel using 128 registers per thread caps residency at 512 threads — just 16 warps, and suddenly the warp scheduler may not have enough warps to hide memory latency. Registers are the currency that buys occupancy; this equation returns as one of the three budgets in this module's capstone lesson.

Add it up: 256 KB × 132 SMs on a flagship die is 33 MB of architectural registers — more memory than the entire RAM of the PC that ran DOOM, spent purely on the working values of threads-in-flight. Meanwhile the same die's L2 is a few tens of MB and its L1s a few KB per thread-slice. The size inversion tells you what a GPU really is: a machine whose working set is not data but thread state. CPUs keep a little state for few threads and cache lots of data; GPUs keep enormous state for thousands of threads and stream the data past it. Once you see the register file as the biggest memory on the die, the whole architecture reads differently.

CPU thinking says registers are the cheapest resource in the machine — use as many as you like. On a GPU that instinct silently halves your performance: every register per thread is stolen from residency. A kernel that doubles its register use may drop from 64 resident warps to 32 — and if 32 isn't enough to hide DRAM latency, your arithmetic units idle even though "nothing changed" in the algorithm. This is why CUDA exposes knobs like -maxrregcount and __launch_bounds__: they tell the compiler to spill values to memory rather than burn registers, trading a little arithmetic-side traffic for warps of latency-hiding. Sometimes the spill wins, sometimes it doesn't — but the trade exists only because registers are occupancy currency.