Learn: The GPU Memory Hierarchy

What you're learning

Where every byte lives on a GPU, how fast you can reach it, and why moving data — not doing math — is what actually bounds most LLM kernels. By the end you should be able to look at a kernel, name which level of the hierarchy dominates its runtime, and explain why FlashAttention exists.

This lesson assumes 01-gpu-architecture-and-simt — SMs, warps, the SIMT execution model, and occupancy. We now zoom in on the thing those warps are constantly waiting on: memory.

1. Learning map

graph TD
    A["Lesson 01<br/>SMs, warps, occupancy"] --> B["The hierarchy<br/>registers → SMEM/L1 → L2 → HBM → host"]
    B --> C["Registers<br/>per-thread, ~1 cycle"]
    B --> D["Shared memory / L1<br/>per-SM scratchpad, banked"]
    B --> E["L2 cache<br/>GPU-wide"]
    B --> F["Global memory / HBM<br/>the bandwidth wall"]
    C --> C1["Register pressure<br/>→ spills → occupancy"]
    D --> D1["Bank conflicts<br/>→ serialization → padding"]
    F --> F1["Coalescing<br/>32 threads → sectors/cache lines"]
    C1 --> G["The memory wall<br/>FLOPs ≫ bandwidth growth"]
    D1 --> G
    F1 --> G
    G --> H["Most LLM kernels are<br/>memory-bound"]
    H --> I["→ Lesson 05: GEMM tiling"]
    H --> J["→ Lesson 06: FlashAttention"]
    H --> K["→ Lesson 03: Roofline"]

    style B fill:#44a,color:#fff
    style G fill:#a44,color:#fff
    style H fill:#4a4,color:#fff

2. Why this matters

Here is the single most important number in GPU performance engineering. An H100 SXM5 delivers roughly 990 TFLOP/s of dense BF16/FP16 tensor-core throughput and reads from HBM3 at roughly 3.35 TB/s.¹ Divide them:

I_{ridge} = \frac{990 \times 1 0 ^{12} FLOP/s}{3.35 \times 1 0 ^{12} B/s} \approx 295 \frac{FLOP}{byte}

This is the ridge point of the roofline (lesson 03). It says: for every byte you pull from HBM, you must do ~300 floating-point operations just to keep the tensor cores fed. If your kernel does fewer FLOPs per byte than that, it is memory-bound — the tensor cores sit idle waiting for data, and buying more compute buys you nothing.

Most of the operations in a transformer — elementwise activations, LayerNorm, softmax, residual adds, and even attention at inference — fall well below 295 FLOP/byte. They are limited by how fast you can move bytes, not how fast you can multiply. So understanding where bytes live and how to move them cheaply is the whole game.

The one-sentence version

On modern accelerators, a FLOP is nearly free and a byte from HBM is expensive — in both time and energy. Performance engineering is mostly the art of not touching HBM.

3. The hierarchy, top to bottom

Every GPU has a memory hierarchy that trades capacity for speed: the fastest storage is tiny and private, the largest is slow and shared. All numbers below are for the H100 (Hopper, SXM5);¹² the shape is the same across NVIDIA generations, only the constants move.

graph TD
    R["<b>Registers</b> — per thread<br/>~256 KB/SM (65,536 × 32-bit)<br/>latency ~1 cycle · BW ~tens of TB/s per SM"]
    S["<b>Shared memory / L1</b> — per SM<br/>up to 228 KB SMEM (256 KB unified L1+SMEM)<br/>latency ~20–30 cycles · BW ~19 TB/s aggregate"]
    L2["<b>L2 cache</b> — GPU-wide<br/>50 MB<br/>latency ~150–250 cycles · BW ~several TB/s"]
    H["<b>Global memory / HBM3</b> — GPU-wide<br/>80 GB<br/>latency ~450–800 cycles · BW ~3.35 TB/s"]
    P["<b>Host DRAM</b> — over PCIe / NVLink<br/>PCIe 5.0 ~64 GB/s · NVLink 4 ~900 GB/s<br/>latency ~µs (thousands of cycles)"]

    R -->|spill| S
    S --> L2
    L2 --> H
    H -->|H2D / D2H| P

    style R fill:#1a7,color:#fff
    style S fill:#2a6,color:#fff
    style L2 fill:#a83,color:#fff
    style H fill:#a44,color:#fff
    style P fill:#666,color:#fff

