enable chromium_experimental_subgroups;

@group(0) @binding(0) var<storage, read_write> output : array<u32>;

@compute @workgroup_size(1)
fn main(@builtin(subgroup_invocation_id) subgroup_invocation_id : u32, @builtin(subgroup_size) subgroup_size : u32) {
  output[subgroup_invocation_id] = subgroup_size;
}