Learn: GPU Architecture & SIMT

What you're learning

How a modern datacenter GPU is physically organized, and the execution model (SIMT) that lets it hide latency with parallelism instead of caches. By the end you should be able to look at a kernel’s register/shared-memory usage and reason about how many warps will be resident per SM — and why an LLM’s forward pass lives almost entirely on this hardware.

This is Lesson 01 of the GPU-systems curriculum. It assumes you know C++, cache hierarchies, and linear algebra, but nothing GPU-specific.

---

1. Learning map

graph TD
    A["CPU design philosophy<br/>(latency-optimized)"] --> B["Why not CPUs for<br/>dense linear algebra?"]
    B --> C["GPU design philosophy<br/>(throughput-optimized)"]
    C --> D["Hardware hierarchy<br/>GPU → GPC → SM → cores"]
    D --> E["SIMT execution model<br/>warps of 32 in lockstep"]
    E --> F["Warp divergence<br/>& predication"]
    E --> G["SW→HW mapping<br/>grid/block/warp/thread"]
    G --> H["Occupancy<br/>active warps / max warps"]
    F --> H
    D --> I["Tensor Cores<br/>MMA units"]
    H --> J["GPU/LLM connection<br/>GEMM + attention"]
    I --> J
    J --> K["Lesson 02:<br/>memory hierarchy"]
    J --> L["Lesson 03:<br/>performance modeling"]

    style C fill:#44a,color:#fff
    style E fill:#4a4,color:#fff
    style J fill:#a44,color:#fff

The spine: design philosophy → what that forces the hardware to look like → the execution model that hardware demands → how you program it → how to keep it fed → why LLMs are a perfect fit.

---

2. Why this matters

Every frontier LLM you’ve heard of was trained on tens of thousands of GPUs for months. Inference for those same models is served on GPUs. A GPU researcher’s job is, bluntly, to make matmuls and attention run closer to the hardware’s peak — and you cannot do that without a mechanical model of why the hardware is shaped the way it is.

The single most important idea in this entire lesson:

The one-sentence thesis

A CPU avoids latency (big caches, out-of-order execution, branch prediction); a GPU hides latency (thousands of threads in flight, so the scheduler always has some warp ready to run while others wait on memory). Everything about GPU architecture follows from that choice.

Dense linear algebra — $C=AB$ where $A,B,C$ are large matrices — has two properties that make it the ideal workload for the “hide, don’t avoid” strategy:

1. Massive independent parallelism. Every output element $C_{ij}={\sum }_k{A}_{ik}{B}_{kj}$ is independent of every other. You have millions of independent units of work — exactly what you need to keep thousands of ALUs busy.
2. High arithmetic intensity (potentially). An $N\times N$ matmul does $\Theta (N^3)$ FLOPs over $\Theta (N^2)$ data, so with good data reuse the ratio of compute to memory traffic is high. This is what lets you become compute-bound instead of memory-bound (Lesson 03).

Hold onto both. They’re why the GPU won.

---

3. CPU vs GPU design philosophy

Start from an unconditional truth: transistors are finite, and you must choose what to spend them on. A chip designer partitions the die between three broad categories — control logic, cache/memory, and arithmetic units (ALUs).

The CPU bet: minimize time-to-result for one thread

A CPU is built to run a single (or few) thread(s) as fast as possible. To do that it spends transistors on avoiding stalls:

• Large caches (a modern server CPU has tens of MB of L2/L3) so most memory accesses never hit DRAM.
• Out-of-order execution, register renaming, deep reorder buffers to find independent instructions to run while one stalls.
• Sophisticated branch predictors and speculative execution so control flow doesn’t bubble the pipeline.
• High clocks (~3–5 GHz) and a few “fat” cores (tens of cores).

The result: excellent performance on branchy, pointer-chasing, latency-sensitive code. But most of the die is control + cache, not arithmetic.

The GPU bet: maximize aggregate throughput

A GPU assumes you have far more parallel work than cores. So instead of making one thread fast, it runs an enormous number of threads and tolerates each one being individually slow:

• Thousands of simple ALUs, grouped into many small cores. An H100 has ~16,896 FP32 “CUDA cores.”¹
• Tiny per-thread control. No branch prediction, no out-of-order execution. Control logic is amortized across many threads that execute together (this is SIMT — Section 5).
• Small caches, huge register file. The register file is enormous precisely so thousands of threads can each keep their state resident and be switched between at zero cost.
• Latency hidden by parallelism, not caches: when warp A stalls on a 400+ cycle DRAM load,² the scheduler instantly issues from warp B, C, D…

graph LR
    subgraph CPU["CPU die budget"]
        C1["Control<br/>(OoO, branch pred)"]
        C2["Large caches"]
        C3["Few ALUs"]
    end
    subgraph GPU["GPU die budget"]
        G1["Tiny control<br/>(shared per warp)"]
        G2["Small caches +<br/>huge register file"]
        G3["Thousands of ALUs"]
    end
    style C3 fill:#a44,color:#fff
    style G3 fill:#4a4,color:#fff