Level	Scope	Capacity (H100)	Latency	Bandwidth	Managed by
Registers	per-thread	256 KB/SM (64K × 4B)³	~1 cycle	~tens of TB/s per SM⁴	compiler
Shared mem / L1	per-SM	up to 228 KB SMEM (256 KB unified)⁵	~20–30 cycles⁴	~19 TB/s aggregate⁴	programmer (SMEM) / HW (L1)
L2 cache	GPU-wide	50 MB²	~150–250 cycles⁴	~several TB/s	hardware
Global / HBM3	GPU-wide	80 GB¹	~450–800 cycles⁴	3.35 TB/s¹	programmer (allocations)
Host DRAM	system	100s of GB	~µs	PCIe5 ~64 GB/s, NVLink4 ~900 GB/s¹	programmer (copies)

Note the two enormous cliffs. Going register → SMEM is ~20–30×. Going SMEM → HBM is another ~20–40× in latency and ~150× in bandwidth. The entire discipline of GPU kernel optimization is: pull data down the pyramid once, reuse it as many times as possible up top, and write it back once.

Latency is hidden, bandwidth is not

A single HBM load takes ~500 cycles, but the SIMT model (lesson 01) hides that latency by swapping in other ready warps while one warp stalls — this is what occupancy buys you. What occupancy cannot fix is bandwidth: if a kernel must move $N$ bytes through a 3.35 TB/s pipe, no amount of parallelism moves them faster. Latency is a scheduling problem; bandwidth is a physics problem.

4. Registers and register pressure

Registers are the only storage a thread reads/writes at full ALU speed with ~1-cycle latency. Each SM has a fixed register file — on H100, 65,536 32-bit registers per SM (256 KB).³ This file is partitioned across all resident threads, which is the direct link back to occupancy from lesson 01.

If a kernel needs $R$ registers per thread, the number of threads that can be resident on one SM is bounded by:

threads_{resident} \leq \frac{65536}{R}

Two examples on H100 (max 2048 resident threads/SM):³

$R = 32$ : $65536/32 = 2048$ threads — register file does not limit occupancy.
$R = 64$ : $65536/64 = 1024$ threads — you’re capped at 50% occupancy by registers alone.
$R = 128$ : $512$ threads — 25% occupancy.

When a kernel needs more registers than the compiler can allocate (either it hit the per-thread cap of 255,³ or you forced higher occupancy with __launch_bounds__), the excess variables spill to “local memory.”

"Local memory" is a lie — it's global memory

Despite the name, local memory is not close to the thread. It is a per-thread private region that physically lives in HBM (cached in L1/L2).³ A register spill therefore turns a ~1-cycle access into a potential ~500-cycle HBM round trip. A spilling inner loop can be catastrophically slow. Check ptxas -v (or Nsight Compute) for spill stores/spill loads — nonzero counts in a hot kernel are a red flag.

There is a genuine tension here: more registers per thread → more values kept in fast storage and fewer recomputes, but → lower occupancy → less latency hiding. The optimum is workload-dependent and is exactly the kind of thing you sweep experimentally.

Quiz: A kernel uses 96 registers/thread and 512 threads/block on H100 (65,536 regs/SM, 2048 max threads/SM). Is occupancy limited by registers?

Answer

Registers needed per block: $96 \times 512 = 49, 152$ . Only $⌊ 65536/49152 ⌋ = 1$ block fits per SM by register budget → 512 resident threads → $512/2048 = 25%$ occupancy from registers. Two blocks would need 98,304 regs > 65,536, so they don’t fit. Registers are the binding constraint here; dropping to ~64 regs/thread would let 2 blocks (1024 threads) co-reside.

5. Shared memory: the programmer-managed scratchpad

Shared memory (SMEM) is on-chip SRAM, physically the same silicon as L1 — on H100 they share a 256 KB block per SM that you split (e.g. 228 KB SMEM / 28 KB L1, configurable via cudaFuncAttributePreferredSharedMemoryCarveout).⁵³ The crucial distinction:

L1 is hardware-managed cache — you don’t control what’s in it.
SMEM is a software-managed scratchpad — you explicitly stage data into it (__shared__ arrays) and it stays until you’re done. It is shared by all threads in a block, making it the primary vehicle for inter-thread cooperation and data reuse.

The canonical use: load a tile of data from HBM into SMEM once, then have all threads in the block read it many times from SMEM instead of hammering HBM. This is the heart of tiled GEMM (lesson 05).

Banks and bank conflicts

To sustain its huge bandwidth, SMEM is split into 32 banks, one per lane of a warp. Banks are 4 bytes (32 bits) wide and interleaved by word:³

bank (addr) = (\frac{addr}{4}) mod 32

The hardware can service one access per bank per cycle. So a warp’s 32 accesses complete in a single transaction if and only if they hit 32 distinct banks (or all hit the same address, which is a broadcast).³⁶ If $k$ threads in a warp hit different addresses in the same bank, those accesses serialize into $k$ transactions — a $k$ -way bank conflict costing $k \times$ the cycles.

graph LR
    subgraph "No conflict: stride 1"
        T0["t0→bank0"]:::ok
        T1["t1→bank1"]:::ok
        T2["t2→bank2"]:::ok
        Tn["...→bank31"]:::ok
    end
    subgraph "2-way conflict: stride 2"
        A0["t0→bank0"]:::bad
        A1["t1→bank2"]:::bad
        A16["t16→bank0"]:::bad
    end
    classDef ok fill:#2a6,color:#fff
    classDef bad fill:#a44,color:#fff

The classic offender is a column access of a 2D SMEM array whose row width is a multiple of 32. Consider float tile[32][32];. Threads reading a column tile[i][c] for fixed c, i = 0..31 access addresses c, 32+c, 64+c, ... — all of which map to bank c mod 32 → a 32-way conflict, a 32× slowdown.

Column-major access into a 32-wide SMEM tile is a 32-way bank conflict

Any time consecutive threads step through SMEM with a stride that shares a common factor with 32, you serialize. Diagonal/transpose patterns are the usual culprits (matrix transpose kernels, attention score tiles).

The fix: padding. Declare the array one element wider than needed — float tile[32][33];. Now row $i$ starts at offset $33 i$ , so column accesses land on addresses $c, 33 + c, 66 + c, \dots$ whose banks are $(33 i + c) mod 32 = (i + c) mod 32$ — all distinct. One wasted column of SMEM buys conflict-free column access. This padding trick is ubiquitous in transpose and GEMM kernels.

Quiz: Why does padding a [32][32] SMEM tile to [32][33] eliminate the column-access bank conflict?

Answer

With width 32, element tile[i][c] is at word offset $32 i + c$ , so its bank is $(32 i + c) mod 32 = c mod 32$ — every row in a column maps to the same bank $c$ → 32-way conflict. With width 33, the offset is $33 i + c$ and the bank is $(33 i + c) mod 32 = (i + c) mod 32$ . As $i$ runs $0..31$ , $(i + c) mod 32$ takes 32 distinct values → each thread hits a different bank → conflict-free. The extra column shifts each row’s bank alignment by one, breaking the alignment that caused the collision. Cost: $32 \times 4 = 128$ bytes of SMEM per tile.

6. Global memory and coalescing

Global memory (HBM) is where your tensors live. Its latency is ~500 cycles and — decisively — its bandwidth is the scarcest resource on the chip. The hardware does not fetch individual bytes; it moves data in fixed-size sectors of 32 bytes (and 128-byte cache lines = 4 sectors).³⁶ How your warp’s 32 threads map onto these sectors determines your effective bandwidth.

Coalescing is the hardware combining the 32 addresses a warp issues in one memory instruction into the minimum number of sector transactions.⁶

Coalesced (ideal): the 32 threads of a warp read 32 consecutive 4-byte words → 128 contiguous bytes → exactly one 128-byte cache line = 4 sectors, all fully used. 100% of fetched bytes are useful.

