WebGPU: Subgroup Size Control

at a glance

why it exists

Without this feature, a compute shader can only read the subgroup size the driver happened to choose (the subgroup_size built-in). On platforms whose GPUs support several subgroup sizes that creates three concrete problems, per the proposal:

1. Untestable code paths — correctness requires a path per possible size, but the developer cannot force a given size to test it.
2. Suboptimal performance — the driver-chosen size is not always the best for a given workload (the motivating use case is AI/compute-heavy shaders tuned per subgroup width).
3. Oversized derived parameters — workgroup sizes and workgroup-memory array sizes must be sized for the largest possible subgroup, wasting resources when a smaller size would do.

The native equivalents already existed: Vulkan's VK_EXT_subgroup_size_control (or Vulkan 1.3) and HLSL's [WaveSize()] (Shader Model 6.6). This feature brings the same control to WebGPU. Metal has no equivalent.

how it works

1. The page requests a device with "subgroup-size-control" in requiredFeatures. Per the spec, enabling this feature at device creation also enables "subgroups" on the device automatically.
2. A WGSL module gains enable subgroup_size_control; — valid only when the device has the feature, and only alongside enable subgroups;.
3. The compute entry point declares @subgroup_size(N). The pipeline will run with exactly that subgroup size, and the subgroup_size built-in always equals N.
4. Validation happens at shader creation (const-expression violations) and pipeline creation (range, workgroup-divisibility, and implementability checks) — see the error taxonomy.

syntax

The feature is a string in the GPUFeatureName enum — present verbatim in the spec's feature index and in Chromium's gpu_supported_features.idl:

// GPUFeatureName (spec feature index §25.23)
"subgroup-size-control"

Feature detection and device request follow the standard WebGPU pattern:

const adapter = await navigator.gpu.requestAdapter();

const requestedFeatures = [];
if (adapter.features.has("subgroup-size-control")) {
  requestedFeatures.push("subgroup-size-control");
} else {
  // Use an alternate code path (driver-chosen subgroup size).
}

const device = await adapter.requestDevice({ requiredFeatures });

inputs

outputs

• A GPUDevice on which device.features contains "subgroup-size-control" and "subgroups" (the latter auto-enabled).
• WGSL modules created on that device may use enable subgroup_size_control; and the @subgroup_size attribute (contract on the attribute page).
• Queryable bounds: adapter.info.subgroupMinSize / subgroupMaxSize delimit the valid request range (those members belong to the "subgroups" feature; requesting subgroup-size-control makes them meaningful on the same device).

errors

context and exposure

• Available wherever WebGPU is: Exposed=(Window, DedicatedWorker, SharedWorker, ServiceWorker), secure contexts only (per the WebGPU spec's GPU/navigator.gpu exposure).
• Adapter-dependent: presence in adapter.features varies by GPU/driver/OS — the same browser build exposes it on one machine and not another. Metal has no native equivalent (proposal: browsers may expose it on Metal “at their own risk”).
• Compute-only effect: the WGSL attribute is restricted to compute shader entry points; vertex/fragment pipelines are unaffected.

lifecycle

examples

Feature-detect, request, and use (JS + WGSL, per the proposal):

const adapter = await navigator.gpu.requestAdapter();
const requestedFeatures = [];
if (adapter.features.has("subgroup-size-control")) {
  requestedFeatures.push("subgroup-size-control");
} else {
  // Fall back to a driver-sized variant of the shader.
}
const device = await adapter.requestDevice({ requiredFeatures });

enable subgroups;
enable subgroup_size_control;

@compute @workgroup_size(64, 1, 1) @subgroup_size(32)
fn main(@builtin(subgroup_invocation_id) sg_id : u32,
        @builtin(subgroup_size) sg_size : u32) {
    // sg_size is guaranteed to be 32 in this pipeline.
}

tests

WebGPU conformance lives in the gpuweb/cts test suite rather than upstream WPT directories. A search of the CTS operation, shader-execution, and validation trees found no dedicated subgroup-size-control tests at fetch time (2026-07-29); the Intent to Ship does not record a WPT/CTS row beyond the standard process. Treat behavior as spec-defined but conformance-unverified in public test suites.

browser compatibility

Interim table. BCD has no entry for this feature — api/GPUSupportedFeatures.json carries feature_subgroups but no feature_subgroup-size-control key, and there is no @subgroup_size WGSL entry (verified 2026-07-29). Rows below are compiled from the linked primary sources, not from BCD.

security and privacy

• No new query surface — the working group deliberately declined new adapter limits (gpuweb#6241), so the feature adds no adapter-info fields. The pre-existing subgroupMinSize/subgroupMaxSize queries (part of "subgroups") already reveal the hardware's subgroup-size bounds.
• Pipeline-creation oracle (minor) — because supported sizes are not enumerable, a page can only probe implementability by attempting pipeline creation and observing success/uncategorized-error. This is a slow, coarse signal of GPU capability, comparable in kind to the existing bounds queries; no spec mitigation is defined.
• No new shader capabilities — the attribute constrains scheduling geometry; it grants no additional memory access, precision, or instruction surface beyond what "subgroups" already provides.
• Availability variance — like every optional GPU feature, presence/absence contributes to adapter fingerprinting entropy already inherent in adapter.features.

specifications

see also

• @subgroup_size attribute reference — the WGSL grammar, parameters, and error taxonomy (this feature's real surface)
• Chrome Platform Status — WebGPU: Subgroup Size Control (API record)
• blink-dev — Intent to Ship: WebGPU: Subgroup Size Control
• MDN — GPUAdapterInfo: subgroupMaxSize (the query side, which MDN covers; MDN has no page for this feature — verified 2026-07-29)
• webstatus.dev — WebGPU subgroups (related; Baseline limited)
• Chrome Platform Showcase — interactive demos for this feature