"GPUs are just faster CPUs" — no

A GPU is slower than a CPU on a single-threaded, branchy task, and it always will be. A single CUDA thread has no branch predictor and a lower effective clock. GPUs win only when you have thousands of independent, mostly-uniform units of work. If your problem doesn’t have that structure, the GPU sits idle waiting on memory with nothing to hide the latency.

Quiz: A colleague says "my kernel is slow, let me add more L2 cache prefetching to fix the latency." Why is that the wrong first instinct on a GPU?

Answer

On a GPU you hide memory latency by having more warps in flight, not by making individual accesses faster. The right first question is “how many warps are resident per SM, and is the scheduler ever idle because all of them are stalled?” If occupancy (Section 7) is low, the fix is to reduce per-thread registers or shared-memory usage so more warps fit — giving the scheduler more work to hide latency behind. Prefetching helps at the margin, but the GPU’s whole design premise is parallel slack, not cache locality.

---

4. The hardware hierarchy

A modern NVIDIA datacenter GPU is a strict containment hierarchy. We’ll use the H100 (SXM5, Hopper, GH100) as the running example with real numbers.¹³

graph TD
    GPU["GH100 GPU<br/>132 SMs enabled"] --> GPC["GPCs<br/>(8 Graphics Processing Clusters)<br/>each ~16-18 SMs"]
    GPC --> SM["Streaming Multiprocessor (SM)<br/>the fundamental unit of scheduling"]
    SM --> SP0["SM sub-partition 0"]
    SM --> SP1["SM sub-partition 1"]
    SM --> SP2["SM sub-partition 2"]
    SM --> SP3["SM sub-partition 3"]
    SM --> SMEM["Shared memory / L1<br/>(up to 228 KB/SM)"]
    SP0 --> WS["Warp scheduler<br/>+ dispatch unit"]
    WS --> CC["32 FP32 CUDA cores"]
    WS --> DP["16 FP64 cores"]
    WS --> INT["32 INT32 cores"]
    WS --> TC["1 Tensor Core<br/>(4th gen)"]
    WS --> SFU["Special Function Unit<br/>(sin, exp, rsqrt)"]
    WS --> RF["Register file slice<br/>(64K 32-bit regs/SM total)"]

    style SM fill:#44a,color:#fff
    style TC fill:#a44,color:#fff

Reading top-down:

• GPU → GPCs. The die is partitioned into ~8 Graphics Processing Clusters.¹ GPCs matter for graphics and for some scheduling/locality concerns but are mostly transparent to compute kernels. What you care about is the SM.
• SM — the Streaming Multiprocessor. This is the unit. It has its own register file, its own shared memory / L1, its own warp schedulers and execution units. A block of your kernel is assigned to exactly one SM and stays there. Think of an SM as a small, independent, massively-multithreaded processor. H100 has 132 SMs (out of 144 physically present on the die; the rest are disabled for yield).¹
• SM sub-partitions (a.k.a. processing blocks). Each SM is divided into 4 sub-partitions. Each sub-partition owns one warp scheduler + dispatch unit, a slice of the register file, and its own execution units. So 4 warp schedulers per SM.³
• Execution units, per SM (Hopper):³
  • 128 FP32 CUDA cores (32 per sub-partition) — the general-purpose SIMD lanes.
  • 64 FP64 cores (16 per sub-partition) — HPC double precision.
  • 64 INT32 cores.
  • 4 Tensor Cores (one per sub-partition, 4th generation) — the matmul engines (Section 8).
  • Special Function Units (SFUs) for transcendentals ($\sin$, $\exp$, $1/\sqrt{x}$) — low throughput, used by e.g. softmax and GELU.
  • Load/Store units for memory operations.
• On-SM memory: a 256 KB combined L1/shared-memory block per SM, of which up to 228 KB can be configured as programmer-managed shared memory (Lesson 02 goes deep here).³⁴
• Register file: 65,536 32-bit registers per SM (256 KB).⁴ This is the resource that, more than anything, gates how many threads can be resident.

Why the SM is the mental unit

When you reason about performance, you almost never think about “the whole GPU.” You think about one SM: how many warps are resident on it, how many registers each thread grabbed from its 64K pool, how much of its 228 KB shared memory a block claimed. The GPU is then just “132 of those, running independently.” Scaling to the full chip is the easy part; keeping one SM busy is the hard part.

Rough peak numbers for H100 SXM5 (boost clock ~1.98 GHz), so you have anchors:⁵¹

Engine	Peak	How it’s computed
FP32 (CUDA cores)	~67 TFLOP/s	$132\times 128\times 2\times 1.98\text{e}9$ (the $\times 2$ is FMA = 2 FLOP)
FP64 (CUDA cores)	~34 TFLOP/s	half the FP32 lanes
TF32 (Tensor Core)	~495 TFLOP/s	—
BF16 / FP16 (Tensor Core)	~989 TFLOP/s dense	~1,979 with 2:4 sparsity
FP8 (Tensor Core)	~1,979 TFLOP/s dense	~3,958 with sparsity
HBM3 bandwidth	~3.35 TB/s	80 GB HBM3

