Tiling groups neighboring cells so that their updated magnetic values can be reused for electric updates. FDTD CUDA's v4 kernel allocates two shared-memory tiles, each with halo space:

extern __shared__ FloatType tile[];
const int width = static_cast<int>(blockDim.x) + 1;
FloatType *tileHx = tile;
FloatType *tileHy = tile + width * (static_cast<int>(blockDim.y) + 1);
const int c =
    (static_cast<int>(threadIdx.y) + 1) * width + static_cast<int>(threadIdx.x) + 1;

The extra row and column hold magnetic values needed from the lower and left neighbors. Threads fill their tile entries, selected edge threads compute halo values, and all threads synchronize before reading neighbors.

Each block can then compute electric updates from its local shared tile instead of launching a separate whole-grid electric phase. Outputs go to separate field arrays, followed by buffer swaps. This prevents one block from overwriting old values that another block still needs.

Source: fdtd_tmz_backend_gpu_v4.cu.