demo · v144

Device Probe

Subgroup size varies wildly: Nvidia warps are 32, AMD wavefronts 32 or 64, Apple subgroups can be 32, Intel 8/16/32. This probe asks your real device for its minSubgroupSize/maxSubgroupSize, lists every subgroup feature your adapter reports, and visualises the lane numbering for a small workgroup so you can see subgroup_invocation_id in action.

WebGPU: checking…

your adapter

subgroup features

subgroup_invocation_id visualisation

a 64-thread workgroup grouped by detected subgroup size. lane 0 of each subgroup is highlighted.

probing…

the WGSL

enable subgroup_id;

@compute @workgroup_size(64)
fn main(
  @builtin(subgroup_invocation_id) inv: u32,
  @builtin(subgroup_size) size: u32,
  @builtin(subgroup_id) sg: u32
) {
  // inv  = 0..size-1 within this subgroup
  // sg   = which subgroup within the workgroup
  // size = detected at runtime
}

see also