SIMT and GPUs

You know a neural network is mostly matrix multiplication. A GPU runs it as a SIMT throughput machine: tens of thousands of scalar threads, grouped in hardware into warps of 32 (NVIDIA) or 64 (AMD CDNA) that share one instruction stream, spread over one to a few hundred multithreaded cores.³⁹ Three design decisions follow from that and set the rest of the chapter:

• Latency is tolerated, not avoided. Instead of large caches and out-of-order execution, each core holds dozens of warp contexts on chip and issues from whichever is ready, with no branch prediction or speculation.³ The needed concurrency follows Little’s law, and registers per thread trade directly against it.
• Bandwidth is managed explicitly. A software-managed scratchpad (shared memory or LDS) and coalesced access patterns are part of the programming model, because DRAM bandwidth per FLOP keeps shrinking.
• Dense math moved into matrix instructions. Since Volta, tensor or matrix cores deliver roughly 8–16× the throughput of the vector lanes at AI precisions, mostly by amortizing per-instruction overhead.⁵⁷

We use NVIDIA’s and AMD’s published architecture documents side by side; terms differ (SM vs CU, warp vs wavefront, shared memory vs LDS, tensor core vs matrix core) but the ideas map one to one. Arithmetic intensity and the roofline come from What the workload needs; precision formats from Number formats.

The execution hierarchy is grid → thread block (workgroup) → warp → thread. Blocks are distributed to SMs with free capacity, all threads of a block run on that one SM, and they share its shared memory and can barrier-synchronize.³ Each block is split into warps in thread-ID order, so control flow that depends on threadIdx / warpSize never diverges.³ For the CPU counterpart, fixed-width and length-agnostic SIMD with explicit masks, see Multicore and vector units.

Physically, an SM is four scheduler partitions. In GV100 each partition has one warp scheduler, one dispatch unit, a 64 KB register file (16,384 32-bit registers), 16 FP32 lanes, 8 FP64 lanes, 16 INT32 lanes and two tensor cores, so a 32-thread FP32 instruction occupies its partition’s FP32 pipe for two cycles.⁴ Ampere and Hopper keep the four-scheduler layout; each scheduler issues one instruction per cycle from its statically assigned warps.³ AMD’s CDNA CU has four 16-lane SIMDs (a wave64 instruction takes four cycles on one SIMD), four warp pools of up to ten slots each (40 per CU; eight per pool on CDNA 2), and an issue arbiter that can send up to five instructions per cycle to different unit types (vector, vector memory, scalar, LDS, branch).⁹

Divergence mechanics

Before Volta, each warp had one program counter and an active mask. On a divergent branch the warp saves its mask, runs one path and then the other with the inactive threads masked off, and restores the full mask where the paths rejoin.⁴ The standard reconvergence point is the branch’s immediate post-dominator, the first instruction every path must pass through, implemented with a per-warp stack.¹³ Lane efficiency for a region is the average fraction of active lanes; nested divergence compounds multiplicatively.

Volta added independent thread scheduling: per-thread program counters and call stacks, with a schedule optimizer that regroups active threads of a warp into SIMT units. Divergent paths can now interleave, which makes intra-warp locks safe, but at any cycle the warp still issues one instruction for the active lanes of one path.⁴ It also broke implicit warp-synchronous code, which must now use explicit __syncwarp() or warp-level primitives.³

AMD handles divergence by predication with a per-wavefront execution mask, while a separate scalar unit runs warp-uniform control flow and values without touching the vector lanes.⁹ A wider wavefront amortizes each issued instruction over more threads but wastes more lanes when threads disagree; AMD’s documentation notes that the 32-thread mode of its graphics (RDNA) parts reduces divergence penalties and register pressure.⁹

