How a GPU actually works — SIMT, warps & the silicon.

A GPU is not a fast CPU. It is a throughput machine that hides memory latency by keeping thousands of threads in flight and switching between them for free. Threads run in lock-step groups of 32 (a warp) on a Streaming Multiprocessor (SM). Understand the warp, the SM, and the memory hierarchy, and every optimization in the rest of this series stops being magic and starts being arithmetic. This is the foundation article — no CUDA code required yet, just the mental model.

Most people who write GPU code start from CPU intuition: a few fast cores, branch prediction, deep caches, out-of-order execution. That intuition is not just unhelpful for GPUs — it is actively wrong. A GPU throws away almost everything a CPU spends transistors on (big caches, branch predictors, speculative execution) and spends those transistors on one thing: raw arithmetic units, thousands of them. The price you pay is that the GPU is dumb about latency. A single thread on a GPU is slow. The trick is that a GPU never runs a single thread — it runs tens of thousands, and uses the sheer number of them to hide the fact that any one of them is stalled waiting on memory.

This article builds the model from the bottom up: the CPU/GPU split, SIMT and the warp, the Streaming Multiprocessor, how warps get scheduled to hide latency, the memory hierarchy that everything bottlenecks on, and finally how this maps onto real silicon — NVIDIA Ampere and Hopper, and AMD RDNA and CDNA. By the end you should be able to look at a kernel and reason about why it is fast or slow before you ever open a profiler.

CPU vs GPU: latency machine vs throughput machine

The cleanest way to understand a GPU is by contrast. A CPU is a latency-optimized device. Its entire design answers the question: "how do I finish this one instruction stream as fast as possible?" So it spends transistors on:

• Large caches (tens of MB of L3) so data is close.
• Branch predictors and speculation so it rarely stalls on a branch.
• Out-of-order execution so independent instructions fill the gaps while one waits.
• High clock speeds (4–5 GHz) and a handful of very wide, very smart cores.

A GPU is a throughput-optimized device. Its design answers a different question: "how do I finish this huge pile of independent work per second?" It does not care how long any single item takes. So it spends transistors on:

• Thousands of simple ALUs instead of a few complex cores.
• Tiny per-thread resources but enormous register files shared across many threads.
• Latency hiding by oversubscription — keep so many threads resident that when some stall on memory, others always have work ready.
• Modest clocks (~1.5–2 GHz) but massive parallel width.

This is why a GPU's single-thread performance is genuinely bad and that is fine by design. The whole machine is a bet that your problem has enough independent parallel work to keep those ALUs busy. Matrix multiply, convolution, attention, ray tracing, particle simulation — these are embarrassingly parallel, which is exactly why GPUs ate them.

The unit of execution: the warp (SIMT)

Here is the single most important idea in GPU programming, and the one CPU intuition gets most wrong: threads do not execute individually. They execute in lock-step groups of 32 called a warp. (AMD calls its group a wavefront; it is 32 or 64 threads depending on architecture. The idea is identical.)

NVIDIA's name for this model is SIMT — Single Instruction, Multiple Thread. It sits between two older ideas:

• SIMD (Single Instruction, Multiple Data), as in CPU vector units: one instruction operates on a fixed-width vector. The programmer or compiler must explicitly pack data into vectors.
• SIMT: you write code as if each thread is independent and scalar — it has its own program counter conceptually, its own registers, its own index. But the hardware executes 32 of those threads together, issuing one instruction to all 32 lanes per cycle.

So SIMT gives you the programming convenience of independent scalar threads with the hardware efficiency of SIMD vector execution. You get to write c[i] = a[i] + b[i] per thread and the hardware quietly runs 32 such adds in one shot.

Warp divergence: the cost of branches

Because all 32 threads in a warp share a single instruction stream, what happens when they hit a branch and disagree? Say half the warp takes the if and half takes the else:

if (threadIdx.x % 2 == 0)
    x = expensive_path_A();   // even lanes want this
