GPU architecture: SMs, warps, the memory hierarchy and tensor cores

Before this lesson: the matrix product as dot products (linear algebra 1) and the sizes of number formats: fp32 is 4 bytes, bf16 2 bytes, fp8 1 byte (linear algebra 13). No machine learning is needed. Time: about 50 minutes, plus the notebook.

---

Why this matters #

Every training step and every generated token is a sequence of GPU kernels. To make them faster, you must know what the hardware does with your code. This lesson explains the questions that come up again and again in ML systems work:

• Why does a matrix multiply in bf16 reach hundreds of TFLOP/s, while a softmax or a LayerNorm reaches only a small fraction of that? (Tensor cores, and memory traffic.)
• Why does FlashAttention (lesson 4) keep its tiles in shared memory, and why are fused kernels (lesson 3) faster? (HBM is far and slow; on-chip memory is close and fast.)
• Why is a small batch sometimes as slow as a larger one? (Waves and the tail effect.)
• Why does an H100 decode an LLM only about 1.6× faster than an A100, not 3×? (Memory bandwidth; lesson 2.)

1. Predict first #

2. The picture #

A CPU and a GPU #

A CPU and a GPU are built for different goals. A CPU core is built for low latency: one thread of work must finish as fast as possible. It uses large caches, branch prediction and out-of-order execution for this. A GPU is built for throughput: the total work per second counts, not the time of one piece.

	CPU (one server chip)	GPU (H100 SXM, 2022)
cores	tens of complex cores	132 SMs with 128 fp32 lanes each: 16,896 "CUDA cores"
one core is	fast, with large caches and branch prediction	simple; many threads share it
a slow memory access is handled by	caches and out-of-order execution	switching to another group of threads
good for	serial code, branches, small tasks	the same operation on millions of values

So a GPU needs many independent pieces of work. Adding two vectors of $10^8$ elements is ideal: each element is one independent piece. A loop where step $k$ needs the result of step $k-1$ is a bad fit.

How a kernel is split #

A kernel is a function that runs on the GPU in many threads at once. You launch it over a grid of thread blocks. A block has up to 1,024 threads. Each thread finds its own piece of work from its block index and its thread index. For a vector add $z = x + y$ with blocks of 256 threads, the thread with block index $i$ and thread index $j$ handles element $256\,i + j$.

The hardware places each block on one SM. The SM cuts the block into warps of 32 consecutive threads. A warp is the unit of scheduling: its 32 threads execute the same instruction at the same time, each on its own data.

Interactive figure: loading…

Waves. An SM can hold several blocks at once (up to a limit, section 3). The blocks that run at the same time form a wave. If a launch has a few more blocks than one wave can hold, the extra blocks run in a second wave on a mostly idle GPU. This is the tail effect.

Waiting for memory: latency hiding #

A load from HBM takes a long time: about 480 clock cycles in a microbenchmark of a Hopper GPU (Luo et al. 2024), that is, a few hundred nanoseconds. A warp that needs the data must wait.

The SM does not wait with it. Each SM has 4 warp schedulers. In every cycle, each scheduler picks one warp that is ready and issues its next instruction. Switching between warps costs nothing, because every resident warp keeps its registers on the SM. So while some warps wait for memory, others compute. This is latency hiding, and it is why an SM needs many resident warps.

Interactive figure: loading…

What a warp does with a branch. All 32 threads of a warp execute the same instruction. If some threads take the if branch and the others take the else branch, the warp runs both paths, one after the other, and switches off the threads that do not belong to the current path. This is warp divergence. NVIDIA calls this execution model SIMT (single instruction, multiple threads).

3. Definitions and main results #

The execution hierarchy #

Level	What it is	Size on an H100
thread	one sequence of instructions with its own registers	up to 255 registers
warp	32 consecutive threads of a block that execute one instruction together	32 threads
thread block	a group of threads that runs on one SM and can share memory	up to 1,024 threads
grid	all blocks of one kernel launch	up to billions of threads
SM (streaming multiprocessor)	a processor with its own cores, registers and shared memory	132 per GPU