addr (t) = base + 4 t, t = 0 \dots 31 \Rightarrow 1 transaction, 128 B, 0 waste

Strided: if thread $t$ reads base + 4 * stride * t, consecutive threads land in different sectors. With stride = 32, all 32 threads touch 32 different 128-byte lines. You fetch $32 \times 128 = 4096$ bytes to deliver the $32 \times 4 = 128$ bytes you actually wanted:

efficiency = \frac{128}{4096} = \frac{1}{32} \approx 3.1%

That is a 32× effective-bandwidth loss from address pattern alone — the DRAM pins move at full 3.35 TB/s, but 97% of the bytes crossing them are thrown away.

graph TD
    subgraph "Coalesced — 1 transaction"
        direction LR
        C["warp reads addr 0,4,8,...,124"] --> CL["1× 128B cache line<br/>100% useful"]
    end
    subgraph "Strided (32) — 32 transactions"
        direction LR
        S["warp reads addr 0,128,256,..."] --> SL["32× 128B lines<br/>3% useful"]
    end
    style CL fill:#2a6,color:#fff
    style SL fill:#a44,color:#fff

Uncoalesced access is the #1 beginner GPU performance bug

Row-major arrays where consecutive threads map to consecutive rows (not columns), Array-of-Structs layouts, and gather/scatter with random indices all destroy coalescing. The fix is usually a data-layout change: Struct-of-Arrays instead of Array-of-Structs, transposing which index the thread ID maps to, or staging through SMEM so the global access pattern is coalesced even if the SMEM pattern isn’t. Always make threadIdx.x (the fastest-varying thread index) map to the fastest-varying (contiguous) memory dimension.

Quiz: A warp of 32 threads each loads one float (4 B). Pattern A: thread $t$ reads element $t$ . Pattern B: thread $t$ reads element $t \times 16$ . Cache-line = 128 B. Bytes fetched and efficiency for each?

Answer

Pattern A (coalesced): addresses span $0..124$ B = one 128 B line. Fetched: 128 B; useful: 128 B; efficiency 100%, 1 transaction.
Pattern B (stride 16): thread $t$ at byte $64 t$ , spanning $0$ to $64 \times 31 = 1984$ B. Two adjacent threads (64 B apart) share a 128 B line, so pairs of threads share lines → about $⌈ 1984/128 ⌉ + 1 \approx 16$ distinct 128 B lines fetched $= 16 \times 128 = 2048$ B for $32 \times 4 = 128$ B wanted. Efficiency $\approx 128/2048 = 6.25%$ , ~16 transactions. Every doubling of stride roughly halves efficiency until you hit one line per thread (the 1/32 floor).

7. The memory wall

Here is the historical trend that makes all of the above matter more every year. Peak compute has grown far faster than memory bandwidth:

GPU (flagship)	Year	Peak FP16/BF16 TC	HBM BW	Ridge (FLOP/byte)
P100	2016	~21 TFLOP/s (FP16)⁷	720 GB/s	~29
V100	2017	125 TFLOP/s⁸	900 GB/s	~139
A100	2020	312 TFLOP/s⁹	2.0 TB/s	~156
H100	2022	990 TFLOP/s¹	3.35 TB/s	~295
B200	2024/25	~2250 TFLOP/s (dense)¹⁰	~8 TB/s	~280

Over P100→H100, compute grew ~47× while bandwidth grew ~4.6×. The ridge point marches right: each generation, a kernel must do more arithmetic per byte to stay compute-bound. Ops that were compute-bound on older hardware silently become memory-bound on newer hardware. This gap — compute scaling faster than the memory feeding it — is the memory wall, and it is why “just add more FLOPs” stopped being a strategy.

The corollary for LLM engineers

As accelerators get faster, a growing fraction of your workload becomes bandwidth-limited. The lever that keeps mattering is arithmetic intensity (FLOPs per byte moved) — and you raise it by reuse (tiling, fusion), not by faster math.

Data movement dominates energy, too

It’s not just time. The energy to move a word from HBM dwarfs the energy of the arithmetic you do with it. Rough, order-of-magnitude figures (process-dependent, but the ratios are the point):