else
    x = expensive_path_B();   // odd lanes want this

The hardware cannot run two different instructions on one warp at the same time. So it serializes: it runs path A with the odd lanes masked off (their results discarded), then runs path B with the even lanes masked off. Both halves of the warp pay for both paths. This is warp divergence, and in the worst case a branch can cut your throughput in half (or worse, with nested branches).

The crucial nuance: divergence is only costly within a warp. If your branch splits cleanly along warp boundaries — say, warps 0–3 take one path and warps 4–7 take the other — there is no penalty at all, because each warp is internally unanimous. This is why a condition like if (blockIdx.x < N) is usually cheap (whole blocks agree) while if (data[threadIdx.x] > 0) on random data is expensive (lanes disagree). We will return to this in article 1.2; for now, just hold the rule: keep threads in a warp doing the same thing.

The Streaming Multiprocessor (SM): where warps live

Warps do not float in space — they run on a Streaming Multiprocessor, the GPU's fundamental compute building block. (AMD calls the equivalent a Compute Unit, CU.) A modern data-center GPU has dozens to over a hundred SMs; an NVIDIA H100 has 132. Each SM is a small, self-contained parallel processor with:

• CUDA cores / ALUs — the arithmetic lanes that actually execute warp instructions (FP32, INT32, and on newer parts FP64 and FP16 units).
• Tensor Cores — specialized matrix-multiply-accumulate units (since Volta) that do the heavy lifting for deep learning. A single Tensor Core instruction multiplies small matrix tiles in one shot.
• Warp schedulers — typically 4 per SM, each able to issue an instruction from a ready warp every cycle.
• A register file — large (256 KB per SM on Ampere/Hopper) and partitioned among all resident threads.
• Shared memory / L1 cache — a fast, on-chip scratchpad (up to ~228 KB per SM on Hopper) that threads in a block use to cooperate. This is the single most important optimization lever and gets its own treatment in article 1.2.
• Load/Store units and Special Function Units (SFUs) for memory ops and transcendentals (sin, exp, rsqrt).

When you launch a kernel, you launch a grid of thread blocks. The GPU's global scheduler assigns whole blocks to SMs. A block stays resident on one SM for its entire life — it never migrates. Within the SM, the block's threads are chopped into warps of 32, and those warps compete for the warp schedulers. Several blocks can be resident on one SM simultaneously, as long as their combined registers and shared memory fit. That co-residency is what creates the pool of warps the SM needs to hide latency — the subject of the next section.

Latency hiding: the whole point

Now the payoff. A read from GPU global memory (HBM) costs roughly 400–800 clock cycles of latency. On a CPU, a stall like that is a disaster — out-of-order execution and caches exist precisely to avoid it. A GPU has none of that machinery. So how does it not grind to a halt every time a thread touches memory?

The answer is zero-cost context switching between warps. Because every resident warp's full state (registers, program counter) lives permanently in the SM's huge register file, the warp scheduler can switch from a stalled warp to a ready one in a single cycle, with no save/restore. There is no context to swap — it is all already on-chip.

This reframes the entire optimization problem. On a CPU you fight latency. On a GPU you hide it — and you hide it by having enough independent warps resident that there is always one ready to run. The ratio of resident warps to the maximum the SM can hold is called occupancy, and it is one of the central knobs in GPU performance. Low occupancy means few warps to switch between, which means memory stalls poke through as idle ALU time. We will define occupancy precisely and learn to tune it in articles 1.2 and 2.1; for now the intuition is enough: more resident warps → better latency hiding → higher throughput (up to a point).

The memory hierarchy: where performance is won or lost

Compute is rarely the GPU bottleneck. Memory almost always is. A modern GPU can do tens of TFLOP/s of arithmetic but can only stream a few TB/s of data from its main memory — so for most kernels, the question is not "how fast can it compute" but "how fast can I feed it." The memory hierarchy, fastest and smallest first:

Level	Scope	Approx. latency	Approx. size	Bandwidth
Registers	per-thread	~1 cycle	256 KB / SM	~enormous
Shared memory / L1	per-block	~20–30 cycles	~228 KB / SM	~10s TB/s
L2 cache	whole GPU	~200 cycles	~50 MB (Hopper)	~several TB/s
Global memory (HBM)	whole GPU	~400–800 cycles	40–192 GB	~2–8 TB/s
Host (CPU) RAM over PCIe/NVLink	system	~microseconds	100s GB–TB	~32–900 GB/s

Two facts drive everything. First, each level down is roughly an order of magnitude slower and larger — exactly like a CPU hierarchy, but with the twist that shared memory is programmer-managed, not an automatic cache. Second, the gap between compute throughput and HBM bandwidth is wide and growing; this gap is what the Roofline model (article 2.1) formalizes, and it is why so much of GPU optimization is really memory optimization in disguise.

Mapping the model onto real silicon

The abstract model — SM, warps, schedulers, hierarchy — is stable across GPU generations. What changes each generation is the counts and special units. A quick tour of the architectures named in this series.

NVIDIA Ampere (A100, RTX 30-series)

Ampere (2020) is the workhorse that trained most of the models you have heard of. Per the NVIDIA A100 whitepaper: 108 SMs, 3rd-generation Tensor Cores adding TF32 (a 19-bit format that gave a near-free speedup for FP32 training) and structured sparsity support (skip half the weights for up to 2× throughput). The A100 introduced 40/80 GB of HBM2e at up to ~2 TB/s and a 40 MB L2. Crucially for this series, Ampere added asynchronous copy (cp.async) — letting a thread launch a global→shared-memory copy that proceeds in the background, a key ingredient in modern fused kernels like FlashAttention.

NVIDIA Hopper (H100, H200)

Hopper (2022) is the current data-center standard. 132 SMs, 4th-gen Tensor Cores with the Transformer Engine and FP8 support (doubling throughput again for LLM training/inference), HBM3 at ~3.35 TB/s (H100) or HBM3e at ~4.8 TB/s (H200), and a 50 MB L2. Hopper added thread block clusters and distributed shared memory — letting blocks on different SMs share data directly — plus the Tensor Memory Accelerator (TMA) for bulk asynchronous memory movement. These features exist specifically to feed the Tensor Cores fast enough.

AMD CDNA (MI250X, MI300X)

CDNA is AMD's data-center compute architecture — the GPU behind the KOG monokernel article and a growing share of HPC and AI clusters. Its building block is the Compute Unit (CU), analogous to an SM, and it executes wavefronts (AMD's warps). The MI300X is notable for 192 GB of HBM3 — far more memory than an H100 — which makes it attractive for serving very large models on a single device. AMD's matrix engines are the equivalent of Tensor Cores, and the software stack is ROCm/HIP rather than CUDA, though HIP is deliberately CUDA-like to ease porting.

AMD RDNA (Radeon, consumer)

RDNA is AMD's graphics-focused line (gaming GPUs). It shares the CU/wavefront vocabulary but is tuned for rendering and shader workloads rather than HPC matrix math — narrower wavefronts (32-wide "wave32" mode by default), smaller register files, no high-throughput FP64. The split between RDNA (graphics) and CDNA (compute) mirrors NVIDIA's quieter split between consumer GeForce and data-center parts: same core ideas, different transistor allocation for different markets.

Concept	NVIDIA / CUDA	AMD / ROCm
Compute block	Streaming Multiprocessor (SM)	Compute Unit (CU)
Lock-step thread group	Warp (32 threads)	Wavefront (32 or 64)
Matrix unit	Tensor Core	Matrix Core
Scratchpad	Shared memory	Local Data Share (LDS)
Programming model	CUDA	HIP / ROCm
Bulk async copy	cp.async / TMA	async LDS load

The execution hierarchy, end to end