Notice the ~15× gap between BF16 Tensor Core (989) and FP32 CUDA core (67). That gap is the entire reason Tensor Cores exist, and the entire reason LLM training uses BF16/FP8. Hold that thought for Section 8.

---

5. The SIMT execution model

Now the core abstraction. SIMT = Single Instruction, Multiple Threads.

The hardware does not schedule individual threads. It groups threads into warps of 32,⁴ and a warp is the atomic unit the warp scheduler issues. All 32 threads (lanes) in a warp share one program counter and execute the same instruction in lockstep — but each lane operates on its own registers and its own data.

This is the transistor payoff from Section 3: because 32 lanes share one instruction fetch/decode and one PC, the control logic is amortized 32×. That’s how you afford thousands of ALUs — you don’t pay for thousands of independent instruction streams.

How SIMT differs from SIMD

You already know SIMD (AVX-512 on a CPU: one instruction operates on a 512-bit vector of 16 FP32 elements). SIMT is similar in spirit but the programming model is scalar-per-thread:

• SIMD: you write vector instructions explicitly; the vector width is exposed in your code.
• SIMT: you write ordinary scalar code from the perspective of one thread; the hardware bundles 32 threads into a warp behind the scenes. Threads can (in principle) follow different control paths — the hardware handles it, at a cost (divergence, below).

The mental model that pays off forever

Write your kernel as “what does thread $i$ do?” — scalar, per-thread. Then remember the hardware executes 32 of those in lockstep as a warp. Almost every GPU performance phenomenon (divergence, memory coalescing, bank conflicts) is a consequence of that 32-wide lockstep bundling.

Each cycle, per warp scheduler

An H100 SM has 4 warp schedulers.³ Each scheduler can have many warps resident (assigned to it) but issues from at most one (or two, via dual-issue) per cycle. Its job: every cycle, pick a resident warp that is ready (its operands are available, not waiting on memory) and issue its next instruction. If a warp stalls on a long-latency load, the scheduler just picks a different ready warp next cycle. This zero-overhead context switch between warps is the latency-hiding mechanism. There is no register save/restore — every resident warp’s registers are physically present in the register file the whole time. That’s why the register file is so huge.

Warp divergence

Because a warp shares one PC, what happens when threads in a warp take different branches?

if (threadIdx.x % 2 == 0) {
    // path A  — even lanes
    x = expensive_A();
} else {
    // path B  — odd lanes
    x = expensive_B();
}

The warp cannot execute two different instructions at once. So the hardware serializes the paths: it runs path A with the odd lanes masked off (predicated inactive, doing no useful work), then runs path B with the even lanes masked off. Both paths execute; each lane’s result is committed only while it is active.

$$t_{\text{divergent}}\approx t_A+t_B\text{vs.}{t}_{\text{uniform}}\approx \max (t_A,t_B)$$

In the worst case — 32 lanes taking 32 different paths — you lose up to 32× throughput on the divergent region, because you’ve serialized what should have been parallel.

Warp divergence is a per-warp phenomenon, not per-block

Divergence only costs you when threads within the same warp (lanes 0–31, 32–63, …) take different paths. A branch on blockIdx.x, or a branch where the condition is uniform across every 32-lane group, costs nothing — every lane in each warp agrees, so no serialization happens. The expensive case is data-dependent branching where neighboring lanes disagree. Design your indexing so that lanes within a warp follow the same control flow (e.g. branch on threadIdx.x / 32, or restructure data so a warp handles a uniform chunk).

Volta+ independent thread scheduling — a caveat

Since Volta, each thread has its own PC and call stack (Independent Thread Scheduling), so lanes can interleave divergent paths and make forward progress independently — this fixed some starvation deadlocks.⁶ But it did not make divergence free: divergent paths still execute at reduced lane utilization. Don’t confuse “no deadlock” with “no cost.” Also: never assume implicit intra-warp synchronization anymore — use __syncwarp() and the explicit-mask *_sync shuffle/vote intrinsics.

Predication

For short branches, the compiler often avoids divergence entirely with predication: instead of branching, it executes both the “then” and “else” instructions on all lanes but uses a per-lane predicate bit to decide whether each instruction actually commits its result. No branch, no PC divergence — just some wasted instruction slots. This is cheaper than a real branch for small bodies (a few instructions), which is why the compiler prefers it. For large branch bodies, real (divergent) control flow wins because predication would execute both full bodies on every lane.

Quiz: Warp divergence. You have if (data[i] > 0) heavy_A(); else heavy_B(); where data[i] is roughly random per element. Warp size 32. Estimate the throughput hit, and give a restructuring that removes it.

Answer