Operation	Energy (approx.)¹¹
FP16/BF16 multiply-add (on-chip)	~1 pJ
Read from register / SMEM	~1–10 pJ
Read from L2	~20–100 pJ
Read a word from HBM/DRAM	~several hundred pJ (100–1000× a FLOP)
Move a word host↔device over PCIe	~nJ (1000s× a FLOP)

So an off-chip access costs 2–3 orders of magnitude more energy than the FLOP it feeds. At datacenter scale this sets the power bill and the thermal envelope: minimizing HBM traffic is simultaneously a speed, power, and cost optimization. This is the deep reason “IO-aware” is the design principle behind modern kernels.

Quiz: An elementwise kernel reads a BF16 tensor, multiplies each element by a scalar, and writes it back. What is its arithmetic intensity, and is it compute- or memory-bound on an H100? What does this imply for a chain of such ops (e.g. bias → GELU → dropout)?

Answer

Intensity. Per element: read 2 bytes + write 2 bytes = 4 bytes moved, for 1 FLOP (the multiply). So $I = 1/4 = 0.25$ FLOP/byte. (Even counting a fused multiply-add as 2 FLOPs, $I \leq 0.5$ .)
Classification. H100’s ridge point is ~295 FLOP/byte (BF16). At $I = 0.25 ≪ 295$ , this kernel sits far out in the memory-bound region — it will run at essentially HBM bandwidth, using a tiny fraction of the tensor cores. Its runtime is $\approx \frac{4 \times N}{3.35 TB/s}$ , independent of how much math you add per element.
Implication — this is the case for fusion. If bias, GELU, and dropout are three separate kernels, you pay the full HBM round-trip (read + write the whole tensor) three times — $\sim$ 12 bytes/element of traffic for ~a few FLOPs. Fusing them into one kernel reads once and writes once ( $\sim$ 4 bytes/element), cutting traffic ~3× and giving a near-3× speedup, because the op is purely bandwidth-bound. This is exactly why frameworks fuse elementwise chains (and why FlashAttention fuses the whole attention op — lesson 06).

8. GPU/LLM connection

Now tie it to transformers directly.

Attention is memory-bandwidth-bound. Standard attention computes $S = Q K^{⊤}$ , $P = softmax (S)$ , $O = P V$ . The score/probability matrices $S, P \in R^{N \times N}$ are $O (N^{2})$ in size. Naively materializing them means writing $S$ to HBM, reading it back for softmax, writing $P$ , reading it back for the $P V$ matmul. For sequence length $N$ and head dim $d$ , you move $O (N^{2})$ bytes but do only $O (N^{2} d)$ FLOPs — arithmetic intensity $\sim d$ (tens to ~128), far below the ~295 ridge. The kernel spends its life waiting on HBM, and the $N^{2}$ HBM traffic is also the memory-capacity bottleneck for long context.

FlashAttention (→ lesson 06) is the direct payoff of everything in this lesson. It never materializes $S$ or $P$ in HBM. Instead it tiles $Q, K, V$ into blocks that fit in SMEM, computes attention block-by-block with an online softmax (running max + rescaling), and keeps the running output in registers/SMEM. HBM traffic drops from $O (N^{2})$ to $O (N^{2} d / M)$ where $M$ is the SMEM size — turning a memory-bound kernel into a compute-bound one and yielding large wall-clock speedups.¹² It is the archetypal IO-aware / kernel-fusion algorithm: same math, dramatically less data movement.

Fusion in general. Elementwise ops (bias, GELU, dropout, residual add, LayerNorm) each have arithmetic intensity near 1 — pure bandwidth. Run them as separate kernels and each pays a full HBM read+write of the activation tensor. Fuse them into one kernel and the intermediate never leaves registers/SMEM: you read once, do all the math, write once. This is why torch.compile, Triton, and hand-written fused kernels win — they cut HBM round-trips, not FLOPs.

GEMM tiling (→ lesson 05) is where §5 and §6 combine. High-performance matmul (a) loads tiles of $A$ and $B$ from HBM into SMEM with fully coalesced global reads, (b) reuses each loaded element across an entire tile of the output (raising arithmetic intensity from $O (1)$ to $O (tile size)$ ), and (c) arranges the SMEM tiles (with padding) to be bank-conflict-free. Coalescing + SMEM tiling + conflict-free layout is literally the recipe that takes a matmul from a few percent of peak to 80–90%+.

