webgpu-wgsl-compute-shaders

To prevent data races in WGSL compute shaders, use workgroup barriers and atomic functions to synchronize access to shared workgroup memory across invocations. Uniform barrier placement ensures deterministic execution for reductions and scans.

Workgroup size in WebGPU compute kernels defines the number of invocations per workgroup using @workgroup_size, while dispatch count in dispatchWorkgroups specifies how many workgroups to execute. Confusing the two leads to incorrect global and local invocation ID mapping.

Set up @compute and @workgroup_size for 2D image processing kernels by applying the @compute attribute and defining a two-dimensional workgroup size. You must then use local and global invocation IDs with proper bounds checks to safely process image pixels.

WGSL supports optional subgroup operations for parallel reductions, but requires proper subgroup feature detection before use. Subgroups optimize compute shader reductions and scans by enabling cooperative invocation grouping beyond standard workgroup memory barriers.

WGSL compute shaders fail or produce non-deterministic outputs when writing to shared workgroup memory without barriers because invocations execute concurrently. Uniform workgroup and storage barriers are mandatory to synchronize read and write operations safely.