With a random per-element condition, essentially every warp will have some lanes taking A and some taking B, so nearly every warp serializes both branches: cost $\approx t_A+t_B$ per warp instead of the ideal (if lanes agreed) $\max (t_A,t_B)$ — roughly a 2× hit here (up to 32× in the fully pathological many-way case, but a 2-way branch caps the serialization at the two paths). It is not averaged away by “half the lanes each” — the warp pays for both paths because it runs them one after the other.

Fix — sort/partition by the predicate. Reorder the data so that all >0 elements are contiguous and all <=0 elements are contiguous (a stream-compaction / partition pass). Then most warps are uniform — all 32 lanes take A, or all take B — and only the one boundary warp diverges. This is a classic trick: pay a cheap partition to convert data-dependent divergence into uniform control flow.

---

6. The software → hardware mapping

CUDA exposes a three-level software hierarchy that maps onto the hardware hierarchy from Section 4. Getting this mapping crisp is essential.

Software (what you write)	Hardware (where it runs)
Grid (all threads of a kernel launch)	the whole GPU
Thread block (blockDim threads)	scheduled onto one SM, entirely
Warp (32 consecutive threads of a block)	one warp scheduler / sub-partition
Thread	one lane; owns private registers

graph TD
    subgraph SW["Software: kernel launch"]
        Grid["Grid<br/>gridDim blocks"] --> B0["Block 0"]
        Grid --> B1["Block 1"]
        Grid --> Bn["Block N"]
        B0 --> T["blockDim threads<br/>= (blockDim/32) warps"]
    end
    subgraph HW["Hardware"]
        SM0["SM 0"]
        SM1["SM 1"]
        SMn["SM 131"]
    end
    B0 -.assigned to.-> SM0
    B1 -.assigned to.-> SM1
    Bn -.assigned to.-> SMn
    style B0 fill:#44a,color:#fff
    style SM0 fill:#4a4,color:#fff

The rules that matter:

1. A block is assigned to exactly one SM and never migrates. The GPU’s block scheduler (GigaThread engine) hands blocks out to SMs as they free up capacity.⁴ You have grids with far more blocks than SMs; they drain through the 132 SMs in waves.
2. A block’s threads are split into warps by consecutive threadIdx. For a 1D block, warp $w$ = threads $[32w,32w+31]$. A block of 256 threads = 8 warps. (For 2D/3D blocks, threads are linearized in row-major order first, then chopped into warps.)
3. Because a whole block lives on one SM, its threads can cooperate cheaply. Two mechanisms, both only available within a block:
   • Shared memory — a fast, programmer-managed scratchpad in the SM’s L1 region (up to 228 KB/SM). All threads in a block read/write it. This is how threads exchange data without round-tripping to DRAM. (Lesson 02.)
   • __syncthreads() — a barrier: every thread in the block waits until all threads reach it. Lets you safely fill shared memory in one phase and consume it in the next. It synchronizes a block, which is possible precisely because the block is co-resident on one SM.
4. Threads get private registers carved from the SM’s 64K-register file. More registers per thread → fewer threads fit → lower occupancy (next section).
5. Blocks are (classically) independent. There’s no cheap global barrier across blocks within a kernel — different blocks may run at different times on different SMs, or not concurrently at all. (Cooperative Groups / grid-wide sync and Hopper’s thread block clusters — which let a small group of blocks on nearby SMs share “distributed shared memory” — relax this;³ we’ll touch clusters in a later lesson. The default mental model stays: blocks are independent.)

Don't assume blocks run concurrently or in order

If your kernel launches 100,000 blocks on a 132-SM GPU, only a limited number are resident at any instant (bounded by occupancy). The rest wait. Any algorithm that assumes block $i$ finishes before block $j$ starts, or that all blocks are live simultaneously, is broken. Inter-block coordination needs atomics, a second kernel launch, or explicit grid-sync primitives — not luck.

Quiz: Block → SM mapping. You launch a kernel with blockDim = 256, gridDim = 1000 on an H100 (132 SMs). (a) How many warps per block? (b) Can all 1000 blocks run at once? (c) If a block allocates 100 KB of shared memory, how many such blocks can co-reside on one SM, ignoring register limits?

Answer

(a) $256/32=8$ warps per block.
(b) No. At most a handful of blocks per SM are resident (bounded by occupancy limits — registers, shared memory, and the hard cap of 32 resident blocks / 64 resident warps per SM). With 132 SMs and, say, 2–4 resident blocks each, only ~260–530 blocks are live at once; the other ~500+ blocks drain through in later waves as SMs free up.
(c) An H100 SM has up to 228 KB of shared memory. $\lfloor 228/100\rfloor =2$ blocks can co-reside based on shared memory alone. Note this also means only $2\times 8=16$ warps are resident from shared-memory pressure — well under the 64-warp max, so this kernel is shared-memory-limited and will have low occupancy (25%). That’s the kind of trade-off Section 7 is about.

---

7. Occupancy

