← WebGPU: Subgroup Size Control

v152 · wgsl · attribute · compute shaders

@subgroup_size attribute

Limited availability

  • Chrome · enabled by default 152
  • Edge · mirrors Chrome (Chromium)
  • Firefox · positive position, support unrecorded
  • Safari · positive position, support unrecorded

Usable when the device grants "subgroup-size-control" (see the feature page); the WGSL subgroup_size_control extension must be enabled together with subgroups.

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 toA 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
ParametersOne parameter: “a const-expression or an override-expression that resolves to an i32 or u32
Source: WGSL — subgroup_size_control extension; WGSL — @subgroup_size attribute

inputs (the size expression)

ConstraintRule
Expression kindconst-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 twoRequired (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
RangeMust 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)
ImplementabilityEven 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”
Source: WGSL — @subgroup_size attribute; proposal — Behavior; gpuweb#6241 (no enumerable size set)

outputs (observable effect)

Source: proposal — Motivation + Behavior; WGSL — @subgroup_size attribute

error taxonomy

ViolationError kindRaised at
enable subgroup_size_control; without the device featureshader-creation errorModule validation
subgroup_size_control enabled without subgroupsshader-creation errorModule validation (“Must be enabled together with the subgroups extension”)
Attribute on a non-compute entry point or other objectshader-creation errorModule validation
Parameter not const-/override-expression, or not i32/u32shader-creation errorModule validation
Non-power-of-two const-expressionshader-creation errorModule validation
Non-power-of-two override-expression valuepipeline-creation errorcreateComputePipeline(Async)
Below subgroupMinSize / above subgroupMaxSizepipeline-creation errorPipeline creation
workgroup_size.x not a multiple of the sizepipeline-creation errorPipeline creation
Register pressure / hardware workgroup-subgroup limitsuncategorized errorPipeline creation — even for valid-looking sizes; have a fallback pipeline
Source: WGSL — @subgroup_size attribute; WGSL — subgroup_size_control extension; proposal — Behavior

context and exposure

Source: proposal — Behavior + Native API Availability; WGSL — subgroup_size_control extension

lifecycle

StageWhat happens
Module scopeenable directives are module-wide; both subgroups and subgroup_size_control must be present in any module using the attribute
Shader creationStatic checks: extension enabled, placement, expression kind, const-expression power-of-two
Pipeline creationOverride-expression values are resolved and checked (power-of-two, min/max range, workgroup divisibility, implementability); the size is baked into the pipeline
DispatchSize 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)
Source: WGSL — @subgroup_size attribute; proposal — Behavior

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.

BrowserSupportEvidence
ChromeEnabled by default from 152 (desktop, Android, WebView)milestone=152 listing; Intent to Ship
EdgeMirrors Chrome (Chromium)Chromium-based; no separate record
FirefoxPositive position; shipped support not recorded in public dataIntent thread (Gecko: Positive, via the WebGPU standardization process)
SafariPositive position; shipped support not recorded in public dataIntent thread (WebKit: Positive, per spec approval)
Source: BCD api/GPUSupportedFeatures.json (no entry, verified 2026-07-29); chromestatus.com/feature/5077657663438848

security and privacy

Source: WGSL — @subgroup_size attribute; gpuweb#6241

see also