GPU Memory Hierarchy (Registers, Shared, Global)

Memory resolution flows sequentially from the instruction issue. If an operand is located in a register, the access latency is approximately 1 cycle. If an operand requires shared memory, access typically takes 20-30 cycles, assuming a perfectly bank-conflict-free access pattern. If an address points to global memory, the hardware routes the request to the L1 cache. A miss in L1 routes to the L2 cache, which spans all SMs. A miss in L2 prompts a global memory read. This final step stalls the warp until the memory controller retrieves the data, unless the scheduler finds other independent instructions or active warps to execute in the interim.

Registers represent the hardest physical ceiling on concurrent execution. If a kernel requires 255 registers per thread, maximum occupancy drops drastically because the 64K register file per SM is exhausted quickly. For shared memory operations, local access is highly efficient at roughly 29 cycles. However, remote DSM access within a cluster of size 2 incurs a higher latency of roughly 181 cycles, which is still significantly faster than traversing to HBM but requires careful pipeline hiding.

FlashAttention relies fundamentally on exact mastery of this hierarchy. By aggressively tiling the Attention matrix (), the algorithm ensures that sequence blocks fit entirely within the fast SRAM (shared memory). This allows the algorithm to complete the reduction without writing intermediate self-attention values back to Global Memory, achieving massive speedups over standard implementations.

To optimize the hierarchy, engineers explicitly tune __launch_bounds__ in CUDA to force the compiler to use fewer registers, sacrificing per-thread state for higher overall SM occupancy. Additionally, utilizing Hopper's Thread Block Clusters allows kernels to leverage DSM for datasets slightly larger than a single SM's shared memory, preventing L2 thrashing. Finally, packing smaller data types (FP8/FP4) using vector registers (e.g., int4, short2) effectively doubles register utilization efficiency without requiring hardware changes.

The ptxas compiler provides explicit output detailing hierarchy usage at compile time, outputting critical logs such as ptxas info : Used 255 registers, 512 bytes smem. Profilers like Nsight Compute highlight L1/L2 hit rates, identify occupancy limiters, and flag excessive local memory usage.

Candidates must be able to calculate theoretical SM occupancy based on register counts and shared memory configurations. Weak engineers commonly confuse "local memory" with "shared memory"; strong engineers explicitly know that local memory is thread-private global memory used for stack frames and register spilling, and architect their kernels to avoid it entirely.