Occupancy = (active warps resident on an SM) / (maximum warps the SM supports). On H100 the max is 64 warps per SM (2048 threads).⁴ If 32 warps are resident, occupancy is 50%.

Why do we care? Back to Section 3: latency is hidden by having other ready warps to issue when one stalls. More resident warps → deeper pool of work → better chance the scheduler always has something to run → memory/pipeline latency stays hidden. Occupancy is the knob for latency hiding.

What limits occupancy

Three hard resource ceilings, per SM. The achievable occupancy is the minimum across all of them:

1. Registers. The SM has 65,536 32-bit registers.⁴ If each thread uses $R$ registers, the max resident threads is $\lfloor 65536/R\rfloor$ (in practice rounded down to allocation granularity). Example: a kernel using $R=64$ regs/thread → $65536/64=1024$ threads = 32 warps = 50% occupancy, even before considering anything else. At $R=32$ → 2048 threads = 64 warps = 100%.
2. Shared memory per block. The SM has up to 228 KB.⁴³ If a block uses $S$ bytes, at most $\lfloor 228\text{KB}/S\rfloor$ blocks fit. (See the quiz above.)
3. Block/warp count caps. Hardware limits: max 32 resident blocks per SM and max 64 resident warps per SM.⁴ Tiny blocks (e.g. 32 threads = 1 warp) hit the 32-block cap at only 32 warps → capped at 50% even with plenty of registers.

$${\text{warps}}_{\text{active}}=\min \left(\underset{\text{register limit}}{\underbrace{\lfloor \frac{65536}{R\cdot 32}\rfloor \cdot \frac{\text{blockWarps}}{\text{blockWarps}}}},\underset{\text{smem limit}}{\underbrace{\lfloor \frac{228\text{KB}}{S}\rfloor \cdot \text{blockWarps}}},\underset{\text{warp cap}}{\underbrace{64}},32\cdot \text{blockWarps}{|}_{\text{block cap}}\right)$$

(Register allocation is per-thread but granularized; the CUDA Occupancy Calculator / cudaOccupancyMaxActiveBlocksPerMultiprocessor does the exact arithmetic — use it rather than doing it by hand in production.)

Why 100% occupancy is NOT the goal

This is the trap.

The occupancy ≠ speed myth

High occupancy is a means to latency hiding, not an end. Pushing occupancy to 100% often means cutting registers per thread, which forces the compiler to spill live values to local memory (which lives in slow DRAM/L2). Now every spilled access is a memory round-trip — you traded a scheduling benefit for a bandwidth cost, and the kernel gets slower. Many of the fastest kernels in the world — cuBLAS GEMM tiles, FlashAttention — run at 30–60% occupancy on purpose: they hoard registers and shared memory to keep a large working set (a matmul tile) on-chip, and they hide latency through instruction-level parallelism (many independent FMAs in flight per thread) rather than through thread-level parallelism (many warps).⁷⁸ Once you have “enough” warps to cover the latency, more warps buy nothing and the register/SMEM you spent to get them would’ve been better spent on the working set.

The right framing: occupancy needs to be high enough to saturate the memory pipeline given the kernel’s latency and ILP — and no higher. That threshold is often well below 100%. Measuring it (via the profiler’s “achieved occupancy” and stall-reason counters) is Lesson 03 material.

Little's Law is the real principle underneath

The number of memory requests you must keep in flight to saturate bandwidth is $N=\text{latency}\times \text{throughput}$. Occupancy is just one way to generate that in-flight parallelism (more warps = more outstanding requests). ILP within a thread is another way.⁷ This is why a low-occupancy, high-ILP kernel can fully saturate the GPU — it satisfies Little’s Law through a different term.

---

8. Tensor Cores

Go back to the H100 peak table: BF16 Tensor Core ~989 TFLOP/s vs FP32 CUDA core ~67 TFLOP/s. A ~15× gap (≈30× with sparsity). Where does it come from, and why does it exist?

What a Tensor Core is

A Tensor Core is a dedicated hardware unit that performs a small matrix-multiply-accumulate (MMA) in one operation:

$$D=A\cdot B+C$$

where $A,B,C,D$ are small tiles (e.g. $16\times 16$-ish, exact shapes vary by generation and instruction). In a single issued instruction, the unit performs hundreds of multiply-adds. Contrast with a CUDA core, which does one FMA (one multiply + one add) per lane per cycle. The Tensor Core packs a whole tile-matmul’s worth of MACs into fixed-function silicon.

Why it exists — the arithmetic

A scalar FMA is $O(1)$ work. An $n\times n$ tile matmul is $O(n^3)$ MACs over $O(n^2)$ inputs. By building a unit that consumes a tile of inputs and internally wires up all $n^3$ multiply-adds, you:

• Amortize instruction overhead massively — one instruction, hundreds of MACs — instead of paying fetch/decode/issue per FMA.
• Amortize register/operand movement — operands are read once and reused across the systolic array of multipliers, instead of each FMA re-reading operands.
• Use lower precision inputs (BF16/FP8) with a higher-precision accumulator (FP32), which is numerically fine for deep learning and lets you pack more MAC units in the same area/power.

