CUDA supplies blockIdx and threadIdx per invocation; blockDim comes from the launch size above. These settings replace the desktop host's launch setup. Dynamic shared bytes supplies the allocation for extern __shared__ arrays and recompiles the shader. Suggested inputs are synthetic: check shapes and counts for your kernel. Whole-block bounds enforce counts divisible by the block size before making supported early returns uniform. Optional feedback maps destination buffers to source buffers for GPU copies after each step. A scalar volume preview uses volume.dimensions for the interior XYZ size and volume.halo for boundary cells on each side. Optional copies run immediately after their entry, with source/target buffer names and byteLength (plus optional sourceOffset/targetOffset in bytes). Additional passes run after the main entry, with their own entry, block, groups and buffer-name bindings. Use the shader-pass selector to compare each generated shader. Edit records, fill patterns, scalars, block counts, and output below.