Table of Contents

Thread blocks (2D FDTD)

A thread block groups CUDA threads that can synchronize and share memory. FDTD CUDA launches a two-dimensional grid of blocks, with each thread assigned a simulation cell:

$$ i=\mathrm{blockIdx.x}\,\mathrm{blockDim.x}+\mathrm{threadIdx.x},\qquad j=\mathrm{blockIdx.y}\,\mathrm{blockDim.y}+\mathrm{threadIdx.y} $$

The launch rounds up the number of blocks to cover the whole grid. Threads in a partial block can lie outside the simulation domain, so the kernels check their coordinates before accessing field arrays. In the tiled kernel, these threads still reach __syncthreads() so the block can synchronize correctly.

The CUDA variant name can specify the block dimensions, such as v4:16x16. The default is 32 by 8. The backend rejects dimensions whose product exceeds the device's maximum threads per block.

Block shape changes how threads access neighboring cells and how much shared memory the tiled kernel needs. Larger blocks also reduce the proportion of halo work, but may limit how many blocks can run concurrently. Compare shapes using the same simulation workload rather than assuming the largest block is fastest.