That’s the 15× gap: it’s the difference between general-purpose SIMD lanes doing one MAC each and fixed-function matmul silicon doing a tile at a time.

How you use it

You rarely write raw Tensor Core instructions. You use them through:

• cuBLAS / cuBLASLt / CUTLASS (library GEMMs — what PyTorch’s matmul dispatches to),
• the WMMA API or inline mma.sync / wgmma PTX (Hopper’s warp-group MMA operates at the granularity of a warp group = 128 threads, feeding a large tile),⁴
• Tensor Memory Accelerator (TMA) on Hopper to stream tiles from global → shared memory asynchronously to keep the Tensor Cores fed (Lesson 02).³

The recurring engineering problem is feeding the Tensor Cores: they consume data so fast that memory bandwidth, not compute, becomes the bottleneck unless you carefully tile through shared memory and registers. That’s the bridge to Lessons 02 and 03.

Precision ladder

The Tensor Core throughput ladder — TF32 (~495) → BF16/FP16 (~989) → FP8 (~1,979), each roughly doubling — is the reason for the industry march down the precision ladder in LLM training and inference. Every step down the ladder is a ~2× compute win, paid for with numerical-stability engineering (loss scaling, per-tensor/per-block scaling factors, keeping sensitive ops in higher precision). Blackwell pushes further to FP4.⁹

Quiz: Your attention kernel spends most of its time in $Q{K}^{\top }$ and $\text{softmax}(\cdot )V$ (two matmuls) plus a softmax in between. Which parts hit the Tensor Cores, which don't, and what does that imply for where the bottleneck lands?

Answer

The two matmuls ($Q{K}^{\top }$ and $PV$) map onto Tensor Cores — they’re GEMMs. The softmax (exp, max-reduction, normalization) does not: exp runs on the SFUs and the reductions on CUDA cores, both far lower throughput than the Tensor Cores. So even though matmuls dominate the FLOP count, the softmax’s memory traffic and lower-throughput ops can become the bottleneck — which is exactly why FlashAttention fuses the whole thing, keeps the intermediate scores in shared memory/registers (never writing the $N\times N$ score matrix to DRAM), and computes softmax online.⁸ It’s a case study in feeding the Tensor Cores and not getting stalled by the non-GEMM glue. (More in Lesson 03 and the attention lessons.)

---

9. GPU/LLM connection — why LLMs live on this hardware

Here’s the payoff. An LLM forward or backward pass is overwhelmingly matrix multiplication. Break down a transformer layer:

• QKV projections, attention output projection, and the MLP (up/down projections) are all GEMMs — dense matrix multiplies. For a model with hidden size $d$ and MLP ratio 4, the MLP alone is two $[B\cdot S,d]\times [d,4d]$-shaped matmuls per layer. Across a full model these projections are the large majority of the FLOPs.
• Attention ($Q{K}^{\top }$ then $\text{softmax}\cdot V$) is two more batched matmuls per head, plus the softmax glue.
• Elementwise ops (LayerNorm/RMSNorm, residual adds, activations like SwiGLU) are memory-bound “glue” between the matmuls.

So a transformer is, mechanically, a long chain of GEMMs with elementwise ops between them. And GEMM is the canonical workload for everything in this lesson:

1. Embarrassingly parallel — every output tile is independent → fills thousands of ALUs, gives the scheduler endless ready warps to hide latency (Sections 3, 5, 7).
2. Uniform control flow — no data-dependent branching in the inner loop → zero warp divergence (Section 5). Every lane does the same FMA on different data. SIMT is perfect for this.
3. Maps directly onto Tensor Cores — GEMM is tiled MMA. The ~989 BF16 TFLOP/s exists to run exactly this (Section 8).
4. High arithmetic intensity when tiled well — $O(N^3)$ compute over $O(N^2)$ data means you can become compute-bound and actually approach that 989 TFLOP/s, instead of starving on bandwidth (Section 2, and Lesson 03’s roofline).

This is not a coincidence — it’s why the field is GPU-dominated. The transformer’s computational profile (dense, uniform, high-intensity linear algebra) is a near-exact match for the workload GPUs were optimized to devour. A CPU running the same GEMMs would be starved: too few ALUs, no Tensor Cores, latency it can only avoid (caches) rather than hide (parallelism), and nowhere near the memory bandwidth.

Two threads this opens, which the next lessons pick up:

Where this goes next

Compute is only half the story. Peak 989 TFLOP/s assumes you can feed the Tensor Cores from HBM at 3.35 TB/s⁵ — and for many real LLM ops (attention, decode-time matmuls with tiny batch, all the elementwise glue) you cannot, so you’re memory-bound, not compute-bound. Understanding when you’re limited by compute vs. bandwidth is the whole game.

