Shared memory is storage available to all threads in one CUDA block. The v4 field-update kernel uses it to hold newly computed $H_x$ and $H_y$ values, letting neighboring threads use those values for the electric update in the same kernel.

After threads fill the tiles and halos, the kernel synchronizes before reading neighboring samples:

__syncthreads();
if (inside) {
    const gpu::HAround h{
        .hy = i + 1 < fields.nx ? tileHy[c] : FloatType(0),
        .hy_left = i > 0 ? tileHy[c - 1] : FloatType(0),
        .hx = j + 1 < fields.ny ? tileHx[c] : FloatType(0),
        .hx_below = j > 0 ? tileHx[c - width] : FloatType(0),
    };
    fields.ez_next[k] = gpu::e_value(fields, coeffs, i, j, boundary, h);
}

All threads, including those outside a partial edge block, reach the barrier. Only valid cells access the field arrays afterward. The zeros substitute for magnetic samples outside the domain.

The two shared tiles occupy:

$$ M=2(b_x+1)(b_y+1)\,\mathrm{sizeof}(\mathrm{FloatType}) $$

This is 2,376 bytes for the default $32\times8$ block with float fields. Storage is private to each block; tiling supplies the halo so the electric update does not depend on another block's progress.