The Memory Hierarchy

A kernel is a plan for moving data. The arithmetic is almost incidental — an H100 can issue 989 TFLOP/s of BF16 matmul but only read 3.35 TB/s from memory, so for every byte you fetch you had better find 295 floating-point operations to do with it. Everything in this part is about where a byte can live, what it costs to get it, and the two ways you can accidentally throw away 8× or 32× of your bandwidth without any warning from the compiler.

Five tiers, from the fastest thing on the chip to the slowest. The numbers are for one H100 SXM5.

* L2 bandwidth is an estimate, not an NVIDIA specification: Luo et al., Table V measure 3,942.4 B/clock aggregate on H800; applying an H100 1.98 GHz clock gives about 7.8 TB/s. HBM bandwidth is the published NVIDIA H100 figure.

Threads do not read bytes from DRAM. The memory system reads sectors, and a sector is 32 bytes, naturally aligned. Four sectors make a 128-byte cache line. The Programming Guide states the rule directly:

So the question for every load instruction is: how many distinct 32-byte sectors do the 32 lanes of this warp touch? The best case is four — 32 lanes × 4 bytes = 128 contiguous bytes. The worst case is thirty-two, and then you are using one eighth of your bandwidth to move data you will throw away.

Three cases are worth internalising. Write the byte address of lane l as base + 4·l and count the sectors it crosses:

• Stride 1, aligned — lane l reads byte 4·l, so the warp spans bytes 0–127: sectors 0, 1, 2, 3. Four sectors, 100% efficiency. This is what y[i] = a*x[i] + y[i] gives you when i = blockIdx.x*blockDim.x + threadIdx.x.
• Stride 1, misaligned by 4 bytes — lane l reads byte 4·l + 4; the span 4–131 now touches sectors 0 through 4. Five sectors, 80% efficiency. Costs you one extra sector per warp; often partially recovered from L2 because the neighbouring warp already pulled that line.
• Stride 8 floats — lane l reads byte 32·l, and since one sector is exactly 32 bytes every lane opens a fresh one: 32 sectors, 12.5% efficiency. Any larger stride behaves the same. This is the column-major access pattern, and it is the single most common reason a "simple" kernel runs 8× slower than it should.

2.2.1 Why the block shape decides this

Recall from Part 1 that warps are cut along the linear thread index, so threadIdx.x varies fastest within a warp. That makes the following two kernels — identical in every other respect — differ by 8×:

// dim3 block(32, 8);  A is row-major, N columns.

// COALESCED: within a warp, tid.y is fixed and tid.x runs 0..31,
// so the 32 lanes read 32 consecutive floats = 4 sectors.
int row = blockIdx.y * 8 + threadIdx.y;
int colM = blockIdx.x * 32 + threadIdx.x;
float v = A[row * N + colM];

// UNCOALESCED: swap the roles. Now the 32 lanes of a warp read
// addresses N*4 bytes apart -> 32 separate sectors.
float w = A[colM * N + row];

# Triton makes the same distinction, just with block pointers.
offs_m = pid_m * BLOCK_M + tl.arange(0, BLOCK_M)
offs_n = pid_n * BLOCK_N + tl.arange(0, BLOCK_N)

# Coalesced: the LAST axis of the offset tensor is contiguous in memory.
a = tl.load(A + offs_m[:, None] * stride_am + offs_n[None, :] * stride_an)

# If stride_an != 1 you have a strided load. Triton will still
# emit it, and it will still cost you sectors.

A lane can load 1, 2, 4, 8 or 16 bytes in a single instruction, provided the address is naturally aligned to that size. Using the wide forms does two things at once: it cuts the instruction count by 4×, and — more importantly — it quadruples the bytes each thread has in flight.

That second effect is Little's Law again. From Part 1: under the stated latency estimate, an H100 needs ~915 kB (893 KiB) of memory traffic in flight to saturate HBM, which is 3.4 bytes per thread at full occupancy. A B200 needs more than 15. A kernel where each thread has one 4-byte load outstanding is structurally incapable of saturating a B200's memory system, no matter how good its access pattern is.