Four rules follow from this design. They explain most of CUDA programming (lesson 3 uses them):

1. A block stays on one SM until it finishes. One SM can hold several blocks.
2. Blocks are independent. The hardware may run them in any order, and not all of them are resident at the same time.
3. Threads of one block can cooperate. They share fast shared memory and can wait for each other at a barrier (__syncthreads()).
4. Threads of different blocks cannot wait for each other inside a normal kernel launch. A block that waits for a block that has not started yet would wait forever (a deadlock).

Latency hiding and Little's law #

How many warps are enough? Take a simple model of one warp scheduler. Each warp issues $c$ instructions (one per cycle), then waits $L$ cycles for a load, and repeats. In one period of $c + L$ cycles, one warp keeps the scheduler busy for $c$ cycles. With $k$ warps:

For $c = 10$ and $L = 500$: one warp gives $10/510 \approx 2\%$, and you need $510/10 = 51$ warps. One scheduler can hold at most 16 warps (64 per SM, 4 schedulers), so this loop cannot hide the latency: it needs more work per load (larger $c$), or several loads in flight per warp.

The same idea, for bytes, is Little's law from queueing theory:

Why it is true. Bytes arrive at rate BW. Each byte stays "in flight" (requested but not yet delivered) for the latency $L$. So at any moment, the bytes in flight are the ones requested in the last $L$ seconds: $\mathrm{BW} \cdot L$.

Numbers for an H100. Take BW $= 3.35$ TB/s and a latency of 400 ns (a round number for a busy memory system; the idle latency above is about 270 ns at 1.8 GHz, and latency grows under load):

Over 132 SMs this is about 10 KB per SM. One warp-wide load of one fp32 value per thread moves $32 \times 4 = 128$ bytes, so each SM needs about $10{,}150 / 128 \approx 80$ such loads in flight at all times. One warp with one load in flight is far from enough. You need many resident warps, or several independent loads per thread (unrolled loops, 16-byte vector loads).

Occupancy #

Occupancy is the number of resident warps on an SM divided by the maximum (64 on an H100). The resident blocks share the SM's resources, and the tightest resource sets the limit (NVIDIA Hopper tuning guide):

Resource per SM (H100)	Limit
32-bit registers	65,536 (at most 255 per thread)
shared memory	228 KB (at most 227 KB per block)
resident threads (warps)	2,048 (64)
resident blocks	32

For a kernel with $t$ threads per block, $r$ registers per thread and $s$ KB of shared memory per block, the number of resident blocks is

Example. Blocks of 256 threads, 64 registers per thread, 40 KB of shared memory per block.

1. Registers: $65{,}536 / (256 \cdot 64) = 4$ blocks.
2. Shared memory: $\lfloor 228 / 40 \rfloor = 5$ blocks.
3. Threads: $2{,}048 / 256 = 8$ blocks.
4. Block limit: 32.
5. The minimum is 4 blocks, limited by registers: $4 \times 8 = 32$ warps, so occupancy is $32/64 = 50\%$.

Interactive figure: loading…

The formula is a small simplification. The hardware allocates registers in chunks per warp and reserves 1 KB of shared memory per block; NVIDIA's occupancy calculator (in Nsight Compute) includes these details. They rarely change the answer by more than one block.

High occupancy is a means, not a goal. What hides latency is enough independent work in flight. A tuned matrix-multiply kernel often runs at low occupancy: each thread keeps many independent operations in flight from its many registers.

The memory hierarchy #

Each level is larger and slower than the one above it. H100 SXM:

