Tiling: Trading On-Chip Storage for Reuse

Core Concept

Tiling stages reusable blocks of data in shared memory or registers so that one DRAM load serves multiple computations.

Why It Affects Performance

Repeated DRAM loads can make matrix and convolution kernels bandwidth-bound. Reuse raises arithmetic intensity, but oversized tiles increase registers, shared memory, synchronization, and tail waste.

Example or Pseudocode

for (int k0 = 0; k0 < K; k0 += BK) {
  load_A_tile_to_shared(); load_B_tile_to_shared();
  __syncthreads();
  for (int k = 0; k < BK; ++k) acc += As[row][k] * Bs[k][col];
  __syncthreads();
}

Profiling and Performance Judgment

Compare global bytes, L2 hit rate, shared transactions, registers/thread, occupancy, and movement on the Roofline plot.

Five-Minute Exercise

Why is a larger K tile not always faster?