CUDA Memory Hierarchy & Access Patterns

Coalesced global memory access, tiling with shared memory, bank conflicts, warp shuffle intrinsics, and asynchronous memory prefetching (TMA). The techniques behind FlashAttention and high-performance GEMM.

Overview

Why memory access decides kernel speed

Memory access patterns are the single most impactful factor in GPU kernel performance. A kernel with poor HBM access patterns can reach only 10% of hardware peak bandwidth, while a well-optimized one reaches 90%+.

Coalescing — Global memory — Wasted cache-line transactions
Shared memory tiling — Global to shared — Repeated HBM reads of reused data
Padding — Shared memory — Bank conflicts
Warp shuffles — Registers — Barrier-heavy shared-memory reductions
cp.async and TMA — Global to shared — Loads that stall compute

Global memory coalescing

When a warp of 32 threads issues a load, the memory controller tries to serve the 32 requests with as few 128-byte cache-line transactions as possible.

Coalesced: all 32 threads read consecutive 4-byte floats from a 128-byte aligned address, so one transaction serves the warp.
The rule: thread i accesses base + i*stride with stride = 1 (unit stride).

> [!WARNING] Uncoalesced access is the single most common cause of GPU kernel underperformance. A 32-way stride access wastes 97% of every cache line fetched and can drop effective HBM bandwidth from 3.35 TB/s to ~100 GB/s.

Shared memory and tiling

Shared memory (SMEM) is a software-managed on-chip cache: the programmer explicitly loads data into it, computes on it and writes results back. Its ~10-30 cycle latency against ~500 cycles for HBM means data reused across a thread block should be loaded once into SMEM and read many times.

Naive kernel — Each element of A and B once per output element, O(N^3) — 0.25 FLOPs/byte
Tiled kernel — Threads cooperatively load tiles of A and B into SMEM — tile_size/2 FLOPs/byte
16x16 tile — Same as tiled — 8 FLOPs/byte, enough to approach the roofline

Bank conflicts

Shared memory has 32 banks, matching the warp size, and each bank serves one 4-byte request per cycle. Two threads in a warp hitting the same bank at different addresses conflict; the same address is a broadcast and does not.

32-element array, stride 32 — All threads on bank 0 — 32-way bank conflict
Array padded by one element, stride 33 — Spread across banks — Conflict-free

Warp shuffle intrinsics

__shfl_sync and its variants (__shfl_up_sync, __shfl_down_sync, __shfl_xor_sync) let threads in a warp exchange registers directly, the lowest-latency communication within a warp.

A warp-level sum with __shfl_down_sync takes ~log2(32) = 5 steps, far faster than a shared-memory reduction that needs barriers. Softmax, layer norm and attention kernels use these heavily for per-row reductions over 128-head-dim vectors.

Asynchronous prefetching and TMA

cp.async — Ampere — DMA-style global-to-shared loads that proceed while CUDA cores compute
Tensor Memory Accelerator (TMA) — Hopper — 2D and 3D strided copies with hardware-managed addressing, removing per-thread addressing overhead

Together they enable the async pipeline pattern of high-performance GEMM.

Memory access in transformers

Autoregressive decoding is memory-bandwidth-bound. A batch-size-1 forward pass reads all model weights (e.g., 140 GB for LLaMA-3 70B) for only O(seq_len) compute per weight.

Maximize batch size — Raises arithmetic intensity
Quantize weights — Fewer bytes loaded per step
Fuse operations — Fewer passes over memory

> [!NOTE] FlashAttention's key insight: attention does O(seq_len^2) work but needs only O(seq_len) memory if tiled correctly, because it never materializes the full NxN attention matrix in HBM.

Interview questions

Walk through a tiled matrix multiplication kernel -- how does it improve performance over the naive version?
Explain FlashAttention and why it dramatically reduces memory usage compared to standard attention.
How do you identify and fix a shared memory bank conflict?