demo · v134

WebGPU Subgroups

WebGPU exposes subgroup intrinsics — SIMD-style ops within a workgroup (broadcast, ballot, shuffle, reductions). Big speedups for parallel reductions and prefix sums. The demo below sums a buffer using a naive workgroup-shared loop and the subgroup intrinsic subgroupAdd, then reports time and result for both paths.

probing WebGPU + subgroups feature...

Try it

elements:

naive workgroup loop

—

subgroupAdd

—

no run yet

The code

// WGSL
enable subgroups;

@compute @workgroup_size(64)
fn main(@builtin(local_invocation_id) lid: vec3u,
        @builtin(workgroup_id) wid: vec3u) {
  let v = input[wid.x * 64 + lid.x];
  let partial = subgroupAdd(v);          // SIMD reduction within subgroup
  if (subgroupElect()) {                 // one lane per subgroup writes
    atomicAdd(&total, partial);
  }
}

Try it

naive workgroup loop

subgroupAdd

The code

see also