demo · v144 · WebGPU
Parallel Reduction
Subgroup id + ballot is the bread-and-butter primitive for parallel reductions (sum / max / prefix-scan) without a workgroup-wide barrier. The shader below uses subgroup_invocation_id to compute a per-subgroup sum and write it out.
Origin trial
The
subgroups feature is in Chrome's WebGPU origin trial and behind --enable-dawn-features=allow_unsafe_apis in canaries. The page tries to request it and falls back to a CPU loop if rejected.navigator.gpu: checking…
click run
the API (WGSL)
enable subgroups;
@compute @workgroup_size(64)
fn main(@builtin(subgroup_invocation_id) sid: u32,
@builtin(subgroup_size) ssz: u32,
@builtin(global_invocation_id) gid: vec3<u32>) {
let v = input[gid.x];
let sum = subgroupAdd(v); // one barrier, whole subgroup
if (sid == 0u) output[gid.x / ssz] = sum;
}