Level	Shared by	Size	Bandwidth	Latency (cycles)
registers	one thread	256 KB per SM (65,536 × 4 B); 34.6 MB in total	feeds the cores every cycle	part of the instruction
shared memory / L1	one block (shared memory), one SM (L1)	256 KB per SM (34.6 MB in total), up to 228 KB of it as shared memory	128 B per cycle per SM, ≈ 30 TB/s in total	≈ 30
L2 cache	all SMs	50 MB	≈ 5.5 TB/s (measured)	≈ 260
HBM3	all SMs	80 GB	3.35 TB/s	≈ 480

For on-chip memory, KB means 1,024 bytes, as in NVIDIA's documents. The 30 TB/s assumes a clock of about 1.8 GHz. NVIDIA does not publish the L2 bandwidth; 5.5 TB/s is a measured value (Austin et al., How to Scale Your Model). The latencies were measured on an H800 PCIe, a Hopper GPU with the same SM design (Luo et al. 2024).

Interactive figure: loading…

The key fact: only HBM is large enough to hold weights, activations and optimizer states. So nearly every kernel reads its inputs from HBM and writes its outputs back to HBM. Those bytes are the "bytes" of the roofline model in lesson 2. Shared memory and registers are fast, but a kernel can use them only for the data it is working on right now. Tiling, fusion and FlashAttention (lessons 3–4) are ways to do more work on that data before it goes back to HBM.

Tensor cores #

A tensor core computes a small matrix multiply-accumulate $D = AB + C$ in one instruction. The inputs $A, B$ are in a low-precision format (tf32, bf16, fp16, fp8 on Hopper, or int8). The accumulator $C, D$ is in fp32 (int32 for int8).

Example. The warp-level instruction mma.sync.m16n8k16 with bf16 inputs multiplies a $16 \times 16$ tile $A$ by a $16 \times 8$ tile $B$ and adds the product to a $16 \times 8$ fp32 tile $C$. Each of the $16 \cdot 8 = 128$ outputs is a dot product of length 16: 16 multiplies and 16 adds. So one instruction does $2 \cdot 16 \cdot 8 \cdot 16 = 4{,}096$ FLOPs. (Hopper also has larger wgmma instructions, issued by a group of 4 warps together.)

The gap. An H100 SXM reaches 989 TFLOP/s dense bf16 on its tensor cores, but 67 TFLOP/s in fp32 on its CUDA cores:

So work that is not a matrix multiply (softmax, normalisation, activations, optimizer updates) is about an order of magnitude slower per FLOP. NVIDIA's product page lists 1,979 TFLOP/s for bf16; that number assumes 2:4 structured sparsity. Dense work gets half: 989. This track always uses dense peaks.

Where the peak numbers come from

A peak is (number of units) × (FLOPs per unit per cycle) × (clock). For the fp32 CUDA cores of an H100 SXM: $16{,}896$ lanes $\times\ 2$ FLOPs (one fused multiply-add) $\times\ 1.98$ GHz (boost clock) $= 66.9$ TFLOP/s. This is the 67 of the datasheet. Real kernels run at the clock the GPU can hold under load, which may be lower, and they never issue a useful FMA in every cycle. A large, well-tuned bf16 matrix multiply usually reaches about 70–80% of the datasheet peak.

Two resources per kernel #

Put the pieces together. A kernel uses two resources:

• arithmetic: FLOPs on the tensor cores (fast, only for matrix multiplies) or the CUDA cores;
• data movement: bytes between HBM and the SMs, at most 3.35 TB/s on an H100.

The slower of the two sets the time. Lesson 2 turns this into the roofline model.

Answers to section 1. (1) About 2%. One warp is busy for $c = 10$ of every $c + L = 510$ cycles: $10/510 \approx 2\%$. You need about 51 warps' worth of independent work to fill the gaps. (2) About 2 ms. 132 blocks fill one wave. The 133rd block runs in a second wave, alone on one SM, while 131 SMs are idle. (3) No. exp is not a matrix multiply, so it runs on the CUDA cores, not the tensor cores. And it does very little arithmetic per byte it reads and writes, so it waits for HBM almost all the time (lesson 2 makes this precise).

