v145 · Web APIs · WebGPU

WebGPU: subgroup_uniformity feature

Chrome 145 adds the subgroup_uniformity WGSL feature, relaxing the strict uniformity analysis for subgroup operations — allowing WGSL shaders to use subgroup built-ins in contexts where strict uniform analysis would incorrectly reject them.

background

WGSL's strict uniformity analysis ensures that values passed to certain operations (like subgroup operations, texture sampling) are provably uniform across a workgroup or subgroup. In practice, this analysis is sometimes more conservative than necessary — it rejects valid shader patterns where the developer knows the values are uniform at runtime.

Enabling subgroup_uniformity as a WGSL feature relaxes the analysis for subgroup operations, allowing more compute shaders to compile that previously failed the analysis check.

concepts

  1. Subgroup Demo

    Checks subgroups feature support on the GPU adapter and shows how to opt in to subgroup_uniformity when creating a device.

  2. WGSL Reference

    Subgroup built-in functions available in WGSL, what uniformity means in this context, and examples of shader patterns that benefit from the relaxed analysis.

  3. Divergence Stress

    Parameterised compute dispatch that lets you sweep the divergence ratio in a workgroup and watch the verdict change. Falls back to a CPU simulator when WebGPU isn't available.

  4. Shader Playground

    WGSL editor with live compilation via createShaderModule() + getCompilationInfo(). Toggle enable subgroups and enable subgroup_uniformity on/off and compile — see driver messages that explain exactly why strict uniformity analysis accepts or rejects each shader pattern.

the change

// Request the subgroups feature when creating a device:
const adapter = await navigator.gpu.requestAdapter();
const device = await adapter.requestDevice({
  requiredFeatures: ['subgroups'],  // enables subgroup built-ins
});

// Enable subgroup_uniformity in WGSL source:
const shaderModule = device.createShaderModule({
  code: `
    enable subgroups;
    enable subgroup_uniformity;  // Chrome 145+: relaxed uniformity analysis

    @compute @workgroup_size(64)
    fn main(@builtin(local_invocation_index) lid: u32) {
      // subgroupBroadcast may require uniform first arg in strict analysis
      // With subgroup_uniformity, this can now compile:
      let val = subgroupBroadcast(lid, 0u);
    }
  `,
});

references