Tying the software and hardware sides together, here is the full chain from your launch to a lane doing arithmetic:

• Grid: the entire launch — all the work for one kernel call.
• Block: a group of threads that can cooperate via shared memory and barriers; assigned whole to one SM.
• Warp: 32 threads that execute in lock-step; the true hardware unit of scheduling.
• Thread: your scalar program, with its own registers and indices.
• Lane: the physical ALU slot a thread occupies within its warp.

The corresponding memory scopes nest the same way: threads have private registers, blocks share shared memory, and the whole grid shares global memory. Getting work and data to line up with this nesting — block-sized chunks that fit in shared memory, warp-sized accesses that coalesce — is the entire game.

Common misconceptions

FAQ

Is a CUDA core the same as a CPU core?

No, and the marketing name is misleading. A "CUDA core" is a single ALU lane — closer to one lane of a CPU's vector unit than to a full CPU core. A whole SM (with its scheduler, registers, and shared memory) is the better analogue to a CPU core, and a CUDA core is one of its many arithmetic lanes.

If warps are 32 threads, why can I launch blocks of any size?

You can request, say, 100 threads per block, but the hardware rounds up to whole warps — 100 threads become 4 warps (128 lanes) with the last 28 lanes permanently masked off and wasted. That is why block sizes should always be multiples of 32; otherwise you are paying for lanes that do nothing.

What is the difference between a warp and a thread block?

A block is a programmer-chosen group (up to 1024 threads) that can cooperate via shared memory and synchronize with barriers. A warp is a fixed 32-thread hardware slice of a block, and it is the unit the scheduler actually issues instructions to. You program in blocks; the hardware executes in warps.

Does AMD's 64-wide wavefront mean its code is twice as parallel?

Not in any meaningful "free speedup" sense — it just means the lock-step group is bigger, so divergence and coalescing are reasoned about 64 lanes at a time instead of 32. Newer AMD architectures default to 32-wide (wave32) precisely because 64-wide made divergence costlier for many workloads.

Why do GPUs use HBM instead of regular GDDR or DDR?

High-Bandwidth Memory stacks DRAM dies vertically and connects them with an extremely wide bus, trading capacity and cost for raw bandwidth. Since GPU kernels are usually bandwidth-bound, paying for HBM's multi-TB/s throughput is worth it. Consumer cards often use GDDR (cheaper, less bandwidth) because gaming is less bandwidth-starved than large-model AI.

Takeaways

• A GPU is a throughput machine: many simple ALUs, latency hidden by oversubscription, not avoided by caches.
• Threads execute in warps of 32 in lock-step (SIMT). The number 32 underlies nearly every performance rule.
• Divergence is a per-warp cost: branches are only expensive when lanes within a warp disagree.
• The SM holds many resident warps and switches between them for free; the ratio of resident to maximum is occupancy.
• The memory hierarchy — registers → shared/L1 → L2 → HBM — is where most kernels actually bottleneck.
• NVIDIA (SM/warp/Tensor Core/CUDA) and AMD (CU/wavefront/Matrix Core/HIP) use different names for the same ideas.

Next up, article 1.2 turns these ideas into concrete rules you can apply to a kernel: memory coalescing, the shared-memory scratchpad, bank conflicts, occupancy math, and how to actually keep a warp from diverging.

References & further reading

• Kirk & Hwu — Programming Massively Parallel Processors — the canonical text on the CUDA execution model and evolving GPU architecture.
• NVIDIA A100 / Ampere Architecture Whitepaper — SM layout, 3rd-gen Tensor Cores, TF32, async copy.
• NVIDIA H100 / Hopper Architecture Whitepaper — Transformer Engine, FP8, thread block clusters, TMA.
• AMD CDNA & RDNA Architecture — Compute Units, wavefronts, Matrix Cores, the compute/graphics split.
• CUDA C++ Programming Guide — the authoritative reference for SIMT, warps, and the memory model.