Add extension subgroup-size-control

This patch adds a new WebGPU and WGSL extension `subgroup-size-control` based on `proposals/subgroup-size-control.md`.
gpuweb · Jiawei-Shao · Feb 27, 2026 · Feb 27, 2026 · Feb 27, 2026 · Feb 27, 2026
commit f17c4e85063506d51a5e045204b0e7c4d1637cd2
diff --git a/spec/index.bs b/spec/index.bs
@@ -142,6 +142,7 @@ 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: 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
@@ -1979,6 +1980,9 @@ interface GPUAdapterInfo {
     readonly attribute DOMString description;
     readonly attribute unsigned long subgroupMinSize;
     readonly attribute unsigned long subgroupMaxSize;
+    readonly attribute unsigned long minExplicitComputeSubgroupSize;
+    readonly attribute unsigned long maxExplicitComputeSubgroupSize;
+    readonly attribute unsigned long maxComputeWorkgroupSubgroups;
     readonly attribute boolean isFallbackAdapter;
 };
 </script>
@@ -2023,6 +2027,22 @@ interface GPUAdapterInfo {
         If the {{GPUFeatureName/"subgroups"}} feature is supported, the maximum
         supported [=builtin-value/subgroup size=] for the [=adapter=].
 
+    : <dfn>minExplicitComputeSubgroupSize</dfn>
+    ::
+        If the {{GPUFeatureName/"subgroup-size-control"}} feature is supported, the minimum
+        value usable for [=attribute/subgroup_size=] in compute shaders on the [=adapter=].
+
+    : <dfn>maxExplicitComputeSubgroupSize</dfn>
+    ::
+        If the {{GPUFeatureName/"subgroup-size-control"}} feature is supported, the maximum
+        value usable for [=attribute/subgroup_size=] in compute shaders on the [=adapter=].
+
+    : <dfn>maxComputeWorkgroupSubgroups</dfn>
+    ::
+        If the {{GPUFeatureName/"subgroup-size-control"}} feature is supported, the maximum
+        number of subgroups per workgroup when [=attribute/subgroup_size=] is used in compute shaders
+        on the [=adapter=].
+
     : <dfn>isFallbackAdapter</dfn>
     ::
         Whether the [=adapter=] is a [=fallback adapter=].
@@ -2070,6 +2090,30 @@ interface GPUAdapterInfo {
         for the property which do not distinguish different devices, but are still usable
         (e.g. use the default value of 128 for all devices).
 
+    1. If {{GPUFeatureName/"subgroup-size-control"}} is supported, set
+        {{GPUAdapterInfo/minExplicitComputeSubgroupSize}} to the smallest supported explicit compute
+        subgroup size. Otherwise, set this value to 4.
+
+        Note: To preserve privacy, the user agent may choose to not support some features or provide values
+        for the property which do not distinguish different devices, but are still usable
+        (e.g. use the default value of 4 for all devices).
+
+    1. If {{GPUFeatureName/"subgroup-size-control"}} is supported, set
+        {{GPUAdapterInfo/maxExplicitComputeSubgroupSize}} to the largest supported explicit compute
+        subgroup size. Otherwise, set this value to 128.
+
+        Note: To preserve privacy, the user agent may choose to not support some features or provide values
+        for the property which do not distinguish different devices, but are still usable
+        (e.g. use the default value of 128 for all devices).
+
+    1. If {{GPUFeatureName/"subgroup-size-control"}} is supported, set
+        {{GPUAdapterInfo/maxComputeWorkgroupSubgroups}} to the largest supported number of
+        subgroups per compute workgroup with explicit subgroup size. Otherwise, set this value to 32.
+
+        Note: To preserve privacy, the user agent may choose to not support some features or provide values
+        for the property which do not distinguish different devices, but are still usable
+        (e.g. use the default value of 32 for all devices).
+
     1. Set |adapterInfo|.{{GPUAdapterInfo/isFallbackAdapter}} to |adapter|.{{adapter/[[fallback]]}}.
 
     1. Return |adapterInfo|.
@@ -3043,6 +3087,7 @@ enum GPUFeatureName {
     "clip-distances",
     "dual-source-blending",
     "subgroups",
+    "subgroup-size-control",
     "texture-formats-tier1",
     "texture-formats-tier2",
     "primitive-index",
@@ -8299,6 +8344,13 @@ dictionary GPUComputePipelineDescriptor
                             [|device|.limits.{{supported limits/maxComputeWorkgroupSizeX}},
                             |device|.limits.{{supported limits/maxComputeWorkgroupSizeY}},
                             |device|.limits.{{supported limits/maxComputeWorkgroupSizeZ}}].
+                        1. If |entryPoint| has a `subgroup_size` attribute with value |subgroupSize|:
+                            1. |subgroupSize| |must| be &ge; |this|.{{GPUDevice/adapterInfo}}.{{GPUAdapterInfo/minExplicitComputeSubgroupSize}}.
+                            1. |subgroupSize| |must| be &le; |this|.{{GPUDevice/adapterInfo}}.{{GPUAdapterInfo/maxExplicitComputeSubgroupSize}}.
+                            1. |entryPoint|'s `workgroup_size` X dimension |must| be a multiple of |subgroupSize|.
+                            1. |entryPoint| |must| use &le;
+                                <code>|subgroupSize| &times; |this|.{{GPUDevice/adapterInfo}}.{{GPUAdapterInfo/maxComputeWorkgroupSubgroups}}</code>
+                                per workgroup.
                     </div>
 
                 1. If any [=pipeline-creation error|pipeline-creation=] [=uncategorized errors=]
@@ -17378,6 +17430,20 @@ 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 [=subgroup_size=] attribute in WGSL to control over compute pipeline subgroup size.
+
+This feature adds no [=optional API surfaces=], but the following entries of {{GPUAdapterInfo}}
+expose bounds for explicit compute subgroup size control:
+- {{GPUAdapterInfo/minExplicitComputeSubgroupSize}}
+- {{GPUAdapterInfo/maxExplicitComputeSubgroupSize}}
+- {{GPUAdapterInfo/maxComputeWorkgroupSubgroups}}
+
+- 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
@@ -1400,6 +1400,7 @@ The [=syntax/attribute=] names are:
 * <a for=attribute lt=blend_src>`'blend_src'`</a>
 * <a for=attribute lt=must_use>`'must_use'`</a>
 * <a for=attribute lt=size>`'size'`</a>
+* <a for=attribute lt=subgroup_size>`'subgroup_size'`</a>
 * <a for=attribute lt=vertex>`'vertex'`</a>
 * <a for=attribute lt=workgroup_size>`'workgroup_size'`</a>
 
@@ -1478,6 +1479,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=]:
 
@@ -1852,6 +1854,10 @@ 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>[[WebGPU#dom-gpufeaturename-subgroup-size-control|"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=].
 </table>
 
 <div class='example wgsl using extensions expect-error' heading="Using hypothetical enable-extensions">
@@ -8472,6 +8478,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/subgroup_size=]
  * [=attribute/workgroup_size=]
 
 WGSL defines the following attributes that can be applied to function
@@ -9170,6 +9177,39 @@ path: syntax/size_attr.syntax.bs.include
 
 </table>
 
+## `subgroup_size` ## {#subgroup-size-attr}
+
+<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.
+    [=shader-creation error|Must not=] be applied to any other object.
+    [=shader-creation error|Must=] only be used when the [=extension/subgroup_size_control=] extension is enabled.
+
+  <tr>
+    <td>Parameters
+    <td>Takes one parameter.
+
+        The parameter [=shader-creation error|must=] be a [=const-expression=] or an [=override-expression=].
+        The parameter [=shader-creation error|must=] be of type [=i32=] or [=u32=].
+
+        A [=shader-creation error=] results if the parameter is a [=const-expression=]
+        that evaluates to a value that is not a power of two.
+
+        A [=pipeline-creation error=] results if the parameter evaluates to a value that is not
+        a power of two, is less than the lower bound specified by the WebGPU API,
+        or exceeds the upper bound specified by the WebGPU API (see [[WebGPU#limits]]).
+
+        A [=pipeline-creation error=] also results if the compute entry point's
+        [=attribute/workgroup_size=] x-dimension is not a multiple of this value,
+        or if the total workgroup size exceeds this value multiplied by the
+        WebGPU API bound on the maximum number of subgroups per workgroup.
+
+</table>
+
 ## `workgroup_size` ## {#workgroup-size-attr}
 
 <pre class=include>
@@ -9339,6 +9379,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/subgroup_size=]
  * [=attribute/workgroup_size=]
 
 <div class='example wgsl global-scope' heading='workgroup_size Attribute'>
@@ -10000,6 +10041,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}