In short: A GPU hides memory latency instead of avoiding it, by keeping thousands of threads grouped into 32-thread warps resident on each streaming multiprocessor and switching between them for free, so the chip is optimised for throughput rather than the speed of any single thread.
Throughput versus latency: the design philosophy
A CPU core spends most of its transistors making one thread fast: deep out-of-order windows, branch predictors, and multi-megabyte caches that keep a DRAM access, roughly 80 ns away, off the critical path. A GPU takes the opposite bet. It assumes the workload has enormous parallelism (a matrix multiply, a pixel shader, a convolution), so it spends area on arithmetic and on the machinery to keep that arithmetic fed, and tolerates a slow individual thread.
SIMT: one instruction, 32 threads
NVIDIA calls the execution model SIMT, single instruction, multiple threads. The programmer writes scalar code for one thread; the hardware gangs 32 consecutive threads into a warp that shares a program counter and fetches and decodes each instruction once. That sharing is the area saving: one fetch/decode unit feeds 32 lanes. (AMD calls the equivalent group a wavefront, 32 or 64 wide.)
The catch is divergence. If half a warp takes the if branch and half takes the else, the hardware runs both paths serially with inactive lanes masked off, so you pay for the sum of the paths. A branch on threadIdx.x % 2 can halve throughput, while a branch on threadIdx.x / 32 (uniform within a warp) costs nothing. Since Volta, each thread has its own program counter state and can reconverge independently, but divergent execution still serialises in time.
Key fact: a warp is 32 threads on NVIDIA hardware. Divergent branches inside a warp execute serially; divergence between warps is free.
Inside a streaming multiprocessor
A GPU is an array of identical streaming multiprocessors (SMs); an H100 SXM exposes 132 of them. Each SM is split into four sub-partitions, each with its own warp scheduler, dispatch unit and 32-lane FP32 datapath, plus Tensor Cores, load/store units and special-function units. The big shared resources are the register file, shared memory/L1, and the connection to L2. For a deeper look at scheduling, see our companion piece on SM warp scheduling.
The register file: why context switches are free
A CPU thread switch saves registers to memory. A GPU never does: every resident thread keeps its registers physically allocated in the SM's register file (64K 32-bit registers, 256 KB, on recent parts). With 2,048 threads resident, that gives each thread an average of 32 registers. Switching from warp A to warp B is just the scheduler selecting different register addresses, zero cycles of overhead.
This produces the central tuning constraint, occupancy. If a kernel uses 128 registers per thread, only 512 threads (16 warps) fit per SM, leaving the scheduler few warps to hide latency with. Nsight Compute reports achieved occupancy and the limiter (registers, shared memory or block size). On the silicon side, such a multi-ported, wide register file is a compiled-memory macro whose area and read-port count drive place-and-route in tools like Synopsys Fusion Compiler or Cadence Innovus.
Hiding latency with warps
A global-memory load takes 400 to 800 cycles. A single warp would sit idle for all of them. But if the SM holds 16 to 64 warps, then whenever warp 0 issues a load and stalls, the scheduler issues from warp 1, then 2, then 3. By the time it cycles back, the data has often arrived. Little's law makes it quantitative: to sustain 3 TB/s with ~600 ns of effective round-trip latency you need roughly 1.8 MB of loads in flight across the chip, which is why GPUs need such massive thread counts.
Shared memory versus L2
L2 is a hardware-managed cache shared by all SMs: transparent, coherent at kernel boundaries, and around 200 cycles away. Shared memory lives inside the SM, takes about 20 to 30 cycles, and is explicitly allocated by the programmer, which is what makes tiling work. In a tiled matrix multiply, a block loads a 32x32 tile of A and B from global memory into shared memory once, then each element is reused 32 times from fast SRAM, cutting DRAM traffic by about 32x.
Shared memory has 32 banks of 4 bytes each. If threads in a warp access different banks, all 32 are served in one cycle; if several hit the same bank at different addresses you get an n-way bank conflict and n serial accesses. The classic fix when transposing a tile is padding the row: declare __shared__ float t[32][33] so column accesses fall into different banks.
| Level | Scope | Approx. latency | Managed by |
|---|---|---|---|
| Registers | Thread | ~1 cycle | Compiler |
| Shared memory | Thread block | 20-30 cycles | Programmer |
| L2 cache | Whole GPU | ~200 cycles | Hardware |
| HBM / GDDR | Whole GPU | 400-800 cycles | Programmer (allocation) |
Key fact: bank conflicts cost cycles inside the SM; uncoalesced accesses cost bandwidth to DRAM. Profile both separately.
Memory coalescing
When the 32 threads of a warp execute a load, the memory system merges their addresses into as few 32-byte sectors (grouped as 128-byte lines) as possible. If thread i reads a[i] as 4-byte floats, the warp touches one contiguous 128-byte line: one transaction. If thread i reads a[i*32] (a column of a row-major matrix), each lane lands in a different line, producing up to 32 transactions and using only 4 of every 32 fetched bytes, about 12% efficiency.
This is why GPU codebases favour struct-of-arrays layouts and why transposing through shared memory (read coalesced, write coalesced, swap inside the tile) is a standard interview question. Our article on HBM memory bandwidth shows why the bus itself is the next bottleneck.
Key fact: an H100 can deliver about 3 TB/s from HBM3, but only if warps issue contiguous, aligned, coalesced accesses.
Putting it together: the interview answer
When asked why a GPU beats a CPU on matrix workloads, structure your answer as: (1) many lanes sharing control through SIMT, (2) a huge register file enabling free warp switches, (3) enough resident warps to cover 400+ cycle DRAM latency, (4) shared-memory tiling and coalesced access to protect bandwidth. Then add the caveat: divergent, pointer-chasing, or low-parallelism code is slower on a GPU. Designers building such blocks at Avecas face the same trade-offs when sizing register files and SRAM macros for AI accelerators, and the course material at CourseTron courses and ChipXpert training builds from these fundamentals.
FAQ
What is a warp and why is it 32 threads?
A warp is the group of threads that the hardware schedules together and that executes one instruction in lockstep. Thirty-two is an NVIDIA design choice that balances SIMD efficiency against divergence cost and the width of the datapath and memory transactions.
Why do GPUs not need big caches like CPUs?
They tolerate latency by switching among thousands of threads, so they spend area on lanes and registers instead. Caches still exist, but mainly to filter bandwidth and enable reuse rather than to cut latency.
What is the difference between shared memory and L2?
Shared memory is on-SM SRAM explicitly managed by the programmer and shared by a thread block; L2 is a hardware-managed cache shared by every SM and is several times slower.
What does coalescing mean?
It means the 32 addresses of a warp's memory instruction fall in the same few aligned 128-byte lines, so they merge into minimal DRAM transactions and use nearly all fetched bytes.
References & Further Reading
Ready to go from theory to silicon? Explore the AI hardware and VLSI tracks in the CourseTron course catalogue and build the skills hiring teams test for.