// 1 instruction, 16 bytes per lane, 128 bytes per QUARTER-warp.
// Requires x to be 16-byte aligned and n divisible by 4.
const float4* x4 = reinterpret_cast<const float4*>(x);
float4 v = x4[i];                 // compiles to LDG.E.128
v.x *= a; v.y *= a; v.z *= a; v.w *= a;

// cudaMalloc always returns >=256-byte-aligned pointers, so the base
// is fine; it is your OFFSETS that break alignment.

One warp issuing LDG.E.128 requests 512 bytes — sixteen sectors, four cache lines — in one instruction. Same efficiency and payload as four separate warp-wide 4-byte coalesced load instructions, one quarter of the instructions, four times the memory-level parallelism per thread. This is why many high-performance kernels use float4 or int4 as their unit of global traffic when alignment and layout permit it.

2.4.1 L1 and shared memory are the same silicon

L1 and shared memory have been one physical array since Volta, split at kernel launch between hardware-managed cache and programmer-managed scratchpad. The array has grown every generation — 96 KB on Volta, 192 KB on A100, 256 KB per SM on Hopper and Blackwell (128 KB on consumer parts) — and so has the scratchpad's maximum share: 164 KB of shared memory per SM on A100, 228 KB on H100 and B200. Selecting the 228 KB carveout leaves 28 KB of L1; because CUDA reserves 1 KB per block, one block can request at most 227 KB of it. That is a real trade, and for a well-tiled kernel it is usually the right one — you are replacing a cache that guesses with a scratchpad that knows.

// Static __shared__ arrays are capped at 48 KB. Above that: dynamic + opt-in.
cudaFuncSetAttribute(kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, 227*1024);

// The carveout is a HINT; percentages round UP to a supported capacity.
cudaFuncSetAttribute(kernel, cudaFuncAttributePreferredSharedMemoryCarveout, 100);

2.4.2 L2 is big, shared, and not uniform

40 MB on an A100, 50 MB on an H100, ~126 MB on a B200 — large enough that for many LLM decode steps the entire activation working set is L2-resident, and only the weights come from HBM. But L2 is physically split into two partitions, each closer to half the GPCs. Crossing that boundary costs measurably more latency and bandwidth, and on a B200 the split is also a die boundary: measured 21 TB/s within a partition against 16.8 TB/s across it.

NVIDIA publishes the 126 MB GB200 L2 capacity in its Blackwell Tuning Guide. The 21/16.8 TB/s partition-bandwidth figures are separate original Chips and Cheese measurements, not NVIDIA specifications.

This is why tile scheduling order matters. A GEMM that walks output tiles in row-major order streams a fresh row of B through L2 for every tile; one that walks them in small 2D groups reuses what is already there. Part 5 measures the difference.

You also get an explicit lever: a persistence window that marks a range of global memory as preferentially retained in L2.

// Set aside part of L2 for persisting accesses (max 75% of L2 on A100/H100).
cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, size);

cudaStreamAttrValue attr;
attr.accessPolicyWindow.base_ptr  = kv_cache;      // hot, reused every step
attr.accessPolicyWindow.num_bytes = bytes;
attr.accessPolicyWindow.hitRatio  = 0.6f;         // <1.0 avoids LRU thrash
attr.accessPolicyWindow.hitProp   = cudaAccessPropertyPersisting;
attr.accessPolicyWindow.missProp  = cudaAccessPropertyStreaming;
cudaStreamSetAttribute(stream, cudaStreamAttributeAccessPolicyWindow, &attr);

At the instruction level, PTX exposes cache operators for the same idea — ld.global.cg to skip L1 entirely, .cs ("streaming", evict-first) for data you will touch once, .L1::evict_last for data you want kept. In CUDA C++ these surface as __ldcg, __ldcs, __stwt, and cuda::annotated_ptr.

