diff --git a/spec/index.bs b/spec/index.bs
index 263e2b9dcb..cb28c57f74 100644
--- a/spec/index.bs
+++ b/spec/index.bs
@@ -1456,6 +1456,7 @@ enum GPUFeatureName {
     "pipeline-statistics-query",
     "texture-compression-bc",
     "timestamp-query",
+    "subgroup-operations"
 };
 </script>
 
@@ -7706,6 +7707,10 @@ The following enums are supported if and only if the {{GPUFeatureName/"pipeline-
         * {{GPUQueryType/"pipeline-statistics"}}
 </dl>
 
+## <dfn dfn-type=enum-value dfn-for=GPUFeatureName>subgroup-operations</dfn> ## {#subgroup-operations}
+
+Issue: Define functionality when the {{GPUFeatureName/"subgroup-operations"}} [=feature=] is enabled.
+
 
 ## <dfn dfn-type=enum-value dfn-for=GPUFeatureName>texture-compression-bc</dfn> ## {#texture-compression-bc}
 
diff --git a/wgsl/index.bs b/wgsl/index.bs
index 106f7666bc..5c6e407ac6 100644
--- a/wgsl/index.bs
+++ b/wgsl/index.bs
@@ -4600,6 +4600,19 @@ See [[#builtin-inputs-outputs]] for how to declare a built-in variable.
       <td>vec3&lt;u32&gt;
       <td width="50%">The [=workgroup_size=] of the current entry point.
 
+  <tr><td>`subgroup_size`
+      <td>compute
+      <td>in
+      <td>u32
+      <td width="50%">The subgroup size of the current entry point.
+
+  <tr><td>`subgroup_invocation_index`
+      <td>compute
+      <td>in
+      <td>u32
+      <td width="50%">The current invocation's subgroup invocation index.
+          Must be in range [0, `subgroup_size`-1].
+
   <tr><td>`sample_index`
       <td>fragment
       <td>in
@@ -5787,6 +5800,37 @@ reduce a shader's memory bandwidth demand.
         See [[#floating-point-conversion]] for edge case behaviour.
 </table>
 
+## Subgroup built-in functions ## {#subgroup-builtin-functions}
+
+<table class='data'>
+  <thead>
+    <tr><td>Subgroup built-in functions<td>SPIR-V
+  </thead>
+  <tr><td>subgroupIsFirst() -&gt; bool<td>OpGroupNonUniformElect
+  <tr><td>subgroupAll(bool) -&gt; bool<td>OpGroupNonUniformAll
+  <tr><td>subgroupAny(bool) -&gt; bool<td>OpGroupNonUniformAny
+  <tr><td>subgroupBallot(bool) -&gt; vec4&lt;u32&gt;<td>OpGroupNonUniformBallot
+  <tr><td>subgroupBroadcastFirst(Integral) -&gt; Integral<td>OpGroupNonUniformBroadcastFirst		
+  <tr><td>subgroupBroadcastFirst(Floating) -&gt; Floating<td>OpGroupNonUniformBroadcastFirst
+  <tr><td>subgroupAdd(Integral) -&gt; Integral<td>OpGroupNonUniformIAdd with Reduce
+  <tr><td>subgroupAdd(Floating) -&gt; Floating<td>OpGroupNonUniformFAdd with Reduce
+  <tr><td>subgroupMul(Integral) -&gt; Integral<td>OpGroupNonUniformIMul with Reduce
+  <tr><td>subgroupMul(Floating) -&gt; Floating<td>OpGroupNonUniformFMul with Reduce
+  <tr><td>subgroupMin(Integral) -&gt; Integral<td>OpGroupNonUniformUMin or OpGroupNonUniformSMin with Reduce
+  <tr><td>subgroupMin(Floating) -&gt; Floating<td>OpGroupNonUniformFMin with Reduce
+  <tr><td>subgroupMax(Integral) -&gt; Integral<td>OpGroupNonUniformUMax or OpGroupNonUniformSMax with Reduce
+  <tr><td>subgroupMax(Floating) -&gt; Floating<td>OpGroupNonUniformFMax with Reduce
+  <tr><td>subgroupAnd(Integral) -&gt; Integral<td>OpGroupNonUniformBitwiseAnd
+  <tr><td>subgroupOr(Integral) -&gt; Integral<td>OpGroupNonUniformBitwiseOr
+  <tr><td>subgroupXor(Integral) -&gt; Integral<td>OpGroupNonUniformBitwiseXor
+  <tr><td>subgroupPrefixAdd(Integral) -&gt; Integral<td>OpGroupNonUniformIAdd with ExclusiveScan
+  <tr><td>subgroupPrefixAdd(Floating) -&gt; Floating<td>OpGroupNonUniformFAdd with ExclusiveScan
+  <tr><td>subgroupPrefixMul(Integral) -&gt; Integral<td>OpGroupNonUniformIMul with ExclusiveScan
+  <tr><td>subgroupPrefixMul(Floating) -&gt; Floating<td>OpGroupNonUniformFMul with ExclusiveScan
+</table>
+
+Note: Subgroup built-in functions exist if "subgroup-operations" is enabled in requestDevice.
+
 # Glossary # {#glossary}
 
 TODO: Remove terms unused in the rest of the specification.