The unifying mental model

Every major LLM kernel optimization — FlashAttention, fused elementwise, tiled GEMM, KV-cache layout, paged attention — is, at bottom, a scheme to do more work per byte pulled from HBM. Once you see the hierarchy, you see that they’re all the same idea.

9. Practice problems

Practice problem 1: Effective bandwidth of a strided copy

Problem: On an H100 (peak HBM BW 3.35 TB/s, 32-byte sectors, 128-byte cache lines), you launch a copy kernel where each thread copies one float (4 B). Threads in a warp access with stride 8 elements (thread $t$ reads element $8 t$ ). Estimate the achievable effective bandwidth as a fraction of peak.

Worked solution:

Warp of 32 threads, thread $t$ at byte offset $32 t$ (since $8 \times 4 = 32$ B stride).

Addresses: $0, 32, 64, \dots, 32 \times 31 = 992$ B — spanning $\approx 1024$ B.

Number of 128 B cache lines touched: $⌈ 1024/128 ⌉ = 8$ lines. (Every 4 threads fall in one line: $128/32 = 4$ .)

Bytes fetched: $8 \times 128 = 1024$ B. Bytes actually used: $32 \times 4 = 128$ B.

Efficiency $= 128/1024 = 12.5%$ .

Effective useful bandwidth $\approx 0.125 \times 3.35 TB/s \approx 0.42 TB/s$ .

The DRAM interface still runs at full speed, but 87.5% of the bytes it delivers are discarded. Reducing the stride to 1 (coalesced) would recover the full ~3.35 TB/s. This is why measuring dram__bytes.sum vs. useful bytes in Nsight Compute is so revealing.

Practice problem 2: Counting bank conflicts

Problem: A block declares __shared__ float s[128]; (32 banks, 4 B each). A warp accesses s[threadIdx.x * 2] for threadIdx.x = 0..31. How many banks are hit, and what is the conflict degree? What if the access were s[threadIdx.x * 2 + 1] — does the degree change? What single change makes it conflict-free?

Worked solution:

Bank of word index $i$ is $i mod 32$ . Access indices: $0, 2, 4, \dots, 62$ .

Banks hit: ${0, 2, 4, \dots, 62} mod 32 = {0, 2, \dots, 30, 0, 2, \dots, 30}$ — the even banks ${0, 2, \dots, 30}$ , 16 distinct banks, each hit by exactly 2 threads → a 2-way bank conflict (2 transactions instead of 1).

For s[2t+1]: indices $1, 3, \dots, 63$ → odd banks ${1, 3, \dots, 31}$ , again 16 banks each hit twice → still 2-way. Adding a constant offset shifts banks but not the conflict degree; the stride is what matters.

Root cause: stride 2 shares factor 2 with 32. Fix: change the access stride to be coprime with 32 (stride 1 → conflict-free), or if a stride-2 logical pattern is required, pad the array so the physical stride becomes odd (e.g. index as s[t * 2 + t/16] style padding, or restructure so consecutive threads read consecutive words). Stride 1 is the clean answer: s[threadIdx.x] hits 32 distinct banks, 1 transaction.

10. Practice & resources

Hands-on exercises

