Tiling Strategies — infra.bea.ai

Tiling is the single most important optimization pattern in GPU programming. Every fast kernel — GEMM, FlashAttention, reduction, convolution — uses tiling to move data from slow HBM into fast shared memory and maximize data reuse. This chapter covers the principles, patterns, and tradeoffs.

1. The Core Idea

Without tiling: every operation reads from HBM (~600 cycle latency). With tiling: load a chunk into shared memory (~30 cycle latency), reuse it many times, then load the next chunk.

Without tiling:                    With tiling:
HBM → Compute → HBM               HBM → Shared Mem → Compute (reuse N times)
HBM → Compute → HBM                                → Compute (reuse)
HBM → Compute → HBM                                → Compute (reuse)
...                                HBM → Shared Mem → Compute (reuse N times)
                                   ...

The math: If each data element is used $R$ times, tiling reduces HBM reads by $R\times$ . For GEMM with tile size $T$ : $R = T$ , so a 32×32 tile gives 32× fewer HBM reads.

2. Tiling for Matrix Multiplication

The Problem

Naive GEMM: each output element $C_{ij}$ reads one row of $A$ and one column of $B$ from HBM:

// Naive: O(M*N*K) HBM reads
C[i][j] = 0;
for (int k = 0; k < K; k++)
    C[i][j] += A[i][k] * B[k][j];  // each A[i][k] and B[k][j] from HBM

For $M=N=K=4096$ : each element reads 4096 floats from A and 4096 from B. Total: $4096^2 \times 2 \times 4096 = 137$ billion bytes from HBM. At 3.35 TB/s, that’s 41 ms — far slower than the 3.4 ms achievable with tiling.

The Tiled Solution

Tiled GEMM

Divide output matrix $C$ into tiles of size $T_M \times T_N$
Each thread block computes one output tile
Loop over K in steps of $T_K$ : a. Cooperatively load a $T_M \times T_K$ tile of $A$ into shared memory b. Cooperatively load a $T_K \times T_N$ tile of $B$ into shared memory c. __syncthreads() — wait for all loads d. Compute partial dot products using shared memory data e. __syncthreads() — wait before loading next tile

__shared__ float As[TILE_M][TILE_K];
__shared__ float Bs[TILE_K][TILE_N];

float acc = 0.0f;
for (int t = 0; t < K; t += TILE_K) {
    // Cooperative load: each thread loads one or more elements
    As[threadIdx.y][threadIdx.x] = A[row][t + threadIdx.x];
    Bs[threadIdx.y][threadIdx.x] = B[t + threadIdx.y][col];
    __syncthreads();

    // Compute: all data in fast shared memory
    for (int k = 0; k < TILE_K; k++) {
        acc += As[threadIdx.y][k] * Bs[k][threadIdx.x];
    }
    __syncthreads();
}
C[row][col] = acc;

Data Reuse Analysis

With tile size $T$ :

Each element of $A$ ‘s tile is used by $T_N$ threads (all threads in the same row of the output tile)
Each element of $B$ ‘s tile is used by $T_M$ threads
Reuse factor: $T$ reads from shared memory instead of HBM

For $T=32$ : HBM reads reduce by 32×. For $T=128$ : 128× reduction.

3. Double Buffering: Overlapping Load and Compute

Naive tiling has a stall: after loading a tile and computing, all threads wait for the next load. Double buffering overlaps loading the next tile with computing the current tile:

__shared__ float As[2][TILE_M][TILE_K];  // two buffers
__shared__ float Bs[2][TILE_K][TILE_N];
int curr = 0;

// Load first tile
load_tile(As[curr], Bs[curr], t=0);
__syncthreads();

for (int t = TILE_K; t < K; t += TILE_K) {
    int next = 1 - curr;

    // Start loading NEXT tile while computing CURRENT tile
    load_tile_async(As[next], Bs[next], t);  // non-blocking

    // Compute with current tile (in shared memory)
    compute(As[curr], Bs[curr], acc);

    __syncthreads();  // wait for both load and compute
    curr = next;
}
// Process last tile
compute(As[curr], Bs[curr], acc);

Benefit: Load latency (~600 cycles) is hidden behind compute. For compute-heavy tiles, this gives near-zero load overhead.

4. Tile Size Tradeoffs

Larger tiles	Smaller tiles
✅ More data reuse → fewer HBM reads	✅ Less shared memory → higher occupancy
✅ Better arithmetic intensity	✅ Less register pressure
❌ More shared memory → fewer blocks per SM	❌ More HBM reads
❌ More registers for accumulators	❌ Lower arithmetic intensity
❌ Tail effects if matrix not divisible	✅ Better for small/irregular matrices

Typical Tile Sizes in Practice

Kernel	Tile size	Why
Simple GEMM	32×32	Fits in shared memory, good starting point
Production GEMM (cuBLAS)	128×128 or 256×128	Maximizes reuse, uses register tiling
FlashAttention	64-128 per block	Balances reuse with SRAM capacity
Reduction	Block size (256-1024)	One tile per block

Register Tiling

Production kernels like cuBLAS use a two-level tiling strategy:

Thread block tile: loaded into shared memory (128×128)
Per-thread tile: each thread accumulates a small sub-tile in registers (8×8)

This maximizes reuse at both levels — shared memory reuse across threads, register reuse within a thread.

5. Tiling Beyond GEMM

The tiling pattern applies everywhere:

Kernel	What’s tiled	What’s reused
GEMM	A and B matrices along K dimension	Each tile element used by many output elements
FlashAttention	Q, K, V blocks along sequence	K/V blocks reused across Q rows
Convolution	Input patches and filter weights	Filter weights reused across spatial positions
Reduction	Input array chunks	Partial sums accumulated in shared memory
Softmax	Input chunks for max/sum	Two-pass (max, then normalize) on same data

The universal principle: If data is read from HBM more than once, it should be tiled into shared memory.

Self-Check

Explain why tiling reduces HBM reads and by how much (reuse factor)
Implement tiled matrix multiplication with shared memory
Explain double buffering and how it hides load latency
Analyze the tile size tradeoff (reuse vs occupancy)
Explain two-level tiling (shared memory + register)
Identify tiling opportunities in a given kernel