• Lesson 02 — 02-gpu-memory-hierarchy: registers → shared memory → L2 → HBM, coalescing, bank conflicts, and why tiling exists. This is how you feed the beast.
• Lesson 03 — performance modeling: the roofline model, arithmetic intensity, and how to decide whether a kernel is compute- or memory-bound before you profile it.

---

10. Practice & resources

Hands-on exercises

Exercise A: Query your device and connect numbers to this lesson

Write a tiny CUDA program that calls cudaGetDeviceProperties and prints:

• multiProcessorCount (should be 132 on H100 SXM), maxThreadsPerMultiProcessor (2048 → 64 warps), maxThreadsPerBlock (1024), warpSize (32).
• regsPerMultiprocessor (65536), sharedMemPerMultiprocessor / sharedMemPerBlockOptin (up to 228 KB on H100 with the opt-in), regsPerBlock.
• maxBlocksPerMultiProcessor (32), memoryClockRate and memoryBusWidth → compute theoretical HBM bandwidth and compare to the ~3.35 TB/s spec.

Then, for each printed number, write one sentence tying it to a section above (e.g. “regsPerMultiprocessor = 65536 → a kernel using 64 regs/thread caps at 32 warps = 50% occupancy”). Also print the peak FP32 FLOP/s using multiProcessorCount × 128 × 2 × clockRate and check you get ~67 TFLOP/s.

Why: this makes the abstract hierarchy concrete on whatever hardware you actually have, and forces the register↔occupancy arithmetic into muscle memory.

Exercise B: Occupancy reasoning from register/SMEM usage

Problem. A kernel is compiled (check with nvcc -Xptxas -v or --resource-usage, or the profiler) and reports 72 registers/thread and 48 KB shared memory per block, launched with blockDim = 128 (4 warps/block), on an H100 (65,536 regs/SM, 228 KB SMEM/SM, 64-warp cap, 32-block cap). Compute achieved occupancy and identify the binding constraint. Then propose two changes to raise it and state their risk.

Worked solution.

• Register limit: $65536/72=910$ threads → floor to allocation granularity, ~896 threads → 28 warps → 7 blocks (128 threads each) → but let’s take threads: 28 warps ≈ 44% occupancy from registers alone. (Register allocation is granular; the occupancy calculator would give the exact figure, ~28 warps.)
• Shared-memory limit: $\lfloor 228/48\rfloor =4$ blocks → $4\times 4=16$ warps → 25% occupancy. ← this is smaller.
• Block/warp caps: 4 blocks is well under 32; 16 warps under 64. Not binding.
• Binding constraint = shared memory → achieved occupancy ≈ 25% (16/64 warps).
• Fix 1: cut shared-memory per block (e.g. smaller tile, or split into two smaller blocks). Dropping to 24 KB → 9 blocks fit → but now the register limit (or 32-block/64-warp cap) binds; recompute. Risk: smaller tiles mean less data reuse → more DRAM traffic (Lesson 02 trade-off).
• Fix 2: reduce registers (via __launch_bounds__ or -maxrregcount) to raise the register-limited ceiling — but this only helps if SMEM stops being the binding constraint first, and it risks register spilling to local memory, which can be a net loss (Section 7’s warning).
• Meta-point: before optimizing, ask whether 25% is even hurting — if the kernel is compute-bound with high ILP (like a GEMM tile), 25% may already saturate the pipeline and none of this matters. Profile the stall reasons first (Lesson 03).

Exercise C (stretch): Provoke and measure warp divergence

Write two versions of a kernel over an array: (1) if (idx % 2) ... else ... (adjacent lanes diverge — worst case), (2) if ((idx / 32) % 2) ... else ... (whole warps agree — no divergence). Make both branches equal, non-trivial work. Time both. Then inspect with the profiler’s “branch efficiency” / “warp execution efficiency” metric (Nsight Compute: smsp__thread_inst_executed_per_inst_executed). Confirm version (1) is ~2× slower with ~50% warp execution efficiency, version (2) is ~100%. This makes Section 5 tangible.

Real resources

• 📖 PMPP — Programming Massively Parallel Processors, 4th ed. (Hwu, Kirk, El Hajj).¹⁰ The canonical textbook. Chapters 1–4 cover exactly this lesson (architecture, data parallelism, scalable execution, memory) in more depth. Read it alongside this curriculum; it’s the single best foundation.
• 📘 NVIDIA CUDA C++ Programming Guide — the authoritative reference. Sections: Programming Model (grid/block/thread), Hardware Implementation (SIMT architecture, warps), and the Compute Capability appendix (the exact per-SM resource limits for Hopper/Blackwell). Bookmark it; you’ll return constantly.
• 🎥 GPU MODE (formerly CUDA MODE) lecture series — community lectures/YouTube + Discord, working through CUDA and kernel optimization hands-on. Excellent for turning book knowledge into kernels. Start from the early lectures.
• 📗 GPU Glossary by Modal — a crisp, hyperlinked reference for exactly the terms in this lesson (SM, warp, SIMT, occupancy, Tensor Core, CUDA core). Great for a fast lookup or sanity-check on terminology.
• 🔧 NVIDIA Hopper Architecture Whitepaper (GH100) — for the real numbers (132 SMs, 4th-gen Tensor Cores, TMA, thread block clusters) straight from the source.

