CUDA Execution Model
Definition
The CUDA execution model is the three-level software hierarchy — grid → thread block → thread (with warps of 32 as the implicit fourth level) — that a kernel launch exposes, and its fixed mapping onto GPU hardware (SMs and lanes). You write scalar per-thread code; the hardware executes it as warps in lockstep (SIMT). Getting this software↔hardware mapping crisp is the foundation for reasoning about occupancy, divergence, and memory behavior.
How it works
| Software (what you write) | Hardware (where it runs) |
|---|---|
| Grid (all threads of a launch) | the whole GPU |
Thread block (blockDim threads) | one [[concepts/streaming-multiprocessor |
| Warp (32 consecutive threads) | one warp scheduler / SM sub-partition |
| Thread | one lane; owns private registers |
Rules that matter:
- A block is assigned to exactly one SM by the GigaThread block scheduler and stays there. Grids typically have far more blocks than SMs (H100: 132), so blocks drain through in waves.
- Threads split into warps by consecutive
threadIdx(row-major linearized for 2D/3D). A 256-thread block = 8 warps. - Because a block is co-resident on one SM, its threads cooperate cheaply via two block-scoped mechanisms:
- Shared memory (
__shared__): fast programmer-managed scratchpad in the SM’s L1 region (up to 228 KB/SM on H100). See gpu-memory-hierarchy. __syncthreads(): a barrier — every thread in the block waits until all arrive. Enables the fill-then-consume pattern (stage a tile into SMEM, sync, read it). Possible only because the block is co-resident.
- Shared memory (
- Threads get private registers from the SM’s 65,536-register file; more registers/thread → fewer resident threads → lower occupancy.
- Blocks are (classically) independent — no cheap global barrier across a grid. Hopper thread block clusters relax this by letting nearby blocks share distributed shared memory.
Do not assume blocks run concurrently or in order — any algorithm requiring block to finish before block is broken without atomics or a second launch.
Why it matters
Every LLM kernel — GEMM tiles, GEMM, FlashAttention — is expressed in this model. Reasoning about a kernel’s speed is reasoning about how blocks map to SMs and how many warps stay resident. Shared memory + __syncthreads() are the primitives behind every tiling scheme that keeps LLM activations on-chip instead of round-tripping to HBM.
Taught in
- 01-gpu-architecture-and-simt — §6 the software→hardware mapping, §7 occupancy.
See also
- streaming-multiprocessor — where a block lands
- warp-simt — the 32-lane lockstep unit
- occupancy — resident warps / max warps
- gpu-memory-hierarchy — where shared memory sits
- shared-memory-bank-conflicts — SMEM access hazard
- gpu-systems-for-llms