Cache-Aware Tiling Strategies
Cache-aware tiling decomposes massive multi-dimensional matrices into smaller geometric tiles that fit perfectly within L1/Shared Memory.
Source: mortalapps.com- Cache-aware tiling decomposes massive multi-dimensional matrices into smaller geometric tiles that fit perfectly within L1/Shared Memory.
- The core purpose is to maximize memory reuse, fetching data from slow HBM once and executing
math operations on it before eviction. - The primary optimization idea is balancing the dimensions of the tile (
) against register file limits and shared memory capacity. - The most important engineering insight is that the ratio of Compute (FLOPs) to Memory Access (Bytes) scales with tile size, making larger tiles definitively better—until they spill out of SRAM.
Why This Matters
Standard matrix multiplication (
) without tiling performs operations using memory accesses, making it fundamentally and unavoidably memory-bound. Cache-aware tiling reduces the memory access complexity to 
, dramatically elevating the operation's Arithmetic Intensity. Revolutionary algorithms like FlashAttention are entirely predicated on exactly this concept.
Core Intuition
Imagine painting a massive billboard. If you dip your brush in a bucket on the ground, climb the ladder, paint a stroke, and climb down, you waste all your time climbing. Tiling is carrying a small palette up the ladder. You fill the palette (the tile), climb up, and paint hundreds of strokes (the math) until the palette is dry. The size of the palette is strictly limited by how much you can physically hold (Shared Memory/Registers).
Technical Deep Dive
A dense GEMM kernel iteratively loads a tile of matrix 
(shape 
) and a tile of matrix 
(shape 
) into Shared Memory. The physical capacity of the hardware dictates the maximum limits: Hopper H100 provides 228 KB of shared memory per SM (max 227 KB usable per thread block). Blackwell B200 (GB100) retains 228 KB per SM, matching H100. Tile sizes are governed by the same 228 KB physical limit. To compute the output tile 
(
), each thread claims a sub-tile, moving values from Shared Memory into its Registers for calculation.
Key Takeaways
Internals / Execution Flow
The execution flow manages the movement of tiles down the hierarchy. First, Producer warps load and tiles from Global Memory into the 228 KB Shared Memory capacity utilizing TMA or cp.async. Next, Consumer warps load sub-tiles of and into the 255 maximum allowed registers per thread via ldmatrix. Tensor Cores perform multiply-accumulate operations (mma/wgmma) on the register data. The kernel slides the 
-dimension window, accumulating the partial sums in the registers. Finally, the fully accumulated tile is written from registers back to global memory.
Performance Implications
Larger tiles yield higher arithmetic intensity. An 
tile requires loading 
elements to perform 
operations. The ratio is 
. Thus, doubling the tile size doubles the compute-to-memory ratio, insulating the execution from the 4.8 TB/s HBM bandwidth limit. However, if the tile size exceeds the 228 KB shared memory limit or the 255 register limit, the compiler introduces register spilling, completely destroying performance by introducing global memory traffic.
Real-World Usage
In FlashAttention, the massive Attention matrix 
is intentionally never written to global memory. The algorithm tiles 
, , and 
such that the Softmax denominator can be incrementally updated on-chip. The block sizes (e.g., 
) are mathematically derived to fit exactly within the local SRAM bounds of the target GPU architecture, ensuring the footprint remains entirely in fast memory.
Interview Guidance
Candidates will frequently be asked: "Derive the optimal block size for FlashAttention on an A100 vs an H100." The answer requires knowing the exact SM shared memory sizes (164 KB A100 vs 228 KB H100) and correctly calculating the byte footprint of the FP16/BF16 
tiles based on those bounds.