Memory Hierarchy and Coalesced Global Access

Core Concept

Registers, shared memory/L1, L2, and DRAM have different capacities and latencies. Global-memory efficiency depends on whether a warp's addresses can be served by a small number of memory transactions.

Why It Affects Performance

Thirty-two lanes may request only 128 useful bytes yet trigger far more traffic because of stride, alignment, or layout, wasting bandwidth and increasing long-scoreboard stalls.

Example or Pseudocode

float v = x[base + threadIdx.x];          // adjacent lanes, coalesced
float w = x[(base + threadIdx.x)*stride]; // scattered when stride is large

For row-major matrices, mapping threadIdx.x to adjacent columns is usually natural.

Profiling and Performance Judgment

Use global load/store efficiency, sectors per request, and L1/L2/DRAM byte counters. Compare the algorithmic minimum traffic with measured traffic.

Five-Minute Exercise

A warp loads FP32 values from x[lane*32]. Why can 128 useful bytes require many transactions?