The CUDA guide states the condition directly: full utilization requires every scheduler to have some instruction to issue at every cycle of the latency period. With maximum-throughput instructions, hiding a latency of $L$ cycles takes $4L$ instructions in flight per SM on devices that issue one instruction per scheduler per cycle across four schedulers.³ For 4-cycle dependent math that is 16 warps per SM, fewer if each warp has independent instructions to issue back to back.³ Volkov frames it as Little’s law, concurrency = latency × throughput: on Maxwell, 6-cycle math at 4 instructions per cycle per SM needs 24 instructions in flight, and 368-cycle memory at the sustainable bandwidth needs about 30 memory instructions in flight per SM.¹¹

Measured V100 latencies give a sense of the spread a scheduler must cover¹²:

Occupancy limits

Resident warps per SM = min(warp slots, register limit, shared-memory limit, block limit × warps per block). From Kepler GK180 (compute capability 3.5) through Volta GV100, NVIDIA’s data-center SMs held 64 warps, 65,536 registers and at most 255 registers per thread.⁴ The register limit has cliffs at block granularity. The CUDA guide’s example: 512-thread blocks at 64 registers fit exactly two blocks ($2 \times 512 \times 64 = 65{,}536$, so 32 warps); at 65 registers only one block fits and the SM drops to 16 warps.³ The compiler balances register count against spills; programmers steer it with __launch_bounds__ or -maxrregcount.³ On AMD CDNA the same pressure applies to vector registers (256–512 KiB per CU) and to the 10 warp slots per SIMD.⁹

Two refinements matter in practice. First, synchronization also stalls warps: a block’s warps waiting at a barrier can idle the SM, which is why having several resident blocks helps.³ Second, occupancy is a means, not a goal. Because Little’s law counts instructions in flight, not warps, independent loads and math within a warp can replace warps.¹¹ Under the hood works through the model and its failure modes.

The Expert view adds a registers-per-thread slider that caps resident warps at $\min(16,\ \lfloor 16{,}384 / (32R) \rfloor)$, one scheduler’s share of a 65,536-register SM; a divergence toggle that runs the math block twice at half the lanes; and a scheduler toggle between greedy-then-oldest (the default) and loose round-robin, both as defined in the GPU-architecture literature.¹⁵ Try:

• Set 16 warps, latency 200 and math 16 (about 90% busy). Now raise registers per thread from 32 to 128. At what register count does the cap start to cost utilization, and by how much?
• Turn divergence on. Issue-slot utilization goes up, while useful lane work goes down. Why?
• With 16 warps, switch to round-robin. The warps fall into step and stall together, the effect Narasiman et al. describe.¹⁴

The model is deliberately simple: fixed latencies, no bandwidth limit, no cache, one load per loop and no independent instructions within a warp. It isolates thread-level latency hiding.

Each level trades capacity for latency, bandwidth and energy. A Horowitz 45 nm estimate puts a 32-bit read from an 8 KB SRAM at about 5 pJ and from DRAM at about 640 pJ, more than 100× apart.²⁰ The levels, with representative published numbers:

H100 figures are NVIDIA’s preliminary announcement specs. Sources: register and shared capacities⁶³⁹, caches and HBM⁵⁶⁷.

Details that bite

• Coalescing granularity. Global accesses are served in 32-byte transactions; the number of transactions per warp instruction sets the bandwidth efficiency. Fully coalesced also requires block and array widths to be multiples of the warp size for 2D access.³
• Bank conflicts. Shared memory has 32 banks of 4 bytes; $n$ threads hitting different addresses in one bank serialize into $n$ requests (a broadcast of the same word is free).³ GEMM and attention kernels swizzle their tile layouts to avoid this.
• Spills. Spilled registers go to local memory, which lives in device memory and has global-memory latency.³
• Achievable bandwidth. Microbenchmarks reached 750 of a theoretical 900 GiB/s on V100 HBM2 (83%).¹² Budget against measured, not pin, bandwidth.

Both vendors grew the scratchpad and made it asynchronous