Coalesced vs. strided copy microbenchmark. Write two CUDA kernels that copy a large float array: one where thread $t$ reads in[t] (coalesced), one where it reads in[t * STRIDE]. Time both with CUDA events, compute achieved GB/s, and plot effective bandwidth vs. STRIDE ∈ {1,2,4,8,16,32}. You should see it fall off roughly as $1/ min (stride, 32)$ and floor near 1/32 of peak. Confirm the coalesced version approaches your GPU’s spec HBM bandwidth.
Transpose kernel, with and without SMEM padding. Implement a matrix transpose that stages through a __shared__ float tile[32][32]. Measure it, then change the declaration to [32][33] and measure again. The padded version should be markedly faster — you just eliminated a 32-way bank conflict on the column write. Verify with Nsight Compute’s l1tex__data_bank_conflicts_pipe_lsu_mem_shared counter.
Nsight Compute deep-dive. Profile any kernel with ncu --set full ./a.out. Read the Memory Workload Analysis section: check global load/store efficiency (%), dram__throughput vs. peak, L2 hit rate, and shared-memory bank-conflict counts. Then check ptxas -v output (or the Nsight “Launch Statistics”) for register count and any spill loads/stores. Learn to state, for a given kernel, whether it’s bound by HBM bandwidth, L2, compute, or occupancy.
Register-pressure sweep. Take a compute kernel, add __launch_bounds__(256, N) for increasing N (forcing higher occupancy → fewer registers), and watch for the point where the compiler starts spilling (ptxas -v spill counts jump). Correlate with measured runtime — you’ll often find a non-monotonic sweet spot.

Real resources

PMPP — Programming Massively Parallel Processors, 4th ed. (Hwu, Kirk, El Hajj).¹³ The memory chapters: Ch. 5 (Memory architecture and data locality — tiling, SMEM), and the coalescing/memory-performance material in Ch. 6. The single best textbook treatment.
CUDA C++ Best Practices Guide — NVIDIA’s official doc; the Memory Optimizations chapter covers coalescing, SMEM/bank conflicts, and the L2/occupancy trade-offs with concrete guidance. https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/
GPU MODE lectures (formerly “CUDA MODE”) — community lecture series with excellent sessions on memory coalescing, shared memory, and profiling; pairs well with PMPP. https://github.com/gpu-mode/lectures
Nsight Compute documentation — how to read Memory Workload Analysis, the metric names (dram__bytes, l1tex__*, sm__*), and the roofline chart it generates. https://docs.nvidia.com/nsight-compute/
CUDA C++ Programming Guide — authoritative reference for the memory model, __shared__, bank definitions, and the L1/SMEM carveout API. https://docs.nvidia.com/cuda/cuda-c-programming-guide/
NVIDIA H100 architecture whitepaper — for the exact SM counts, register file, SMEM/L1 sizes, L2 size, and HBM3 bandwidth quoted above.

11. What’s next

Three threads, in order:

03-performance-modeling-roofline — we’ve been hand-waving “the ridge point”; the roofline model makes it rigorous. Given a kernel’s arithmetic intensity, predict whether it’s memory- or compute-bound before you write it, and read the roofline chart Nsight generates.
05-optimizing-gemm (upcoming) — apply coalescing + SMEM tiling + bank-conflict-free layouts to take a matmul from a few percent of peak to near-peak. The concrete synthesis of §5, §6.
06-memory-wall-and-flashattention (upcoming) — the flagship IO-aware kernel: why standard attention is HBM-bound and how tiling + online softmax fixes it.

And back to the foundation: 01-gpu-architecture-and-simt — occupancy and latency hiding, which this lesson’s register/bandwidth story depends on.

Reference topic: gpu-systems-for-llms | Filed: 2026-09-02

References

