gpuweb · Jiawei-Shao · Feb 27, 2026 · Feb 27, 2026 · Feb 27, 2026 · Feb 27, 2026
diff --git a/correspondence/index.bs b/correspondence/index.bs
@@ -301,15 +301,15 @@ User agents are not required to use these formulas and may expose whatever they
         <td>`viewportBoundsRange[0]` (= -2 &times; `max(maxViewportDimensions[0..1])`)
         <td><p class=issue>*No documented limit?*
         <td>-32768 = `D3D12_VIEWPORT_BOUNDS_MIN`
-        
+
         Note: equal to -2 &times; `maxTextureDimension2D`
     <tr>
         <th>Max Viewport Bounds (implied)
         <td>[#373](https://github.com/gpuweb/gpuweb/issues/373)
         <td>`viewportBoundsRange[1]` (= 2 &times; `max(maxViewportDimensions[0..1])` - 1)
         <td><p class=issue>*No documented limit?*
         <td>32767 = `D3D12_VIEWPORT_BOUNDS_MAX`
-        
+
         Note: equal to 2 &times; `maxTextureDimension2D` - 1
 </table>
 
@@ -325,3 +325,32 @@ Warning:
 [Imagination](https://github.com/gpuweb/gpuweb/issues/3631#issuecomment-1498747606) drivers.
 On these drivers, the combined limit may need to be ignored.
 </p>
+
+## `subgroup-size-control` ## {#subgroup-size-control}
+
+The `subgroup-size-control` feature allows the use of the WGSL `subgroup_size` attribute
+in compute shaders to request a specific subgroup size for pipeline creation.
+
+No explicit limits (such as min/max subgroup size or max workgroup subgroup count) are exposed
+on `GPUAdapterInfo` for this feature
+([#6241](https://github.com/gpuweb/gpuweb/issues/6241)).
+Instead, if the implementation cannot create a pipeline with the requested subgroup size,
+it results in an uncategorized error during pipeline creation.
+At least one power-of-two subgroup size between `subgroupMinSize` and `subgroupMaxSize` must
+be supported.
+
+Note:
+The native APIs do expose limits related to explicit subgroup size control:
+- Vulkan: `VkPhysicalDeviceSubgroupSizeControlProperties::{minSubgroupSize, maxSubgroupSize, maxComputeWorkgroupSubgroups}`
+- D3D12: `D3D12_FEATURE_DATA_D3D12_OPTIONS1::{waveLaneCountMin, waveLaneCountMax}`
+
+These are not surfaced in WebGPU because:
+- The D3D12 `waveLaneCountMax` is not reliable according to [DirectXShaderCompiler Wiki](https://github.com/microsoft/DirectXShaderCompiler/wiki/Wave-Intrinsics/#caps-flags).
+- The D3D12 `waveLaneCountMin` may differ from the actual minimum subgroup sizes used in fragment shaders on some Intel GPUs.
+- `maxComputeWorkgroupSubgroups` only exists on Vulkan and has no D3D12 equivalent.
+
+Metal does not natively support explicit subgroup size control. According to the
+[Metal documentation](https://developer.apple.com/documentation/apple-silicon/porting-your-metal-code-to-apple-silicon#Determine-the-SIMD-Group-Size-at-Runtime):
+"The size of a SIMD group varies between different GPUs, particularly Mac GPUs. Don't assume the size of SIMD groups."
+Browsers may choose to expose the feature on Metal+Apple Silicon with a single constant subgroup
+size, but they do so at their own risk since the SIMD width is not guaranteed.
diff --git a/proposals/subgroup-size-control.md b/proposals/subgroup-size-control.md
@@ -27,46 +27,25 @@ This feature has not been approved by the working group yet.
 | Metal | Not Supported | According to [Metal document](https://developer.apple.com/documentation/apple-silicon/porting-your-metal-code-to-apple-silicon#Determine-the-SIMD-Group-Size-at-Runtime): <br>"The size of a SIMD group varies between different GPUs, particularly Mac GPUs. Don't assume the size of SIMD groups." |
 
 
-Note that on Vulkan we need `computeFullSubgroups == VK_TRUE` because we should set `VK_PIPELINE_SHADER_STAGE_CREATE_REQUIRE_FULL_SUBGROUPS_BIT` when creating the compute pipelien to ensure the subgroup sizes must be launched with all invocations active in the compute stage.
+Note that
+* On Vulkan we need `computeFullSubgroups == VK_TRUE` because we should set `VK_PIPELINE_SHADER_STAGE_CREATE_REQUIRE_FULL_SUBGROUPS_BIT` when creating the compute pipeline to ensure the subgroup sizes must be launched with all invocations active in the compute stage.
+* Metal does not natively support controlling subgroup size, but the browsers may choose to expose the feature on Metal at their own risk.
 
-TODO: Shall we support this extension on the newer Apple silicons with a single acceptable subgroup size?
+2. No additional explicit limits are exposed on `GPUAdapterInfo`.
 
-2. Three new limitations for the WGSL attribute `subgroup_size`.
+Previously, three limits were proposed (`explicitComputeSubgroupMinSize`, `explicitComputeSubgroupMaxSize`, `maxComputeWorkgroupSubgroups`) to expose the range of valid subgroup sizes and workgroup subgroup counts. The working group decided not to expose them ([#6241](https://github.com/gpuweb/gpuweb/issues/6241)). The reasons:
+ * Most users of this feature are expert developers coding for a specific architecture who already know which subgroup sizes to use.
+ * The native API limits do not always accurately represent the range of usable sizes (e.g. on Intel Gen12, `waveLaneCountMin` is 8 but the minimum requestable compute subgroup size is 16; on D3D12, `waveLaneCountMax` is not reliable).
+ * `maxComputeWorkgroupSubgroups` only exists on Vulkan and has no D3D12 equivalent.
 
-(1) `minExplicitComputeSubgroupSize` specifies the minimum value that can be used as the attribute `subgroup_size`.
-
-| Platform | Implementation |
-|----------|------|
-|Vulkan|`VkPhysicalDeviceSubgroupSizeControlPropertiesEXT::minSubgroupSize` |
-|D3D12|`D3D12_FEATURE_DATA_D3D12_OPTIONS1::waveLaneCountMin` |
-
-(2) `maxExplicitComputeSubgroupSize` specifies the maximum value that can be used as the attribute `subgroup_size`.
-
-| Platform | Implementation |
-|----------|------|
-|Vulkan|`VkPhysicalDeviceSubgroupSizeControlPropertiesEXT::maxSubgroupSize` |
-|D3D12|`D3D12_FEATURE_DATA_D3D12_OPTIONS1::waveLaneCountMax` |
-
-(3) `maxComputeWorkgroupSubgroups` limits the total workgroup size when the attribute `subgroup_size` is used.
-
-| Platform | Implementation |
-|----------|------|
-|Vulkan|`VkPhysicalDeviceSubgroupSizeControlProperties.maxComputeWorkgroupSubgroups` |
-|D3D12| Not supported |
-
-Note that we need new limitations instead of the existing `subgroupMinSize` and `subgroupMaxSize` is because:
-1. D3D12 runtime validates `[WaveSize]` with `waveLaneCountMin` and `waveLaneCountMax`
-2. On D3D12 we don't always use `waveLaneCountMin` as `subgroupMinSize` because on some Intel GPUs, it is possible to run some pixel shaders with wave lane count 8, while on that platform `waveLaneCountMin` is 16, meaning in compute shaders the wave lane count will always be at least 16.
-3. On D3D12 we don't always use `waveLaneCountMax` as `subgroupMaxSize` because in [D3D12 document](https://github.com/Microsoft/DirectXShaderCompiler/wiki/Wave-Intrinsics#:~:text=UINT%20WaveLaneCountMax) "the WaveLaneCountMax queried from D3D12 API is not reliable and the meaning is unclear.
+Instead, if the implementation cannot create a pipeline with the requested subgroup size (e.g. due to the size being outside the supported range, register pressure, or hardware-specific workgroup subgroup count limits), it results in an **uncategorized error** during pipeline creation. The implementation must support at least one power-of-two subgroup size between `subgroupMinSize` and `subgroupMaxSize` (from `GPUAdapterInfo`).
 
 
 # Behavior
  * The attribute `subgroup_size` is restricted to `compute` shaders (in HLSL `[WaveSize()]` is [only supported in compute shaders](https://microsoft.github.io/DirectX-Specs/d3d/HLSL_SM_6_6_WaveSize.html#hlsl-attribute)).
  * The parameter must be a const-expression or an override-expression that resolves to an `i32` or `u32`.
  * The parameter must be must be a power-of-two (required by [D3D12](https://microsoft.github.io/DirectX-Specs/d3d/HLSL_SM_6_6_WaveSize.html#allowed-wave-sizes)).
- * The parameter must be greater than or equal to the `minExplicitComputeSubgroupSize` on the current adapter (required by [D3D12](https://microsoft.github.io/DirectX-Specs/d3d/HLSL_SM_6_6_WaveSize.html#runtime-validation)).
- * The parameter must be less than or equal to the `maxExplicitComputeSubgroupSize` on the current adapter (required by [D3D12](https://microsoft.github.io/DirectX-Specs/d3d/HLSL_SM_6_6_WaveSize.html#runtime-validation)).
- * The total workgroup size (`workgroupSize.x * workgroupsize.y * workgroupsize.z`) must be less than or equal to the product of the attribute `subgroup_size` and `maxComputeWorkgroupSubgroups` (required by [Vulkan](https://docs.vulkan.org/refpages/latest/refpages/source/VkPipelineShaderStageCreateInfo.html#VUID-VkPipelineShaderStageCreateInfo-pNext-02756)).
+ * If the implementation cannot create a pipeline with the requested subgroup size (e.g. due to the size being outside the supported range, register pressure, or hardware-specific workgroup subgroup count limits), it results in an uncategorized error during pipeline creation. The implementation must support at least one power-of-two subgroup size between `subgroupMinSize` and `subgroupMaxSize`.
  * `workgroupSize.x` must be a multiple of the attribute `subgroup_size` (required by [Vulkan](https://docs.vulkan.org/refpages/latest/refpages/source/VkPipelineShaderStageCreateInfo.html#VUID-VkPipelineShaderStageCreateInfo-pNext-02757)).
 
 # WGSL Specification

diff --git a/spec/index.bs b/spec/index.bs
@@ -144,6 +144,9 @@ spec: WGSL; urlPrefix: https://gpuweb.github.io/gpuweb/wgsl/#
             text: dual_source_blending; url: extension-dual_source_blending
             text: subgroups; url: extension-subgroups
             text: primitive_index; url: extension-primitive_index
+            text: subgroup_size_control; url: extension-subgroup_size_control
+        for: attribute
+            text: subgroup_size; url: subgroup-size-attr
         for: language-extension
             text: readonly_and_readwrite_storage_textures; url: language_extension-readonly_and_readwrite_storage_textures
             text: packed_4x8_integer_dot_product; url: language_extension-packed_4x8_integer_dot_product
@@ -3059,6 +3062,7 @@ enum GPUFeatureName {
     "texture-formats-tier2",
     "primitive-index",
     "texture-component-swizzle",
+    "subgroup-size-control",
 };
 </script>
 
@@ -17449,6 +17453,8 @@ This feature adds the following [=optional API surfaces=]:
 
 Allows the use of the subgroup and quad operations in WGSL.
 
+Note: This feature is automatically enabled by {{GPUFeatureName/"subgroup-size-control"}}.
+
 This feature adds no [=optional API surfaces=], but the following entries of {{GPUAdapterInfo}}
 expose real values whenever the feature is available on the adapter:
 - {{GPUAdapterInfo/subgroupMinSize}}
@@ -17554,6 +17560,19 @@ This feature adds the following [=optional API surfaces=]:
 - New {{GPUTextureViewDescriptor}} dictionary members:
     - {{GPUTextureViewDescriptor/swizzle}}
 
+<h3 id=dom-gpufeaturename-subgroup-size-control data-dfn-type=enum-value data-dfn-for=GPUFeatureName>`"subgroup-size-control"`
+</h3>
+
+Allows the use of the [=attribute/subgroup_size=] attribute in WGSL to control compute pipeline subgroup size.
+
+Enabling {{GPUFeatureName/"subgroup-size-control"}} at device creation will enable
+{{GPUFeatureName/"subgroups"}}.
+
+This feature adds no [=optional API surfaces=].
+
+- New WGSL extensions:
+    - [=extension/subgroup_size_control=]
+
 # Appendices # {#appendices}
 
 ## Texture Format Capabilities ## {#texture-format-caps}

diff --git a/wgsl/index.bs b/wgsl/index.bs
@@ -289,6 +289,7 @@ spec: WebGPU; urlPrefix: https://gpuweb.github.io/gpuweb/#
     type: enum-value
         for: GPUFeatureName
             text: "primitive-index"; url: dom-gpufeaturename-primitive-index
+            text: "subgroup-size-control"; url: dom-gpufeaturename-subgroup-size-control
 </pre>
 
 # Introduction # {#intro}
@@ -1410,6 +1411,7 @@ The [=syntax/attribute=] names are:
 * <a for=attribute lt=size>`'size'`</a>
 * <a for=attribute lt=vertex>`'vertex'`</a>
 * <a for=attribute lt=workgroup_size>`'workgroup_size'`</a>
+* <a for=attribute lt=subgroup_size>`'subgroup_size'`</a>
 
 ### Built-in Value Names ### {#builtin-value-names}
 
@@ -1488,6 +1490,7 @@ The [=enable-extension=] names are:
 * <a for=extension lt=dual_source_blending>`'dual_source_blending'`</a>
 * <a for=extension lt=subgroups>`'subgroups'`</a>
 * <a for=extension lt=primitive_index>`'primitive_index'`</a>
+* <a for=extension lt=subgroup_size_control>`'subgroup_size_control'`</a>
 
 The valid [=language extension=] names are listed in [[#language-extensions-sec]] but in general have the same form as an [=identifier=]:
 
@@ -1864,6 +1867,12 @@ The valid [=enable-extensions=] are listed in the following table.
       <td>The built-in variable [=built-in values/primitive_index=] is valid to use in the WGSL
           module. Otherwise, using [=built-in values/primitive_index=] will result in a
           [=shader-creation error=].
+    <tr><td><dfn noexport dfn-for="extension">`subgroup_size_control`</dfn>
+      <td>{{GPUFeatureName/"subgroup-size-control"}}
+      <td>The attribute [=attribute/subgroup_size=] is valid to use in the WGSL module.
+        Otherwise, using [=attribute/subgroup_size=] will result in a [=shader-creation error=].
+        [=shader-creation error|Must=] be enabled together with the [=extension/subgroups=] extension.
+        Otherwise, enabling [=extension/subgroup_size_control=] will result in a [=shader-creation error=].
 </table>
 
 <div class='example wgsl using extensions expect-error' heading="Using hypothetical enable-extensions">
@@ -8518,6 +8527,7 @@ The [=return type=], if specified, [=shader-creation error|must=] be [=construct
 WGSL defines the following attributes that can be applied to function declarations:
  * the [=shader stage attributes=]: [=attribute/vertex=], [=attribute/fragment=], and [=attribute/compute=]
  * [=attribute/workgroup_size=]
+ * [=attribute/subgroup_size=]
 
 WGSL defines the following attributes that can be applied to function
 parameters and return types:
@@ -9215,6 +9225,41 @@ path: syntax/size_attr.syntax.bs.include
 
 </table>
 
+## `subgroup_size` ## {#subgroup-size-attr}
+
+<pre class=include>
+path: syntax/subgroup_size_attr.syntax.bs.include
+</pre>
+
+<table class='data builtin'>
+  <caption><dfn noexport dfn-for="attribute">`subgroup_size`</dfn> Attribute</caption>
+  <tr>
+    <td style="width:10%">Description
+    <td>Specifies the subgroup size for a compute shader invocation.
+
+    [=shader-creation error|Must=] only be applied to a [=compute shader stage|compute shader=] entry point function.
+
+  <tr>
+    <td>Requires
+    <td>[=shader-creation error|Must=] only be used when the [=extension/subgroup_size_control=] extension is enabled.
+
+  <tr>
+    <td>Parameters
+    <td>[=shader-creation error|must=] be a [=const-expression=] or an [=override-expression=] that [=type rules|resolves=] to an [=i32=] or [=u32=].<br>
+
+If the value is not a power of two, then:
+* It is a [=shader-creation error=] if the expression is a [=const-expression=].
+* It is a [=pipeline-creation error=] if the expression is an [=override-expression=].
+
+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.
+
+Pipeline creation may result in an [=uncategorized error=] due to register pressure
+or specific hardware limitations even when the value is between
+{{GPUAdapterInfo/subgroupMinSize}} and {{GPUAdapterInfo/subgroupMaxSize}}.
+
+</table>
+
 ## `workgroup_size` ## {#workgroup-size-attr}
 
 <pre class=include>
@@ -9385,6 +9430,7 @@ It will stabilize in a finite number of steps.
 WGSL defines the following attributes that can be applied to entry point declarations:
  * the [=shader stage attributes=]: [=attribute/vertex=], [=attribute/fragment=], and [=attribute/compute=]
  * [=attribute/workgroup_size=]
+ * [=attribute/subgroup_size=]
 
 <div class='example wgsl global-scope' heading='workgroup_size Attribute'>
   <xmp highlight=wgsl>
@@ -10119,6 +10165,9 @@ Each is described in detail in subsequent sections.
   <tr><td style="width:10%">Description
       <td>
       The [=subgroup size=] of current invocation's subgroup.
+
+      For compute shaders with a [=attribute/subgroup_size=] attribute, this value
+      is equal to the specified attribute value.
 </table>
 
 ##### `subgroup_id` ##### {#subgroup-id-builtin-value}

diff --git a/wgsl/syntax.bnf b/wgsl/syntax.bnf
@@ -174,6 +174,10 @@ workgroup_size_attr :
 | '@' 'workgroup_size' '(' expression ',' expression ',' expression ',' ? ')'
 ;
 
+subgroup_size_attr :
+  '@' 'subgroup_size' '(' expression ',' ? ')'
+;
+
 vertex_attr :
   '@' 'vertex'
 ;
@@ -202,6 +206,7 @@ attribute :
 | must_use_attr
 | size_attr
 | workgroup_size_attr
+| subgroup_size_attr
 | vertex_attr
 | fragment_attr
 | compute_attr

diff --git a/wgsl/syntax/attribute.syntax.bs.include b/wgsl/syntax/attribute.syntax.bs.include
@@ -31,6 +31,8 @@
 
     <span class="choice">|</span> [=syntax/workgroup_size_attr=]
 
+    <span class="choice">|</span> [=syntax/subgroup_size_attr=]
+
     <span class="choice">|</span> [=syntax/vertex_attr=]
 
     <span class="choice">|</span> [=syntax/fragment_attr=]

diff --git a/wgsl/syntax/subgroup_size_attr.syntax.bs.include b/wgsl/syntax/subgroup_size_attr.syntax.bs.include
@@ -0,0 +1,5 @@
+<div class='syntax' noexport='true'>
+  <dfn for=syntax>subgroup_size_attr</dfn> :
+
+    <span class="choice"></span> <a for=syntax_sym lt=attr>`'@'`</a> `'subgroup_size'` <a for=syntax_sym lt=paren_left>`'('`</a> [=syntax/expression=] <a for=syntax_sym lt=comma>`','`</a> ? <a for=syntax_sym lt=paren_right>`')'`</a>
+</div>
diff --git a/wgsl/wgsl.recursive.bs.include b/wgsl/wgsl.recursive.bs.include
@@ -58,6 +58,8 @@
  | `'@'` `'location'` `'('` [=recursive descent syntax/expression=] `','` ? `')'`
 
  | `'@'` `'size'` `'('` [=recursive descent syntax/expression=] `','` ? `')'`
+
+ | `'@'` `'subgroup_size'` `'('` [=recursive descent syntax/expression=] `','` ? `')'`
 </div>
 
 <div class='syntax' noexport='true'>