Ampere added an asynchronous copy from global to shared memory that bypasses the register file, so loads for the next tile overlap math on the current one.⁵ Hopper added a Tensor Memory Accelerator that moves whole multi-dimensional tiles by descriptor, thread block clusters, and distributed shared memory that lets a block read a neighboring SM’s shared memory about 7× faster than exchanging through global memory.⁶ AMD’s CDNA 4 grew the LDS from 64 KB to 160 KB per CU with doubled read bandwidth and lets it load directly from L1, explicitly to feed matrix multiplication.⁸ Both are moves toward the explicitly managed, tile-oriented memory of dataflow designs.

A matrix instruction amortizes the per-instruction overhead (fetch, decode, operand fetch, about 30 pJ in a 45 nm estimate) over many MACs: HFMA carries 2,000% overhead, HDP4A 500%, HMMA 22%, IMMA 16%.²⁰ Dally attributes NVIDIA’s ~1000× single-chip inference gain from 2012 to 2023 to number formats (~16×), complex instructions (~12.5×), process (~2.5×) and sparsity (~2×), with model efficiency on top.²⁰ By that accounting, complex instructions contributed about five times as much as a decade of process scaling.

How the instructions are organized

• Warp-level MMA (Volta–Ampere). The warp collectively holds $A$, $B$ and $C$ fragments in registers and issues an MMA; on Volta each tensor core performs 64 FP16 FMAs per clock, 512 per SM.⁴ Ampere doubled per-SM dense throughput to 1,024 FP16 FMAs per clock with four larger tensor cores, and added 2:4 structured sparsity for another 2× on eligible weights.⁵
• Warp-group, asynchronous MMA (Hopper). MMA instructions span a warp group, execute asynchronously, can read operands directly from shared memory, and can reassign register capacity among warp groups. NVIDIA recommends reaching them through CUTLASS or cuBLAS rather than hand PTX.³
• MFMA (AMD CDNA). Wavefront-wide matrix instructions executed by matrix cores that run independently of the VALU, so vector work can overlap.⁹ CDNA 3 does 2,048 FP16 FLOPs per clock per CU in matrix form versus 256 in vector form; CDNA 4 doubles the matrix rate and adds MXFP6/MXFP4.⁷⁸

Precision is part of the contract: inputs in FP16, BF16, FP8 or smaller, accumulation usually in FP32. Markidis et al. measured up to 83 TFLOPS on V100 tensor cores and showed that the error from half-precision inputs can be reduced with extra computation when an application needs it.¹⁸

The feeding problem

Faster math raises the bar for everything feeding it. Matrix throughput grew much faster than register and memory bandwidth, so each generation added data-movement machinery: asynchronous copies (Ampere), TMA and shared-memory operands (Hopper), larger LDS with direct loads (CDNA 4).⁵⁶⁸ The CDNA 3 whitepaper describes improved register-source caching so that each vector register read can feed more downstream vector or matrix operations.⁷

The programmability argument has three parts.

• A scalar-thread ISA. SIMT lets ordinary per-thread code, branches included, run on the device; divergence costs performance, not correctness.³ Independent thread scheduling extended that to fine-grained locking within a warp.⁴
• Scalable decomposition. Independent blocks let one binary scale across SM counts; Hopper’s clusters relax that slightly: a cluster’s blocks are guaranteed to be scheduled concurrently on a group of SMs, which is what makes distributed shared memory possible.⁶
• Libraries and compilers as the real interface. Peak tensor throughput increasingly requires architecture-specific features (asynchronous warp-group MMA, TMA descriptors) that NVIDIA recommends using through CUTLASS and the CUDA-X libraries.³ Tile-level compilers like Triton target the same hardware from a higher level and reach parity with cuBLAS and cuDNN on the kernels they evaluated.¹⁹

A sketch of the structure every fast GPU matrix multiply shares, simplified from CUTLASS’s hierarchy:

for each output tile (one thread block, e.g. 128x128):
  acc[warp tile] = 0                       # accumulators live in registers
  prefetch A[k=0], B[k=0] -> shared buf 0  # async copy / TMA
  for k in K-steps:
    prefetch A[k+1], B[k+1] -> shared buf (k+1)%2
    for each warp (e.g. 64x64 warp tile):
      load fragments from shared buf k%2 -> registers
      acc += MMA(fragA, fragB)             # tensor / matrix cores
    barrier
  write acc -> global memory (fused epilogue: bias, activation)