NVIDIA, NVIDIA H100 Tensor Core GPU Datasheet, 2023 — H100 SXM5: ~989–990 TFLOP/s BF16/FP16 dense tensor-core throughput, 80 GB HBM3 at 3.35 TB/s, NVLink 4 at 900 GB/s. https://resources.nvidia.com/en-us-gpu-resources/h100-datasheet-24306 ↩ ↩² ↩³ ↩⁴ ↩⁵ ↩⁶
NVIDIA, NVIDIA H100 Tensor Core GPU Architecture whitepaper (GH100/Hopper), 2022 — 50 MB L2 cache on H100 SXM5, per-SM memory configuration, and full memory-subsystem description. https://resources.nvidia.com/en-us-tensor-core ↩ ↩²
NVIDIA, CUDA C++ Programming Guide — memory model, __shared__, 32 shared-memory banks (4 B wide) and the bank-conflict rule, coalescing / 32-byte sectors, cudaFuncAttributePreferredSharedMemoryCarveout, “local memory” residing in device DRAM, plus the Compute Capability 9.0 appendix (65,536 registers/SM, 255 registers/thread, 2048 resident threads/SM). https://docs.nvidia.com/cuda/cuda-c-programming-guide/ ↩ ↩² ↩³ ↩⁴ ↩⁵ ↩⁶ ↩⁷ ↩⁸ ↩⁹
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 register/SMEM/L2/global memory latencies and bandwidths on H100/H800. https://arxiv.org/abs/2402.13499 ↩ ↩² ↩³ ↩⁴ ↩⁵
M. Andersch et al., “NVIDIA Hopper Architecture In-Depth,” NVIDIA Technical Blog, 2022 — 256 KB combined L1/SMEM per SM with up to 228 KB configurable as shared memory. https://developer.nvidia.com/blog/nvidia-hopper-architecture-in-depth/ ↩ ↩²
NVIDIA, CUDA C++ Best Practices Guide — “Memory Optimizations”: global-memory coalescing, shared-memory bank conflicts, and L2/occupancy trade-offs. https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/ ↩ ↩² ↩³
NVIDIA, NVIDIA Tesla P100 whitepaper (GP100/Pascal), 2016 — ~21.2 TFLOP/s FP16, 720 GB/s HBM2. https://images.nvidia.com/content/pdf/tesla/whitepaper/pascal-architecture-whitepaper.pdf ↩
NVIDIA, NVIDIA Tesla V100 GPU Architecture whitepaper (GV100/Volta), 2017 — 125 TFLOP/s tensor-core FP16, 900 GB/s HBM2. https://images.nvidia.com/content/volta-architecture/pdf/volta-architecture-whitepaper.pdf ↩
NVIDIA, NVIDIA A100 Tensor Core GPU Architecture whitepaper (GA100/Ampere), 2020 — 312 TFLOP/s BF16/FP16 dense (624 with sparsity); A100 80 GB SXM at ~2.0 TB/s HBM2e. https://images.nvidia.com/aem-dam/en-zz/Solutions/data-center/nvidia-ampere-architecture-whitepaper.pdf ↩
NVIDIA, NVIDIA Blackwell Architecture Technical Brief, 2024 — B200: ~2.25 PFLOP/s FP16/BF16 dense (4.5 with 2:4 sparsity), ~8 TB/s HBM3e. https://resources.nvidia.com/en-us-blackwell-architecture ↩
M. Horowitz, “Computing’s Energy Problem (and what we can do about it),” IEEE International Solid-State Circuits Conference (ISSCC), 2014 — canonical per-operation energy figures showing DRAM access costs orders of magnitude more than an arithmetic op. ↩
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 — tiling + online softmax reduces attention HBM traffic from $O (N^{2})$ to $O (N^{2} d / M)$ . https://arxiv.org/abs/2205.14135 ↩
W. Hwu, D. Kirk, I. El Hajj, Programming Massively Parallel Processors: A Hands-on Approach, 4th ed., Morgan Kaufmann, 2022 — Ch. 5 (memory architecture and data locality / tiling) and Ch. 6 (performance considerations, coalescing). ↩

Gunjan Dhanuka — Learning Notes

Explorer

Learn: The GPU Memory Hierarchy

Learn: The GPU Memory Hierarchy

1. Learning map

2. Why this matters

3. The hierarchy, top to bottom

4. Registers and register pressure

5. Shared memory: the programmer-managed scratchpad

Banks and bank conflicts

6. Global memory and coalescing

7. The memory wall

Data movement dominates energy, too

8. GPU/LLM connection

9. Practice problems

10. Practice & resources

Hands-on exercises

Real resources

11. What’s next

References

Graph View

Table of Contents

Backlinks

Gunjan Dhanuka — Learning Notes

Explorer

Learn: The GPU Memory Hierarchy

Learn: The GPU Memory Hierarchy

1. Learning map

2. Why this matters

3. The hierarchy, top to bottom

4. Registers and register pressure

5. Shared memory: the programmer-managed scratchpad

Banks and bank conflicts

6. Global memory and coalescing

7. The memory wall

Data movement dominates energy, too

8. GPU/LLM connection

9. Practice problems

10. Practice & resources

Hands-on exercises

Real resources

11. What’s next

References

Footnotes

Graph View

Table of Contents

Backlinks