4. Worked example #

One kernel from launch to finish. Add two fp32 vectors, $z = x + y$, with $n = 2^{24} = 16{,}777{,}216$ elements, on an H100. Use blocks of 256 threads, one element per thread. Each thread uses few registers (say 16) and no shared memory.

Step 1: the grid. Blocks $= \lceil n / 256 \rceil = 2^{24}/2^8 = 65{,}536$. Each block has $256/32 = 8$ warps, so the grid has $65{,}536 \times 8 = 524{,}288$ warps. $n$ is a multiple of 256, so no thread is idle in the last block.

Step 2: occupancy. Registers allow $\lfloor 65{,}536 / (256 \cdot 16) \rfloor = 16$ blocks; shared memory sets no limit; threads allow $2{,}048/256 = 8$; the block limit is 32. So 8 blocks are resident per SM: 64 warps, occupancy 100%.

Step 3: waves. One wave holds $132 \times 8 = 1{,}056$ blocks. Then $65{,}536 / 1{,}056 = 62.06$, so the launch needs 63 waves. The last wave has $65{,}536 - 62 \times 1{,}056 = 64$ blocks. With 63 waves, the tail costs little (at most 1 wave in 63).

Step 4: enough bytes in flight? Each thread issues two independent loads ($x_i$ and $y_i$) before it needs either. Per SM: $64$ warps $\times\ 2$ loads $\times\ 128$ B $= 16{,}384$ B $\approx 16$ KB in flight. Little's law asked for about 10 KB per SM. So this kernel can keep HBM busy.

Step 5: time. The kernel reads $x$ and $y$ and writes $z$: $3 \times 4 \times 2^{24} = 201$ MB. At 3.35 TB/s this takes at least $201 \times 10^6 / 3.35 \times 10^{12} = 60$ µs. The arithmetic is $n = 1.68 \times 10^7$ additions; even on the CUDA cores (67 TFLOP/s) this takes 0.25 µs. The time is set by data movement, 240 times over. Lesson 2 calls such a kernel memory-bound.

5. In code #

These lines recompute the worked example and the occupancy example. The notebook turns each of them into an exercise.

import math

SMS, MAX_WARPS, REGS, SMEM_KB, MAX_THREADS, MAX_BLOCKS = 132, 64, 65_536, 228, 2_048, 32

def occupancy(threads, regs, smem_kb):
    """Resident blocks per SM and occupancy on an H100 (simple model)."""
    limits = {"registers": REGS // (threads * regs),
              "shared memory": SMEM_KB // smem_kb if smem_kb > 0 else math.inf,
              "threads": MAX_THREADS // threads,
              "blocks": MAX_BLOCKS}
    blocks = min(limits.values())
    return blocks, blocks * threads // 32 / MAX_WARPS, min(limits, key=limits.get)

print(occupancy(256, 64, 40))   # (4, 0.5, 'registers')
print(occupancy(128, 32, 64))   # (3, 0.1875, 'shared memory')

# Little's law: bytes in flight = bandwidth x latency
in_flight = 3.35e12 * 400e-9
print(in_flight, in_flight / SMS, in_flight / SMS / 128)   # 1340000.0 10151.5... 79.3...