Each level of the hierarchy is a reuse level: thread-block tiles in shared memory, warp tiles in registers, MMA fragments in the tensor core.¹⁷

Ridge points from the table (dense FP16 FLOPs per HBM byte): H100 $\approx 1{,}000\,\mathrm{T} / 3\,\mathrm{T} \approx 330$; MI300X $\approx 1{,}307 / 5.3 \approx 250$. Both are well above the intensity of batch-1 decode (each FP16 weight is read once per token for one multiply and one add: about 1 FLOP per byte), which is why decode is memory-bound on any GPU. The designs reach their numbers differently: NVIDIA with fewer, larger SMs and a large L2 on one die; AMD with more, smaller CUs across eight compute chiplets, a small L2 per die and a large memory-side cache on the I/O dies.⁶⁷ Earlier anchors from the same documents: V100 had 80 SMs, 6 MB L2 and 900 GB/s⁴; A100 108 SMs, 40 MB and 1,555 GB/s.⁵ Measured rather than peak figures belong to Comparing chips, and how GPUs are joined into larger systems to Scale-up fabrics.

• Generality tax. Work that can’t be expressed as matrix instructions pays the full per-instruction overhead: about 2,000% of the math energy for an FP16 FMA versus 22% for a matrix instruction (45 nm estimates).²⁰ Systolic and dataflow designs remove more of that overhead and more of the flexibility.
• Occupancy versus per-thread state. Registers buy reuse (bigger GEMM tiles, higher arithmetic intensity) and cost warps. CUTLASS notes that accumulators alone take at least half the register budget, so GEMMs run at relatively low occupancy and must hide latency by pipelining instead.¹⁷
• Warp width. Wider warps (64) amortize issue and scheduling over more lanes but waste more on divergence and partial tails; AMD splits the difference by using wave64 on CDNA and wave32 on RDNA.⁹
• Scheduling policy. Fair round-robin makes warps hit long-latency loads together; greedy policies stagger them but can starve warps and disturb cache locality.¹⁴¹⁵
• Memory capacity per device. HBM capacity (80–192 GB in the table above) bounds model and KV-cache size per GPU, which pushes large models across many GPUs and onto the interconnect; see The memory wall and Why the network looks this way.
• Software coupling. Peak throughput depends on architecture-specific instructions reached through vendor libraries and templates; portability across vendors leans on compilers like Triton and on HIP’s close match to CUDA.³¹⁹¹⁰
• Power density. 700–750 W per package drives liquid cooling and rack power design; see Power and cooling.⁶⁷

1. A Little’s-law model of one scheduler

Model one scheduler issuing at most one instruction per cycle. Each warp loops over one load of latency $L$ followed by $C$ dependent math instructions of latency $A$. A warp that is never kept waiting completes a loop in

so one warp sustains $I/P$ instructions per cycle. With $N$ warps and no contention, issue-slot utilization is

That is Little’s law with concurrency counted in warps: throughput (1 per cycle) × latency per instruction ($P/I$) = warps needed.¹¹ Worked example with V100-like numbers ($L \approx 375$, $A = 4$)¹²: with $C = 8$, $P = 407$ and $I = 9$, so about 46 warps per scheduler, nearly three times the 16 slots a quarter-SM has. With $C = 32$, $P = 503$ and $I = 33$, so 16 warps suffice. This is the CUDA guide’s qualitative statement made quantitative: low arithmetic intensity needs more warps.³

Real kernels close the gap with memory-level parallelism. If each warp issues $k$ independent loads before using any of them, the load latency is paid once per $k$ loads: $P \approx L + A k C$, $I = k(1 + C)$, and the warps needed fall roughly $k$-fold. Volkov’s central point is that occupancy and instruction-level parallelism are interchangeable sources of the same concurrency.¹¹ The model also ignores bandwidth: once $N \cdot (\text{bytes per load}) / P$ exceeds the memory system’s throughput, latency rises with load (queueing), and no number of warps helps. That is the roofline’s memory roof.¹¹

