← WebGPU: Subgroup Size Control
v152 · wgsl · attribute · compute shaders
@subgroup_size attribute
The @subgroup_size WGSL attribute pins a compute shader's subgroup size: the pipeline runs with exactly the declared size, and the subgroup_size built-in always equals it. Enabled by the subgroup_size_control WGSL extension, which the "subgroup-size-control" GPU feature unlocks.
syntax
The extension is an enable extension (WGSL §4.1.1) — not a language feature, so it does not appear in navigator.gpu.wgslLanguageFeatures:
enable subgroups;
enable subgroup_size_control;
The attribute grammar (WGSL §12.14, verbatim):
subgroup_size_attr :
'@' 'subgroup_size' '(' expression ',' ? ')'
| Applies to | A compute shader entry point function only — “Must only be applied to a compute shader entry point function”; applying it to any other object (or to vertex/fragment entry points) is invalid |
|---|---|
| Requires | “Must only be used when the subgroup_size_control extension is enabled” — and the extension itself requires the "subgroup-size-control" device feature and a co-enabled subgroups extension |
| Parameters | One parameter: “a const-expression or an override-expression that resolves to an i32 or u32” |
inputs (the size expression)
| Constraint | Rule |
|---|---|
| Expression kind | const-expression or override-expression resolving to i32/u32. Override-expressions let one shader source serve several subgroup sizes, fixed per pipeline at creation |
| Power of two | Required (D3D12 constraint). A non-power-of-two const-expression is a shader-creation error; a non-power-of-two override-expression value is a pipeline-creation error |
| Range | Must not exceed GPUAdapterInfo.subgroupMaxSize or be below subgroupMinSize — violation is a pipeline-creation error |
| Workgroup divisibility | “A pipeline-creation error results if the x-dimension of the entry point's workgroup_size is not a multiple of the subgroup_size value” (Vulkan VUID-derived) |
| Implementability | Even an in-range, power-of-two, divisible size can fail: “Pipeline creation may result in an uncategorized error due to register pressure or specific hardware limitations” |
outputs (observable effect)
- The compute pipeline executes with the declared subgroup size on every dispatch.
- The
subgroup_sizebuilt-in value inside the shader always equals the attribute value (proposal: “the built-in valuesubgroup_sizewill always be the value of the attribute”). - Static knowledge of the size unlocks compile-time optimizations: exact unroll counts for subgroup operations, workgroup-memory layouts sized to the real (not maximum) subgroup count, and per-size tuned variants selected via override-expressions.
error taxonomy
| Violation | Error kind | Raised at |
|---|---|---|
enable subgroup_size_control; without the device feature | shader-creation error | Module validation |
subgroup_size_control enabled without subgroups | shader-creation error | Module validation (“Must be enabled together with the subgroups extension”) |
| Attribute on a non-compute entry point or other object | shader-creation error | Module validation |
Parameter not const-/override-expression, or not i32/u32 | shader-creation error | Module validation |
| Non-power-of-two const-expression | shader-creation error | Module validation |
| Non-power-of-two override-expression value | pipeline-creation error | createComputePipeline(Async) |
Below subgroupMinSize / above subgroupMaxSize | pipeline-creation error | Pipeline creation |
workgroup_size.x not a multiple of the size | pipeline-creation error | Pipeline creation |
| Register pressure / hardware workgroup-subgroup limits | uncategorized error | Pipeline creation — even for valid-looking sizes; have a fallback pipeline |
context and exposure
- Compute shaders only — mirroring HLSL
[WaveSize()], which is likewise compute-only. - Requires a device with
"subgroup-size-control"granted (see the feature context: secure contexts, Window + worker exposure, adapter-dependent). - The attribute composes with the subgroup built-ins (
subgroup_invocation_id,subgroup_size) and subgroup operations from thesubgroupsextension.
lifecycle
| Stage | What happens |
|---|---|
| Module scope | enable directives are module-wide; both subgroups and subgroup_size_control must be present in any module using the attribute |
| Shader creation | Static checks: extension enabled, placement, expression kind, const-expression power-of-two |
| Pipeline creation | Override-expression values are resolved and checked (power-of-two, min/max range, workgroup divisibility, implementability); the size is baked into the pipeline |
| Dispatch | Size is fixed for the pipeline's lifetime; retargeting another size means compiling another pipeline (override-expressions make this a constant override, not a source edit) |
examples
Fixed size (const-expression) — proposal example:
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 == 32 in this pipeline; workgroup is two full subgroups.
}
Per-target variants from one source (override-expression):
enable subgroups;
enable subgroup_size_control;
override sg_size : u32 = 32u; // default; overridable per pipeline
@compute @workgroup_size(128, 1, 1) @subgroup_size(sg_size)
fn main(@builtin(subgroup_invocation_id) sg_id : u32) {
// One source; pipelines built with sg_size = 16, 32, 64, ...
}
// JS: build a per-size pipeline, falling back if creation fails
const module = device.createShaderModule({ code: wgslSource });
let pipeline;
try {
pipeline = await device.createComputePipelineAsync({
layout: "auto",
compute: { module, entryPoint: "main", constants: { sg_size: 32 } },
});
} catch {
pipeline = await device.createComputePipelineAsync({
layout: "auto",
compute: { module, entryPoint: "main" }, // default override value
});
}
Source: fixed-size example verbatim from the proposal — Example usage; override-expression pattern derived from the WGSL parameter rules (gendn-derived, labelled as such)
browser compatibility
Interim table. No BCD entry exists for this attribute or extension (2026-07-29); rows compile the linked primary sources.
| Browser | Support | Evidence |
|---|---|---|
| Chrome | Enabled by default from 152 (desktop, Android, WebView) | milestone=152 listing; Intent to Ship |
| Edge | Mirrors Chrome (Chromium) | Chromium-based; no separate record |
| Firefox | Positive position; shipped support not recorded in public data | Intent thread (Gecko: Positive, via the WebGPU standardization process) |
| Safari | Positive position; shipped support not recorded in public data | Intent thread (WebKit: Positive, per spec approval) |
security and privacy
- The attribute is compile-time metadata: it grants no memory access, new instructions, or side channels beyond the
subgroupsfeature it accompanies. - Pipeline-creation success/failure probing (see the feature-level privacy notes) applies at the per-size level: a page can binary-search implementable sizes across a few pipeline creations — a coarse GPU-capability signal, not user data.
- Deterministic sizing can actually reduce timing variance attributable to driver choice, but nothing in the feature hides GPU identity; adapter fingerprinting considerations are unchanged from WebGPU generally.
see also
- WebGPU: Subgroup Size Control — the GPU feature, device contract, and compat picture
- WGSL spec — @subgroup_size · WGSL spec — subgroup_size_control extension
- gpuweb/gpuweb PR #5578 (merged 2026-06-23) · issue #5545
- Chrome Platform Status — 5077657663438848