# Shared memory (2D FDTD)

**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:

```cpp
__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; [[fdtd-cuda-tiling|tiling]] supplies the halo so the electric update does not depend on another block's progress.