Shared memory is a scratchpad: no tags, no eviction, no misses. It is organised as 32 banks, each 4 bytes wide, each serving one 32-bit word per clock. Successive 32-bit words live in successive banks, so:

A warp-wide access is conflict-free if the 32 lanes hit 32 different banks — or the same word, which is broadcast for free. Anything else serialises:

So an N-way conflict costs exactly N×. The pathological case is 32-way, and you hit it with the most natural code imaginable.

2.5.1 Why the column access is 32-way

With __shared__ float tile[32][32], element tile[r][c] lives at word index 32r + c, so its bank is (32r + c) mod 32 = c. The row index has no effect on the bank. Now run the two natural warp accesses lane by lane, say for row/column 7:

A warp reading a column asks one bank for 32 different words, and gets them one per cycle.

2.5.2 Padding: the cheap fix that breaks later

Declare tile[32][33] instead. Now tile[r][c] is at word 33r + c, bank (33r + c) mod 32 = (r + c) mod 32. Reading column c gives banks c, c+1, c+2, … — all 32 distinct. Conflict gone, for the price of 3.1% of your shared memory.

This one-float pad is right for the hand-written FP32 access above. It does not compose with the standard tensor-core layouts: a 33-float row is 132 bytes, so every row after the first loses the 16-byte alignment expected by ldmatrix, wgmma descriptors and TMA layouts. Padding is not categorically forbidden if the padded stride preserves every required alignment, but the naive +1 fix does not. Which brings us to the usual alternative.

A swizzle permutes addresses within a power-of-two block instead of expanding the block. The trick is to XOR the 16-byte-chunk index with the row index: for any fixed row, c ⊕ r is a permutation of that row's chunks — apply it twice and you are back where you started — so every element still has exactly one home and the array stays dense. And because the permutation moves whole 16-byte chunks and never touches the bits inside a chunk, 16-byte alignment survives.

CUTLASS expresses the whole family with three integers:

M is how many low bits to leave alone, B is how many bits participate, S is the shift. Every swizzle mode that TMA and wgmma accept uses M = 4 (the low 16 bytes are never permuted) and S = 3 (the pattern repeats every 8 rows):

The 128 B mode is the one you want by default: its atom is 8 rows × 128 bytes, which matches both a full cache line on the global side and a full conflict-free warp access on the shared side. In C it is one line:

The Swizzle<B, M, S> parameterisation is CUTLASS’s (cute::Swizzle); the same three non-trivial patterns are the CU_TENSOR_MAP_SWIZZLE_32B/64B/128B values accepted by cuTensorMapEncodeTiled.

The figure below is the important one in this part. It shows a 128-byte-swizzled 8×8 grid of 16-byte chunks, and what happens when a tensor-core load asks for the same logical chunk from all 8 rows — which is exactly what ldmatrix.m8n8 does.

One constraint to know before you meet it: with 128 B swizzle the innermost TMA box dimension must be ≤ 128 bytes, which for BF16 means at most 64 elements. A BLOCK_K = 128 BF16 K-major tile therefore needs a split or reoriented box, a different contiguous dimension, or no swizzle. A smaller non-zero swizzle does not help: the Driver API limits tighten to 64 and 32 bytes for the 64 B and 32 B modes. Violating this constraint is a common cause of CUDA_ERROR_INVALID_VALUE out of cuTensorMapEncodeTiled.

Historically, staging a tile from global to shared memory meant a round trip through registers: LDG into a register, STS out to shared. That burns registers, burns issue slots, and forces the thread to wait.

Ampere added cp.async, which copies global → shared without touching the register file, asynchronously, with group-based completion tracking. Hopper added TMA, which goes further: a single thread describes a multidimensional tile copy, and dedicated hardware generates all the addresses, handles the out-of-bounds edges, applies the swizzle, and signals a barrier when the bytes have landed.

