demo · v134

Reduction: subgroups vs workgroup memory

The GPUweb explainer’s headline microbench: a sum-reduction across a workgroup of 256 lanes. The naive implementation needs a workgroup-shared array plus log2(N) rounds of barrier + half-write — that’s 8 barriers on a 256-lane group. With subgroupAdd() the same reduction is a single intrinsic per subgroup (32 lanes on most hardware) then a tiny cross-subgroup combine. Run both kernels below on the same input and compare timing.

probing WebGPU + chromium-experimental-subgroups…
Subgroups are still gated behind chrome://flags/#enable-webgpu-developer-features and the "chromium-experimental-subgroups" feature must be requested. Where unavailable, the page reports it and falls back to the workgroup-memory version only.

Run a reduction

2^20 = 1,048,576

workgroup memory + barriers

var<workgroup> partials: array<f32, 256>;

@compute @workgroup_size(256)
fn reduce_wg(@builtin(local_invocation_id) lid: vec3<u32>, @builtin(workgroup_id) wid: vec3<u32>) {
  let g = wid.x * 256u + lid.x;
  partials[lid.x] = inp[g];
  workgroupBarrier();
  var stride: u32 = 128u;
  while (stride > 0u) {
    if (lid.x < stride) { partials[lid.x] = partials[lid.x] + partials[lid.x + stride]; }
    workgroupBarrier();
    stride = stride / 2u;
  }
  if (lid.x == 0u) { out[wid.x] = partials[0]; }
}
— ms

subgroups

enable subgroups;
@compute @workgroup_size(256)
fn reduce_sg(@builtin(local_invocation_id) lid: vec3<u32>, @builtin(workgroup_id) wid: vec3<u32>) {
  let g = wid.x * 256u + lid.x;
  let lane_sum = subgroupAdd(inp[g]);
  // first lane of each subgroup writes to a small shared array
  if (subgroupElect()) { partials[lid.x / subgroup_size] = lane_sum; }
  workgroupBarrier();
  if (lid.x < 256u / subgroup_size) {
    let total = subgroupAdd(partials[lid.x]);
    if (lid.x == 0u) { out[wid.x] = total; }
  }
}
— ms

see also