---

11. What’s next

Three threads, in order:

1. 02-gpu-memory-hierarchy (next lesson) — registers → shared memory → L2 → HBM3, memory coalescing, shared-memory bank conflicts, and tiling. This is how you feed the ALUs and Tensor Cores you just learned about. Without it, all that compute starves.
2. 03-performance-modeling-roofline (planned) — the roofline model, arithmetic intensity, and deciding compute-bound vs memory-bound before profiling. Turns “it’s slow” into a quantitative diagnosis.
3. 06-memory-wall-and-flashattention (planned) — the flagship case study that ties together SIMT, occupancy, shared-memory tiling, and Tensor Cores to make attention fast. Everything in this curriculum converges here.

---

Topic hub: gpu-systems-for-llms | Lesson 01 of the GPU-systems curriculum | Filed: 2026-09-02

---

References

Footnotes

1. NVIDIA, NVIDIA H100 Tensor Core GPU Architecture whitepaper (GH100/Hopper), 2022 — full GH100 die has 144 SMs (132 enabled on SXM5), 8 GPCs, and the per-SM/peak-throughput tables. https://resources.nvidia.com/en-us-tensor-core ↩ ↩² ↩³ ↩⁴ ↩⁵

2. W. Luo, R. Fan, Z. Li, D. Du, Q. Wang, X. Chu, “Benchmarking and Dissecting the NVIDIA Hopper GPU Architecture,” arXiv:2402.13499, 2024 — microbenchmarked global-memory (DRAM) load latency on H100/H800. https://arxiv.org/abs/2402.13499 ↩

3. M. Andersch et al., “NVIDIA Hopper Architecture In-Depth,” NVIDIA Technical Blog, 2022 — per-SM composition (4 sub-partitions/warp schedulers, 128 FP32 / 64 FP64 / 64 INT32 CUDA cores, 4 fourth-gen Tensor Cores, 256 KB L1/SMEM with up to 228 KB SMEM), TMA, and thread block clusters / distributed shared memory. https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/ ↩ ↩² ↩³ ↩⁴ ↩⁵ ↩⁶ ↩⁷ ↩⁸

4. NVIDIA, CUDA C++ Programming Guide — “SIMT Architecture” and “Hardware Implementation” (warp size 32, GigaThread block scheduler, warp-group MMA), plus the Compute Capability 9.0 appendix (65,536 32-bit registers/SM, 255 registers/thread, 228 KB max SMEM/block, 64 resident warps and 32 resident blocks per SM). https://docs.nvidia.com/cuda/cuda-c-programming-guide/ ↩ ↩² ↩³ ↩⁴ ↩⁵ ↩⁶ ↩⁷ ↩⁸ ↩⁹

5. NVIDIA, NVIDIA H100 Tensor Core GPU Datasheet, 2023 — H100 SXM5 boost clock ~1.98 GHz, ~67 TFLOP/s FP32, ~989 TFLOP/s BF16/FP16 dense (1,979 with 2:4 sparsity), ~1,979 TFLOP/s FP8 dense (3,958 sparse), 80 GB HBM3 at 3.35 TB/s. https://resources.nvidia.com/en-us-gpu-resources/h100-datasheet-24306 ↩ ↩²

6. NVIDIA, NVIDIA Tesla V100 GPU Architecture whitepaper (GV100/Volta), 2017 — introduction of Independent Thread Scheduling (per-thread PC and call stack); see also the CUDA C++ Programming Guide, “Independent Thread Scheduling” and __syncwarp()/*_sync intrinsics. https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf ↩

7. V. Volkov, “Better Performance at Lower Occupancy,” GPU Technology Conference (GTC), 2010 — latency can be hidden via instruction-level parallelism at low occupancy; high occupancy is not required for peak throughput. https://www.nvidia.com/content/gtc-2010/pdfs/2238_gtc2010.pdf ↩ ↩²

8. T. Dao, D. Y. Fu, S. Ermon, A. Rudra, C. Ré, “FlashAttention: Fast and Memory-Efficient Exact Attention with IO-Awareness,” arXiv:2205.14135, 2022. https://arxiv.org/abs/2205.14135 ↩ ↩²

9. NVIDIA, NVIDIA Blackwell Architecture Technical Brief, 2024 — fifth-gen Tensor Cores with FP4 support. https://resources.nvidia.com/en-us-blackwell-architecture ↩

10. W. Hwu, D. Kirk, I. El Hajj, Programming Massively Parallel Processors: A Hands-on Approach, 4th ed., Morgan Kaufmann, 2022 — Chapters 1–4 (heterogeneous computing, data-parallel execution, scalable execution, memory architecture). ↩