GPU Architecture Overview — SMs, Warps, and Massively Parallel Execution

A throughput machine, not a fast CPU

A CPU core is built to finish one instruction stream as quickly as possible. It spends most of its transistor budget on the machinery that hides latency for a single thread: deep out-of-order windows, branch predictors, speculative execution, and megabytes of cache per core. A GPU spends almost none of its budget there. It has no meaningful out-of-order execution, no speculation, and a cache-per-thread figure measured in bytes. What it has instead is an enormous number of execution lanes and enough threads resident on-chip that whenever one stalls, another is ready to issue.

That single design choice explains nearly everything else. Registers are huge and statically partitioned, because thread state must stay resident rather than be swapped. Instructions are issued for 32 threads at a time, because per-thread control logic would eat the transistor budget. Latency is measured in hundreds of cycles and tolerated rather than avoided. This article is the map: each section describes one layer at the same altitude and points at the article that goes deep on it.

The SM is the unit of execution

A GPU die is not one big processor. It is an array of streaming multiprocessors - SMs - that share an L2 cache and a pool of high-bandwidth memory. An H100 has 132 SMs, each with 128 FP32 lanes plus tensor cores and load/store units - roughly 16,900 lanes, against something like 64 cores on a top-end server CPU. But the lane count is the least interesting number here. The SM is what matters, because it is the unit that owns resources and makes scheduling decisions.

When you launch a grid, a hardware work distributor hands thread blocks to SMs that have enough free resources to host them. A block lands on exactly one SM and stays there until every one of its threads retires. Blocks never migrate and are never partially resident: either the SM can supply the block's registers, its shared-memory allocation, and a slot for each of its warps, or the block waits. Everything a GPU programmer tunes - occupancy, tiling, register pressure, shared-memory layout - is a negotiation over that one SM's fixed resource budget.

So you can reason about one SM and then multiply: a kernel that keeps one SM saturated keeps all of them saturated, provided the launch has enough blocks to go around. The internal structure of a modern SM - processing partitions, warp schedulers, tensor cores as a separate issue pipe, asynchronous copy engines, thread-block clusters - is the H100 SM article.

Warps and SIMT

