GPU Execution Model
A GPU hides memory latency with parallelism instead of caches: thousands of threads in flight, scheduled in warps of 32, pinned to streaming multiprocessors that switch between warps for free whenever one stalls. Every performance conversation in an AI infra loop, from occupancy to why decode is slow, rests on this one mechanism.
TL;DR: Software launches a grid of thread blocks; hardware pins each block to one streaming multiprocessor and runs it as warps of 32 threads in lockstep. An SM keeps up to 64 warps resident and switches between them at zero cost whenever one waits on memory, which is how a GPU covers a 500-cycle HBM load without the caches and branch predictors a CPU relies on. Blocks cannot synchronize with each other, and that constraint is what lets one binary scale from 20 SMs to 132.
The two hierarchies, and how they map
A CUDA kernel is launched as a grid of blocks, each block holding up to 1,024 threads. That is the software view, and it is the only one the programmer writes against. The hardware view is a chip full of streaming multiprocessors (an H100 SXM has 132 of them), each with its own register file, shared memory and warp schedulers.
The mapping between the two is rigid in one place and free in another. Rigid: a block is assigned to exactly one SM and stays there until it finishes; the block's threads are carved into warps of 32, and the warp is the unit the scheduler issues instructions to. Free: which SM a block lands on, and in what order blocks run, is the hardware's choice. Nothing in a correct kernel may depend on it.
Two consequences follow, and interviewers listen for both.
Blocks are independent by construction. There is no instruction that makes block 5 wait for block 3; __syncthreads() only synchronizes threads inside one block, through that block's shared memory. A design that assumes ordering between blocks deadlocks the moment the scheduler runs them in the other order, or runs one before the other has been placed at all. Grid-wide synchronization means ending the kernel and launching another (or cooperative groups, with strict limits on how many blocks may be resident at once).
Warps execute in lockstep. Thirty-two lanes receive one instruction. If half the lanes take one branch of an if and half take the other, the SM runs both paths serially with the inactive lanes masked, so a divergent warp can cost twice the time of a uniform one. This is why "branch on threadIdx modulo 2" is a classic anti-pattern and "branch on warp index" is fine.
Latency hiding is the whole idea
A load from HBM costs several hundred cycles. A CPU covers that with large caches, prefetchers and out-of-order execution, all of which spend silicon on making one thread fast. An SM spends its silicon on holding many threads at once. Up to 64 warps (2,048 threads) can be resident per SM on Hopper, and when the warp at the front of the queue stalls on a load, the scheduler issues the next ready warp on the following cycle. There is no context switch in the OS sense: every resident warp already has its registers allocated, so switching costs nothing.
The arithmetic decides how many warps can be resident, and it is worth doing in the room. An H100 SM has 65,536 32-bit registers. A kernel compiled to 64 registers per thread can hold 65,536 ÷ 64 = 1,024 threads, which is 32 warps, half the maximum. Trim it to 32 registers per thread and 2,048 threads fit. Neither number appears in the source; the compiler chooses, -Xptxas -v reports it, and this is why "my kernel got slower after I added three local variables" is a real bug report with a real mechanism (see Occupancy and Register Pressure).
The practical rule that falls out: a GPU wants tens of thousands of threads in flight, not a few fast ones. An H100 with 132 SMs at 2,048 threads each holds 270,336 threads resident. A kernel that launches 4,096 threads leaves 98% of the machine idle no matter how tight its inner loop is, and this is the single most common reason a "GPU version" of some code is slower than the CPU version.
What it costs, and where it shows up
| Resource per SM (H100) | Limit | What it caps |
|---|---|---|
| Resident threads | 2,048 (64 warps) | how many loads can be in flight |
| Registers | 65,536 × 32-bit | threads per SM, via registers per thread |
| Shared memory | up to 228 KB configurable | blocks per SM, via shared bytes per block |
| Resident blocks | 32 | small-block kernels hit this before the thread cap |
Each limit is a ceiling on residency, and residency is what buys latency hiding. A kernel that uses 100 KB of shared memory per block fits two blocks per SM; if each block is 256 threads, that is 512 resident threads, a quarter of the maximum, and the memory system is starved of outstanding loads regardless of how efficient each load is.
The same model explains the two regimes every LLM engineer lives in. A decode step touches every weight once per token, so it is one enormous stream of loads with almost no arithmetic between them; the SM is issuing loads as fast as HBM can serve them and nothing else, which is the memory-bound regime. A prefill GEMM over thousands of tokens keeps the tensor cores busy for hundreds of cycles per tile of data loaded, so the loads hide behind arithmetic. Same hardware, opposite bottleneck, and the execution model is how you tell which one you are in.
What interviewers are listening for
The warm-up question at NVIDIA and the chip companies is "walk me through grid, block, warp and SM," and the marks are not for reciting the hierarchy. They are for three sentences: blocks pin to SMs and cannot synchronize with each other; a warp is 32 threads in lockstep and divergence costs serial execution; the SM hides latency by switching among resident warps, so residency (occupancy) is a performance variable. Candidates who describe the org chart without saying why it is shaped that way get the follow-up "so when is a GPU slower than a CPU?" and usually cannot answer it. The answer: too little parallelism to fill the machine, branch-heavy or pointer-chasing code that divergence and latency both punish, or a problem dominated by the PCIe transfer in and out.
The follow-up they hold in reserve is the register one. "Your kernel runs at 40% occupancy; is that bad?" The strong answer is that occupancy is a means, not an end: if the kernel already issues enough loads to saturate HBM at 40%, more warps buy nothing, and a version that spills registers to reach 100% will be slower. The number to look at is achieved memory throughput against the roofline, not occupancy on its own.
Common misconceptions
- Blocks map to cores. Blocks map to SMs. The "CUDA core" count in marketing is ALU lanes, not independent processors; an H100 has 132 SMs, not 16,896 processors.
- More threads is always faster. Only until the memory system is saturated or the register file is exhausted. Past that point extra threads add scheduling pressure and spills.
- GPU threads are like CPU threads. They are not independently scheduled; the warp is. Reasoning per thread leads you wrong on divergence and on coalescing.
- Block 0 runs first. Nothing guarantees any order. Correct kernels do not depend on it, and the ones that do fail on a different GPU model.
Key takeaways
- Grid, block, thread are software; GPU, SM, warp are silicon. A block pins to one SM; a warp of 32 is the scheduling unit.
- The GPU hides memory latency with resident warps switched at zero cost, not with caches. Residency is capped by registers, shared memory and thread limits per SM.
- Blocks cannot synchronize with each other; grid-wide sync means ending the kernel.
- Divergence inside a warp serializes; divergence across warps is free.
- Occupancy is a lever for latency hiding, not a goal. Measure achieved bandwidth against the roofline instead.