# The worked example: vector add of n = 2**24 fp32 values, blocks of 256 threads
n, t = 2**24, 256
blocks = math.ceil(n / t)
per_sm, _, _ = occupancy(t, 16, 0)
per_wave = SMS * per_sm
print(blocks, blocks * t // 32, per_sm, math.ceil(blocks / per_wave))   # 65536 524288 8 63
print(3 * 4 * n / 3.35e12 * 1e6, n / 67e12 * 1e6)   # 60.09... µs (bytes) vs 0.25... µs (FLOPs)

6. Where it shows up #

• Every PyTorch operation on a GPU launches one or more kernels. torch.matmul calls cuBLAS kernels that use tensor cores; an elementwise op like x * 2 launches a simple kernel like the vector add of section 4.
• FlashAttention (lesson 4) keeps tiles of $Q$, $K$ and $V$ in shared memory, so the large attention matrix never goes to HBM.
• Wave quantization. NVIDIA's performance guide warns that a matrix multiply whose number of output tiles is just above a multiple of the SM count wastes most of a wave. A tiny change of a layer width can make a kernel noticeably faster or slower.
• LLM decoding reads every weight from HBM for a few tokens, so it is limited by HBM bandwidth (lesson 2, lesson 12).
• ML for science. Stencil codes for weather and climate grids and sparse solvers do little arithmetic per byte. On GPUs they are limited by HBM bandwidth, like the vector add, not by the tensor cores.

7. Common mistakes #

• Thinking a warp's threads run independently. They share one instruction stream. Divergent branches run one after the other, so the cost is the sum of both paths.
• Thinking all blocks run at the same time. Only one wave of blocks is resident; the rest wait. Code must not assume an order of blocks or wait for another block.
• Maximising occupancy for its own sake. Occupancy is a means to hide latency. Enough independent work in flight is the goal; tuned matmul kernels are fast at low occupancy.
• Using the sparse peak. Product pages list bf16 with 2:4 sparsity (1,979 TFLOP/s on an H100). Dense work gets at most half (989).
• Expecting tensor-core speed for every FLOP. Only matrix multiplies (and convolutions written as matrix multiplies) use tensor cores. Elementwise ops, reductions and softmax use the CUDA cores, about 15× slower per FLOP.
• Forgetting that data lives in HBM. Shared memory and registers are fast, but the data must first come from HBM and go back to HBM. Those bytes usually set the time.

8. Check yourself #

1. A kernel uses 255 registers per thread and blocks of 128 threads. How many warps can be resident on one H100 SM, and what is the occupancy?
   Answer

   Registers allow $\lfloor 65{,}536/(128 \cdot 255) \rfloor = \lfloor 2.008 \rfloor = 2$ blocks. Shared memory sets no limit, threads allow 16 blocks. So 2 blocks of 4 warps: 8 warps, occupancy $8/64 = 12.5\%$. (With per-warp register chunks, 255 rounds up to 256 registers per thread; the answer is the same.)
2. An A100 has 2.039 TB/s of HBM bandwidth and 108 SMs. With a latency of 400 ns, how many bytes must be in flight in total and per SM?
   Answer

   $2.039 \times 10^{12} \times 400 \times 10^{-9} = 816$ KB in total, and $816{,}000/108 \approx 7.6$ KB per SM.
3. A kernel launches 264 blocks on an H100, and only one block fits on an SM at a time. Another launch has 265 blocks. Compare their times.
   Answer

   264 blocks are exactly 2 waves of 132. 265 blocks need 3 waves, and the third wave has one block. If every block takes the same time, the second launch takes about 1.5× as long.
4. How many FLOPs does one mma.sync.m16n8k16 instruction do, and why?
   Answer

   $2 \cdot 16 \cdot 8 \cdot 16 = 4{,}096$. The output tile has $16 \cdot 8 = 128$ entries; each is a dot product of length 16, which is 16 multiplies and 16 adds.
5. Why does a GPU switch between warps at no cost, while a CPU context switch is expensive?
   Answer

   Every resident warp keeps its registers on the SM, so the scheduler only has to pick a different warp in the next cycle. A CPU context switch must save and restore the registers of a thread in memory.
6. A kernel reads 1 GB and writes 1 GB on an H100 and does $10^9$ FLOPs on the CUDA cores. Which resource sets its time, and what is the minimum time?
   Answer

   Data: $2 \times 10^9 / 3.35 \times 10^{12} = 597$ µs. Arithmetic: $10^9 / 67 \times 10^{12} = 15$ µs. Data movement sets the time: at least about 0.6 ms.

Recall cardsAnswer in your head first, then reveal. Add the cards to your review deck to see them again just before you would forget.

1. What is a warp?

   Show answer

2. What is an SM, and how many does an H100 SXM have?

   Show answer

3. What happens when the threads of one warp take different branches of an if/else?

   Show answer

4. State Little's law for memory.

   Show answer

5. Why does an SM need many resident warps?

   Show answer

6. Define occupancy and name its limits on an H100 SM.

   Show answer

7. Compute the occupancy for blocks of 256 threads, 64 registers per thread and 40 KB of shared memory on an H100.

   Show answer

8. Give the HBM size and bandwidth of an H100 SXM.

   Show answer

9. What does a tensor core compute?

   Show answer

10. How much faster are H100 tensor cores (bf16) than its CUDA cores (fp32)?

    Show answer

Add these cards to my reviewOpen the review →

9. Interview angle #

• "What is a warp, and what is warp divergence?" — 32 threads that execute one instruction together. If they branch differently, the warp runs both paths in sequence, so divergent code can cost the sum of the paths.
• "Why does a GPU need so many threads?" — To hide memory latency: HBM loads take hundreds of cycles, and an SM covers them by switching between resident warps. Little's law gives the amount of data that must be in flight, about 1.3 MB for an H100.
• "What limits occupancy, and is higher occupancy always better?" — Registers, shared memory, threads and blocks per SM. No: it only has to be high enough to hide latency; register-heavy matmul kernels are fast at low occupancy.
• "Why is a softmax so much slower than a matmul per FLOP?" — It does not use tensor cores (about 15× fewer FLOP/s), and it does few FLOPs per byte, so it waits for HBM.
• "Can two blocks synchronise?" — Not inside a normal launch: blocks are independent and may not be resident together. Use a second kernel launch, a cooperative launch, or (Hopper) a thread block cluster.

10. Practice #

The notebook for this lesson runs in Colab, with NumPy only. It has four exercises, each with a check cell and a folded solution:

1. Write occupancy(threads, regs, smem_kb): resident blocks, occupancy and the limiting resource, checked on the lesson's examples and edge cases (a launch that does not fit).
2. Little's law: the bytes in flight for a given bandwidth and latency, and the loads per SM.
3. A warp-divergence cost model: the time of each warp for a given branch pattern, and why sorting the data can remove divergence.
4. Waves and the tail effect: the number of waves of a grid, and the fraction of the GPU used.

11. Watch and read #

Getting Started with CUDA and Parallel Programming | NVIDIA GTC 2025 SessionNVIDIA Developer · GTC 2025 (a CUDA architect) · 42 min

How a GPU differs from a CPU, why it needs so many threads, and how memory latency shapes the design.

Lecture 4 Compute and Memory BasicsThomas Viehmann · GPU MODE lecture 4 · 57 min

Warps, scheduling, occupancy and the memory types of an SM, with CUDA examples.

• Stanford CS336 (2025), lecture 5: GPUs (on YouTube only; it cannot be embedded here) and the course page.
• Austin et al., How to Scale Your Model, How to Think About GPUs: SMs, the memory hierarchy and a table of GPU specs.
• NVIDIA, Hopper tuning guide (occupancy limits, shared memory) and the CUDA C++ programming guide (the programming model: grids, blocks, warps).
• NVIDIA, GPU performance background: a short guide to SMs, tensor cores and waves for deep learning.
• Luo et al. (2024), Benchmarking and Dissecting the Nvidia Hopper GPU Architecture: the latency and bandwidth measurements used in this lesson.
• Modal, GPU Glossary: short definitions of every term in this lesson.
• Sasha Rush, GPU Puzzles, puzzles 1–8: threads, guards, blocks and shared memory, in Python, in Colab.