Threads are not scheduled individually. The hardware groups them into warps of 32 consecutive threads (AMD's equivalent is the wavefront, historically 64 wide), and a warp is the smallest thing a scheduler can issue to. All 32 lanes execute the same instruction in the same cycle on different data. NVIDIA calls this SIMT - single instruction, multiple thread - to distinguish it from CPU SIMD, where the vector width is baked into the instruction encoding. In SIMT you write ordinary scalar code for one thread and the hardware supplies the vectorization.

Each cycle a warp scheduler picks one eligible warp from its roster and issues that warp's next instruction. Eligible means operands ready, target pipe accepting, not parked on a barrier or a pending memory return. Warps that fail the test are skipped at no cost, because every warp's state already lives in the register file - there is no context to save. Warps that pass issue in order; there is no out-of-order rescue within a warp. How schedulers choose among eligible warps, and what the stall taxonomy looks like, is the warp scheduling article.

Divergence is the tax on the model

The model breaks down when threads in the same warp want to go different ways. If half a warp takes the if and half takes the else, the hardware cannot issue both in one cycle. It executes one side with the other lanes masked off, then the other side, then reconverges: the branch costs the sum of both paths rather than the maximum. A 32-way switch in which every lane picks a different case runs at roughly 1/32 of peak, and a loop with a data-dependent trip count runs until its longest lane finishes. Since Volta each thread has its own program counter, so divergent threads make independent forward progress and constructs that used to deadlock now work - a correctness improvement, not a performance one. The fix is always the same: arrange the data so that threads adjacent in the thread index take the same path.

The memory hierarchy, register to HBM

Five levels, each roughly an order of magnitude larger and slower than the one above. The numbers below are the right shape for recent datacenter parts; exact figures move every generation, and the ratios matter far more than the absolutes.

The register file is the largest and fastest storage on the chip, and it is statically partitioned: the compiler fixes how many registers each thread of a kernel needs, and that allocation is reserved for the thread's entire lifetime. Ask for too many and fewer warps fit; ask for too few and the compiler spills to local memory, which despite the name is backed by cache and ultimately by device memory - register pressure is its own tuning problem.

Shared memory is the level with no CPU analogue: a software-managed scratchpad carved out of the same on-chip SRAM array as L1, scoped to a thread block, explicitly addressed by your code. It is what makes tiling work - load a tile once from HBM, reuse it many times at on-chip latency. It is also banked, so access patterns can collide; see the shared memory article.

At the bottom, HBM bandwidth is enormous but only if you ask for it correctly. The memory system services a warp's 32 loads by coalescing them into as few cache-line transactions as it can: 32 threads reading consecutive floats is one or two transactions, the same threads reading with a large stride is up to 32, and you get a fraction of the bandwidth you paid for. That mechanism is coalesced memory access; the hierarchy in full is the memory hierarchy article.

Occupancy is how latency gets hidden

Put the two previous sections together and the central mechanism falls out. An HBM load takes many hundreds of cycles. Nothing in the SM speculates past it. The only thing that keeps the lanes busy in the meantime is another warp that already has its operands. Occupancy is the ratio of warps resident on an SM to the hardware maximum - 64 on recent datacenter generations - and it is the direct measure of how much latency the SM can absorb.

Three resources cap it, and whichever runs out first binds. Registers: the per-SM register file divided by per-thread demand gives a ceiling on resident warps. Shared memory: a block reserves its allocation for its lifetime, so a large tile means fewer co-resident blocks. Warp and block slots: hard architectural limits. Reading which one binds is the first diagnostic step when a kernel underperforms.

Two misreadings are worth naming. Higher occupancy is not automatically better - a register-rich kernel at 25% occupancy with lots of independent work per thread can beat a register-starved kernel at 75%, because instruction-level parallelism within a warp hides latency too. And occupancy is a capacity for hiding latency, not a guarantee: if every resident warp is stalled on the same dependency, a full SM issues nothing. Achieved occupancy is the number that matters, not the theoretical one. Full treatment: the occupancy article.

The host-device model and kernel launch

A GPU is a coprocessor. The host CPU owns the process, allocates device memory, moves data across the interconnect, and enqueues work; the device executes kernels. Most of the surprising behaviour in GPU programs comes from that boundary rather than from the kernels.

cudaMalloc(&d_x, n * sizeof(float));                 // device allocation
cudaMemcpyAsync(d_x, h_x, bytes, H2D, stream);        // enqueue a copy
saxpy<<<blocks, 256, smem_bytes, stream>>>(n, a, d_x, d_y);  // enqueue a kernel
cudaMemcpyAsync(h_y, d_y, bytes, D2H, stream);        // enqueue a copy back
cudaStreamSynchronize(stream);                        // the only blocking call

Every call but the last is asynchronous. The launch returns to the CPU as soon as the command is queued, and the kernel may not have started. Work within one stream executes in issue order; work in different streams may overlap. The consequences: a wall-clock timer around a launch measures the enqueue, not the kernel; an error surfaced by one call may have originated in an earlier one; and any model that assumes launch is free breaks down for kernels shorter than tens of microseconds - exactly the regime small decode-step kernels live in. Streams, events, and CUDA graphs are covered separately.

The link itself is the other asymmetry. Host-to-device transfers cross PCIe at roughly 64 GB/s per direction on a Gen5 x16 slot, against terabytes per second inside the device: about a fiftyfold gap. That is why the discipline reduces to "get the data on-device and keep it there", and why a model that does not fit in HBM is a structural problem rather than a tuning one. The host path is the PCIe article; the much faster device-to-device path is NVLink and NVSwitch.

What fast means on a GPU

Peak FLOP/s is close to useless as a performance target. The number that predicts what a kernel can achieve is arithmetic intensity: floating-point operations performed per byte moved from HBM. Divide peak throughput by memory bandwidth and you get the device's balance point - hundreds of TFLOP/s over a few TB/s puts it well above 100 FLOPs per byte for reduced-precision matrix math. A kernel below that point is memory-bound and cannot reach peak however the math is scheduled; above it, compute-bound. This is the roofline model, and it sorts almost every GPU workload cleanly.

The sort is stark for AI. A large matrix multiply has intensity proportional to its tile dimension, so it sits comfortably compute-bound - which is why tensor cores exist. Elementwise operations, activations, normalizations, and softmax read and write roughly as many bytes as they do FLOPs: intensity near 1, hopelessly memory-bound, running at a few percent of peak by construction. Autoregressive decoding at batch size 1 is the extreme - every weight is read from HBM to do one vector-matrix product, so step time is model bytes divided by bandwidth.

That framing tells you which optimizations can possibly help. For memory-bound work, move fewer bytes: fuse kernels so intermediates never reach HBM, quantize weights, or restructure to keep tiles on-chip the way FlashAttention does. For compute-bound work, keep the math pipe fed: better tiling, asynchronous copy and software pipelining, larger batches. Applying a compute-bound remedy to a memory-bound kernel is the most common wasted week in GPU optimization, and a profiler settles the question in minutes.

Where the map goes next

Everything above describes one device. Production systems add two more axes. Scaling up: when a model exceeds one GPU's HBM, work is split across devices and the interconnect becomes part of the architecture - NVLink within a node, InfiniBand across nodes, and NCCL collectives as the programming interface. Scaling down: when a workload does not need a whole GPU, one device is subdivided by MIG, time slicing, or MPS - see GPU sharing strategies.

The vocabulary here is NVIDIA's because CUDA is the dominant ecosystem, but the structure is not vendor-specific. AMD's CDNA parts have compute units, wavefronts, LDS, and HBM filling the same roles under ROCm - see AMD Instinct. Google's TPU takes a different bet, replacing the many-SM design with large systolic arrays and a compiler that schedules memory movement statically instead of relying on runtime warp swapping. Reading those against this map shows which parts of GPU architecture are physics and which are one company's choices.