9. The memory hierarchy and copy traffic

“Memory-bound” can describe bandwidth saturation, latency, poor transactions, cache thrashing, spills, or too little parallelism. These require different fixes.

Figure 9.1 — Capacity and scope grow downward; locality and latency improve upward. Always associate a byte count with one level.

Bandwidth and arithmetic intensity

Effective bandwidth is

$$BW_{effective}=\frac{B_{useful,read}+B_{useful,write}}{t}.$$

Arithmetic intensity at a chosen hierarchy level is

$$I=\frac{\text{operations}}{\text{bytes transferred at that level}}.$$

The roofline bound is

$$P \le \min(P_{peak}, I\cdot BW_{peak}).$$

Always name the byte level. DRAM arithmetic intensity and L1 arithmetic intensity are not interchangeable.

Explicit copies versus streaming kernels

“Memory-bound” also conflates explicit copies with traffic that kernel instructions generate on their own; separate the two before deciding what to optimize.

In a representative graph-replayed training trace, one input D2D copy per microstep was negligible. The apparent “copy storm” was mostly quantization, pointwise, and normalization kernels repeatedly streaming activation or weight buffers. Optimizing cudaMemcpyAsync would not address that traffic.

Coalescing and sector utilization

On modern CUDA GPUs, warp requests are serviced in aligned sectors. Sectors exist because the memory controller and the cache hierarchy above it move data in fixed-size chunks regardless of how many bytes a thread actually asked for; there is no hardware path that fetches just the four bytes a thread wants. When a warp's 32 threads request small, scattered addresses, each request can still trigger a full-sector fetch, so the warp pays for many complete sectors while using only a handful of bytes from each — that is the physical reason low sector efficiency wastes bandwidth. Compare requested useful bytes with transferred sectors. If each local or global load uses only 1–4 bytes of a 32-byte sector, the cache and memory system moves much more data than the program consumes.

For useful bytes $B_u$ and transferred bytes $B_t$,

$$\eta_{sector}=\frac{B_u}{B_t}.$$

Low efficiency suggests changing lane-to-data mapping, vectorization, structure-of-arrays layout, or tail handling. High efficiency with saturated DRAM suggests reducing total bytes or fusing adjacent passes.

Vectorization deserves a concrete description rather than just a name in that list. Vectorized memory access means loading or storing through a wider type — float4, int4, or an equivalent 8- or 16-byte-wide access — so each thread issues one instruction for one full, naturally aligned chunk instead of two, four, or more instructions for scattered smaller pieces. This helps in two distinct ways that are worth separating: it cuts the instruction count feeding the LSU pipe, the same issue-pressure argument ch. 8 makes for memory-bound stalls, and, when the underlying data layout is already contiguous per thread, it directly raises $\eta_{sector}$ by construction, because a full-width aligned request is exactly what a sector fetch is built to serve. Vectorization does not fix a fundamentally scattered or misaligned access pattern — it makes an already-contiguous pattern cheaper to issue, so confirm the layout supports it before reaching for wider loads as the fix for a low sector-efficiency number.

Shared-memory bank conflicts

Sector efficiency governs global and local traffic; shared memory has an analogous, on-chip bottleneck. Shared memory is physically built from multiple independent banks precisely so that up to 32 simultaneous accesses — one per lane, one per bank — can be serviced together in a single cycle, rather than one address at a time. A bank conflict is what happens when two or more lanes in the same request address the same bank: the bank can only honor one of them per cycle, so the hardware serializes the colliding accesses into extra wavefronts instead of servicing the whole warp at once. Nsight Compute reports excessive wavefronts/conflicts for shared loads and stores.

The classic transpose fix pads a square shared tile:

__shared__ float tile[TILE][TILE + 1];

The extra column changes the bank mapping of column accesses. Verify rather than assuming: modern instructions, element widths, and architectures alter the exact pattern.

Padding is simple but wastes the padding bytes on every row and can waste more once tile widths grow or vector-width accesses are involved. High-performance library kernels more often reach for shared-memory swizzling: a permuted address mapping — an XOR or similar bit-manipulation of the address, chosen so that colliding indices are pushed to distinct banks — that avoids the same conflicts without giving up any shared-memory capacity to padding. The tradeoff is that a correct swizzle depends on the exact access pattern, element width, and tile shape, and hand-deriving one is easy to get subtly wrong in a way that only shows up as a wrong answer or a conflict that reappears under a different tile size. Rather than deriving swizzle arithmetic from scratch, verify a candidate swizzle against a known-good library implementation such as CUTLASS's shared-memory layouts, or use the library's layout directly.

A memory decision table

Combine bandwidth, cache, and stall evidence into one table before picking a fix:

For a memory-bound activation kernel already sustaining roughly 89% of DRAM peak, micro-optimizing arithmetic has little ceiling. Fusion that avoids a read and write is the high-leverage experiment.

That "high DRAM throughput" row is also the signal that separates a kernel worth pipelining from one that is not. Once DRAM throughput confirms the kernel is actually bandwidth-bound rather than merely latency-stalled with bandwidth to spare, software pipelining — issuing the next tile's load before consuming the current tile's data, so load latency overlaps with compute on data already resident, as described in ch. 8's stall-to-fix section — is the technique that lets a kernel approach the DRAM ceiling instead of stalling between each tile's load and use. Pipelining does not reduce bytes moved, so it does not help the "reduce bytes, fuse, compress" row above; it helps a kernel spend the bytes it does move without adding idle cycles around each transfer.