// Each thread copies 16 B global -> shared, bypassing registers.
unsigned smem_addr = __cvta_generic_to_shared(&smem[stage][tid*4]);
asm volatile("cp.async.cg.shared.global [%0], [%1], 16;\n"
             :: "r"(smem_addr), "l"(&gmem[idx]));

asm volatile("cp.async.commit_group;\n");       // close this stage
asm volatile("cp.async.wait_group 2;\n");       // keep 2 stages in flight
__syncthreads();

// One thread issues the whole tile copy; hardware does the rest.
if (threadIdx.x == 0) {
    mbarrier_init(&bar[stage], 1);
    mbarrier_expect_tx(&bar[stage], TILE_BYTES);      // byte-counting barrier
    cp_async_bulk_tensor_2d(smem[stage], tensor_map, k_off, m_off, &bar[stage]);
}
// Everyone waits on the phase bit; no per-thread address math at all.
mbarrier_wait_parity(&bar[stage], phase);

Both exist to enable the same structure: a multi-stage software pipeline, where the copy for tile i+2 is in flight while the math for tile i is running. This is the single most important structural idea in modern GEMM and attention kernels, and it is why shared memory is allocated as a circular buffer of 3–6 stages rather than one tile.

Hopper added a level between block and grid. A thread block cluster is up to 8 blocks (16 non-portably) guaranteed to be co-scheduled on one GPC, and — the actual point — every block in a cluster can load, store and do atomics directly on the shared memory of its peers, over a dedicated SM-to-SM network. This is distributed shared memory (DSMEM).

Measured on an H800: SM-to-SM latency 180 cycles, a 32% reduction versus going through L2, and 3.27 TB/s of throughput at cluster size 2. A size-8 cluster on an H100 gives a kernel 1.8 MB of single-hop, software-managed memory.

Source: Luo et al., Benchmarking and Dissecting the NVIDIA Hopper GPU Architecture, §III-D3; measured on an H800 PCIe.

The practical payoff for GEMM is TMA multicast: one block issues a TMA load and the hardware delivers the same tile into the shared memory of every block in the cluster. Two SMs cooperating on one B-tile load halves the L2 traffic for that operand — worth 704 → 734 TFLOP/s in one benchmark in this H100 GEMM worklog. Treat that delta as configuration-specific, not a universal multicast gain.

__global__ void __cluster_dims__(2, 1, 1) kernel(
    const float* input, float* output) {
    namespace cg = cooperative_groups;
    cg::cluster_group cluster = cg::this_cluster();

    extern __shared__ float smem[];
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    smem[threadIdx.x] = input[i];

    cluster.sync();                     // every peer exists and its SMEM is initialized
    float* peer = cluster.map_shared_rank(smem, cluster.block_rank() ^ 1);
    float v = peer[threadIdx.x];        // ~180 cycles, no L2 round trip
    output[i] = v;
    cluster.sync();                     // no block exits while a peer still reads its SMEM
}

Everything above compresses into six questions to ask of any kernel:

1. Does a warp's global load span 128 contiguous bytes? If not, count the sectors. Fix the block shape or transpose the data before you fix anything else.
2. Are you loading 16 bytes per thread per instruction? If you are loading 4, you have a quarter of the memory-level parallelism you could have, and on Blackwell that is disqualifying.
3. Does a warp's shared-memory access hit 32 distinct banks? If it reads a column of a power-of-two-strided array, it does not.
4. If tensor cores are involved, does the shared layout satisfy their alignment? The naive one-element padding fix breaks the standard ldmatrix, wgmma and TMA layouts; a compatible swizzle usually avoids both the conflict and the extra stride.
5. Is there anything in flight while you compute? A single-stage loop pays full memory latency on every iteration.
6. Are you re-reading from HBM what L2 or SMEM already has? That is a scheduling-order question, not a bandwidth question — Part 5.

Next: how to tell, before writing a line of code, whether a kernel can be fast — and which of the two ceilings it will hit.