GPU Architecture · Reference Card

The Register File & Occupancy

Where a thread's state lives, who decides how much it gets, and why that caps how many threads run at once.

The register file at a glance

What: Fastest storage on the GPU. Holds each thread's working data (its scalar variables). A per-SM resource, physically split across the SM's 4 sub-partitions — not owned by individual CUDA cores.
Size: 65,536 32-bit registers = 256 KB per SM. Unchanged since Kepler (2012) — same on H100 and RTX 5080.
Register width: 32 bits (4 bytes). A 64-bit value (double, pointer) uses two.
Max per thread: 255 registers. Hardware ceiling.
Aggregate: H100: 132 SMs × 256 KB ≈ 33 MB of register file on the die — larger than most CPUs' L3.

Who decides register count — and how to see it

The compiler backend ptxas fixes a per-thread register count for each kernel at compile time. Every thread of that kernel uses the same number. You don't declare it; you influence it.

See it: nvcc -Xptxas -v (or --resource-usage) prints Used N registers per kernel.
Cap it: --maxrregcount=N, or __launch_bounds__(maxThreads, minBlocks) on the kernel.
Overflow: Exceeding the budget spills to "local memory" — which lives in DRAM (cached in L1), despite the name. Slow. Avoid.

Why it caps occupancy

Occupancy = resident warps ÷ the SM's maximum warps. Registers are one of two budgets that limit it. Example on H100 (65,536 registers/SM, max 64 resident warps = 2,048 threads):

Registers / thread	Threads that fit	Resident warps	Occupancy
32	2,048	64	100%
64	1,024	32	50%
128	512	16	25%
255	~257	~8	~12%

More registers per thread → fewer threads resident → a smaller backlog of warps → less latency-hiding headroom. High occupancy isn't always the goal (a register-hungry kernel can still be fast), but you should always know which budget you're spending.

Two independent budgets set occupancy

1. The data register file — 256 KB/SM, carved per thread.
2. The warp/control slots — a fixed count of resident-warp slots per SM (H100: 64). A kernel hits whichever it exhausts first. (Shared memory per block is a third; see the memory lessons.)

Control state lives elsewhere

The 256 KB register file holds data. A thread's control state sits in dedicated scheduler hardware, one slot per resident warp — separate budget, never in the data RF:

Program counter(s) — where execution is.
Active / predicate mask — which of the 32 lanes are on for this instruction (how SIMT does branches: mask lanes off rather than jump).
Convergence stack & barrier state — for divergence and __syncthreads().
7 predicate registers per thread — small per-thread bits.

Pre-Volta

One PC per warp; threads couldn't make independent progress.

Volta+ (incl. H100, RTX 5080)

Independent Thread Scheduling: each thread has its own PC and call-stack state in the scheduler.

Why the context switch is free

Both a thread's register slice and its warp's control slot stay allocated the entire time its block is resident on the SM. Switching warps = the scheduler indexing into structures that are already there. Nothing is saved or restored — unlike a CPU, which must spill one thread's registers to memory to run another.

⌂ Home Glossary →