2. Computing occupancy

For an NVIDIA SM with $W_{\max}$ warp slots, register file $R_{\mathrm{SM}}$, and shared memory $S_{\mathrm{SM}}$:

• warps per block $w = \lceil T / 32 \rceil$ for $T$ threads per block³
• blocks by registers $\lfloor R_{\mathrm{SM}} / (T r) \rfloor$ for $r$ registers per thread (allocation granularity ignored)
• blocks by shared memory $\lfloor S_{\mathrm{SM}} / s \rfloor$ for $s$ bytes per block
• blocks by slots $\lfloor W_{\max} / w \rfloor$, and the architecture’s block limit

Resident blocks are the minimum of these; occupancy is resident warps over $W_{\max}$. With $R_{\mathrm{SM}} = 65{,}536$, $T = 512$ and $r = 64$, two blocks fit; at $r = 65$ one fits.³ The sim applies the register rule per scheduler: 16,384 registers (the 64 KB per-partition file of GV100⁴) and 16 slots, so $N_{\max} = \min(16,\ \lfloor 512 / r \rfloor)$.

3. Divergence and lane efficiency

For a branch where a fraction $p$ of a warp’s threads take path $X$ ($n_X$ instructions) and the rest take $Y$ ($n_Y$), a mixed warp issues $n_X + n_Y$ instructions while each thread needs only one path. Lane efficiency over the region is

which is 50% whenever the two paths have equal length, whatever the split. The sim’s divergence toggle is that case: issue slots fill faster ($2C$ math per load) but useful lane work halves, so utilization alone overstates performance. The reconvergence stack implements this: on divergence it pushes entries holding the reconvergence PC (the immediate post-dominator) and each path’s mask, executes the top entry until its PC reaches the reconvergence point, then pops.¹³ Fung et al.’s dynamic warp formation regrouped threads from different warps that took the same path, recovering 20.7% on average in their benchmarks for an estimated 4.7% area.¹³

4. Warp scheduling policies

Two baselines from the GPU-architecture literature¹⁵:

• Loose round-robin (LRR): warps take turns; if a warp cannot issue on its turn, the next ready one goes.
• Greedy-then-oldest (GTO): keep issuing from one warp until it stalls, then pick the oldest ready warp (lowest thread IDs among warps that arrived together).

LRR’s fairness has a cost: warps progress in lockstep, reach their long-latency loads together and leave no one to issue while they wait.¹⁴ In the sim, 16 warps with $L = 200$ and $C = 16$ should hide latency ($\lceil 264/17 \rceil = 16$), and GTO comes close; LRR leaves roughly 40% of slots empty because all sixteen warps stall in the same window. Narasiman et al.’s two-level scheduler splits warps into fetch groups, round-robin within a group, so groups reach their stalls at different times; combined with large warps it gained 19.1%.¹⁴ Greedy policies also preserve each warp’s cache locality, which Rogers et al. push further by reducing the number of warps allowed to issue when they would thrash the L1.¹⁵

5. GEMM at low occupancy

A tiled GEMM with an $M_t \times N_t$ block tile and $K$-step $k_t$ loads $(M_t + N_t)\, k_t$ elements per step and does $M_t N_t k_t$ MACs, an intensity of $M_t N_t / (M_t + N_t)$ MACs per element loaded: 64 for a $128 \times 128$ tile. Larger tiles mean more accumulator registers per thread, which is why the accumulators take at least half of each thread’s register budget and why GEMMs run at relatively low occupancy.¹⁷ Latency is hidden instead by double buffering in shared memory and in registers: the next $K$-step’s loads are in flight while the current one is multiplied, providing the memory-level parallelism of section 1 without extra warps.¹⁷ FlashAttention applies the same idea to attention, choosing tile sizes from the SRAM capacity to minimize HBM accesses.¹⁶