diff --git a/CHANGELOG.md b/CHANGELOG.md
index 6a1f20c76..e8ab28d78 100644
--- a/CHANGELOG.md
+++ b/CHANGELOG.md
@@ -219,6 +219,12 @@ By @cwfitzgerald in [#8162](https://github.com/gfx-rs/wgpu/pull/8162).
 
 - [wgsl-in] Allow a trailing comma in `@blend_src(…)` attributes. By @ErichDonGubler in [#8137](https://github.com/gfx-rs/wgpu/pull/8137).
 
+### Documentation
+
+#### General
+
+- Clarify that subgroup barriers require both the `SUBGROUP` and `SUBGROUP_BARRIER` features / capabilities. By @andyleiserson in TBD.
+
 ## v26.0.4 (2025-08-07)
 
 ### Bug Fixes
diff --git a/naga/src/valid/mod.rs b/naga/src/valid/mod.rs
index 24938684a..426b3d637 100644
--- a/naga/src/valid/mod.rs
+++ b/naga/src/valid/mod.rs
@@ -131,13 +131,26 @@ bitflags::bitflags! {
         const CUBE_ARRAY_TEXTURES = 1 << 15;
         /// Support for 64-bit signed and unsigned integers.
         const SHADER_INT64 = 1 << 16;
-        /// Support for subgroup operations.
-        /// Implies support for subgroup operations in both fragment and compute stages,
-        /// but not necessarily in the vertex stage, which requires [`Capabilities::SUBGROUP_VERTEX_STAGE`].
+        /// Support for subgroup operations (except barriers) in fragment and compute shaders.
+        ///
+        /// Subgroup operations in the vertex stage require
+        /// [`Capabilities::SUBGROUP_VERTEX_STAGE`] in addition to `Capabilities::SUBGROUP`.
+        /// (But note that `create_validator` automatically sets
+        /// `Capabilities::SUBGROUP` whenever `Features::SUBGROUP_VERTEX` is
+        /// available.)
+        ///
+        /// Subgroup barriers require [`Capabilities::SUBGROUP_BARRIER`] in addition to
+        /// `Capabilities::SUBGROUP`.
         const SUBGROUP = 1 << 17;
-        /// Support for subgroup barriers.
+        /// Support for subgroup barriers in compute shaders.
+        ///
+        /// Requires [`Capabilities::SUBGROUP`]. Without it, enables nothing.
         const SUBGROUP_BARRIER = 1 << 18;
-        /// Support for subgroup operations in the vertex stage.
+        /// Support for subgroup operations (not including barriers) in the vertex stage.
+        ///
+        /// Without [`Capabilities::SUBGROUP`], enables nothing. (But note that
+        /// `create_validator` automatically sets `Capabilities::SUBGROUP`
+        /// whenever `Features::SUBGROUP_VERTEX` is available.)
         const SUBGROUP_VERTEX_STAGE = 1 << 19;
         /// Support for [`AtomicFunction::Min`] and [`AtomicFunction::Max`] on
         /// 64-bit integers in the [`Storage`] address space, when the return
@@ -206,7 +219,11 @@ bitflags::bitflags! {
     #[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
     #[derive(Clone, Copy, Debug, Default, Eq, PartialEq)]
     pub struct SubgroupOperationSet: u8 {
-        /// Elect, Barrier
+        /// Barriers
+        // Possibly elections, when that is supported.
+        // https://github.com/gfx-rs/wgpu/issues/6042#issuecomment-3272603431
+        // Contrary to what the name "basic" suggests, HLSL/DX12 support the
+        // other subgroup operations, but do not support subgroup barriers.
         const BASIC = 1 << 0;
         /// Any, All
         const VOTE = 1 << 1;
diff --git a/naga/tests/naga/wgsl_errors.rs b/naga/tests/naga/wgsl_errors.rs
index 6e5059e2e..249d77c22 100644
--- a/naga/tests/naga/wgsl_errors.rs
+++ b/naga/tests/naga/wgsl_errors.rs
@@ -1166,6 +1166,16 @@ fn validation_error(
         .map_err(|e| e.into_inner()) // TODO(https://github.com/gfx-rs/wgpu/issues/8153): Add tests for spans
 }
 
+/// Check that a shader validates successfully.
+///
+/// In a few tests it is useful to check conditions where a validation error
+/// should be absent alongside conditions where it should be present. This
+/// wrapper is less confusing than `validation_error().unwrap()`.
+#[track_caller]
+fn no_validation_error(source: &str, caps: naga::valid::Capabilities) {
+    validation_error(source, caps).unwrap();
+}
+
 #[test]
 fn int64_capability() {
     check_validation! {
@@ -3585,6 +3595,7 @@ fn issue7165() {
         fn invalid_return_type(a: Struct) -> i32 { return a; }
     ";
 
+    // We need the span for the error, so have to invoke manually.
     let module = naga::front::wgsl::parse_str(shader).unwrap();
     let err = naga::valid::Validator::new(
         naga::valid::ValidationFlags::all(),
@@ -3834,6 +3845,171 @@ fn const_eval_value_errors() {
     assert!(variant("f32(abs(-9223372036854775807 - 1))").is_ok());
 }
 
+#[test]
+fn subgroup_capability() {
+    // Some of these tests should be `check_extension_validation` tests that
+    // also check handling of the enable directive, but that handling is not
+    // currently correct. https://github.com/gfx-rs/wgpu/issues/8202
+
+    // Non-barrier subgroup operations...
+
+    // ...in fragment and compute shaders require [`Capabilities::SUBGROUP`]`.
+    for stage in [naga::ShaderStage::Fragment, naga::ShaderStage::Compute] {
+        let stage_attr = match stage {
+            naga::ShaderStage::Fragment => "@fragment",
+            naga::ShaderStage::Compute => "@compute @workgroup_size(1)",
+            _ => unreachable!(),
+        };
+        check_one_validation! {
+            &format!("
+                {stage_attr}
+                fn main() {{
+                    subgroupBallot();
+                }}
+            "),
+            Err(naga::valid::ValidationError::EntryPoint {
+                stage: err_stage,
+                source: naga::valid::EntryPointError::Function(
+                    naga::valid::FunctionError::MissingCapability(Capabilities::SUBGROUP)
+                ),
+                ..
+            }) if *err_stage == stage
+        }
+    }
+
+    // ...in fragment and compute shaders require *only* [`Capabilities::SUBGROUP`]`.
+    for stage in [naga::ShaderStage::Fragment, naga::ShaderStage::Compute] {
+        let stage_attr = match stage {
+            naga::ShaderStage::Fragment => "@fragment",
+            naga::ShaderStage::Compute => "@compute @workgroup_size(1)",
+            _ => unreachable!(),
+        };
+        no_validation_error(
+            &format!(
+                "
+                {stage_attr}
+                fn main() {{
+                    subgroupBallot();
+                }}
+            "
+            ),
+            Capabilities::SUBGROUP,
+        );
+    }
+
+    // ...in vertex shaders require both [`Capabilities::SUBGROUP`] and
+    // [`Capabilities::SUBGROUP_VERTEX_STAGE`]`. (But note that
+    // `create_validator` automatically sets `Capabilities::SUBGROUP` whenever
+    // `Features::SUBGROUP_VERTEX` is available.)
+    for cap in [Capabilities::SUBGROUP, Capabilities::SUBGROUP_VERTEX_STAGE] {
+        check_validation! {
+            "
+                @vertex
+                fn main() -> @builtin(position) vec4<f32> {{
+                    subgroupBallot();
+                    return vec4();
+                }}
+            ":
+            Err(_),
+            cap
+        }
+    }
+    no_validation_error(
+        "
+            @vertex
+            fn main() -> @builtin(position) vec4<f32> {{
+                subgroupBallot();
+                return vec4();
+            }}
+        ",
+        Capabilities::SUBGROUP | Capabilities::SUBGROUP_VERTEX_STAGE,
+    );
+
+    // Subgroup barriers...
+
+    // ...require both SUBGROUP and SUBGROUP_BARRIER.
+    for cap in [Capabilities::SUBGROUP, Capabilities::SUBGROUP_BARRIER] {
+        check_validation! {
+            r#"
+                @compute @workgroup_size(1)
+                fn main() {
+                    subgroupBarrier();
+                }
+            "#:
+            Err(naga::valid::ValidationError::EntryPoint {
+                stage: naga::ShaderStage::Compute,
+                source: naga::valid::EntryPointError::Function(
+                    naga::valid::FunctionError::MissingCapability(required_caps)
+                ),
+                ..
+            }) if *required_caps == Capabilities::SUBGROUP | Capabilities::SUBGROUP_BARRIER,
+            cap
+        }
+    }
+
+    // ...are never supported in vertex shaders.
+    check_validation! {
+        r#"
+            @vertex
+            fn main() -> @builtin(position) vec4<f32> {
+                subgroupBarrier();
+                return vec4();
+            }
+        "#:
+        Err(naga::valid::ValidationError::EntryPoint {
+            stage: naga::ShaderStage::Vertex,
+            source: naga::valid::EntryPointError::ForbiddenStageOperations,
+            ..
+        }),
+        Capabilities::SUBGROUP | Capabilities::SUBGROUP_BARRIER | Capabilities::SUBGROUP_VERTEX_STAGE
+    }
+
+    // ...are never supported in fragment shaders.
+    check_validation! {
+        r#"
+            @fragment
+            fn main() {
+                subgroupBarrier();
+            }
+        "#:
+        Err(naga::valid::ValidationError::EntryPoint {
+            stage: naga::ShaderStage::Fragment,
+            source: naga::valid::EntryPointError::ForbiddenStageOperations,
+            ..
+        }),
+        Capabilities::SUBGROUP | Capabilities::SUBGROUP_BARRIER
+    }
+
+    // The `subgroup_id` built-in...
+
+    // ...in compute shaders requires [`Capabilities::SUBGROUP`]`.
+    check_one_validation! {
+        "
+            @compute @workgroup_size(1)
+            fn main(@builtin(subgroup_id) subgroup_id: u32) {{
+            }}
+        ",
+        Err(naga::valid::ValidationError::EntryPoint {
+            stage: naga::ShaderStage::Compute,
+            source: naga::valid::EntryPointError::Argument(
+                _,
+                naga::valid::VaryingError::UnsupportedCapability(Capabilities::SUBGROUP)
+            ),
+            ..
+        })
+    }
+
+    // ...in compute shaders requires *only* [`Capabilities::SUBGROUP`]`.
+    no_validation_error(
+        "
+        @compute @workgroup_size(1)
+        fn main(@builtin(subgroup_id) subgroup_id: u32) {{
+        }}
+        ",
+        Capabilities::SUBGROUP,
+    );
+}
+
 #[test]
 fn subgroup_invalid_broadcast() {
     check_validation! {
diff --git a/wgpu-types/src/features.rs b/wgpu-types/src/features.rs
index 94a9a91cb..397de544d 100644
--- a/wgpu-types/src/features.rs
+++ b/wgpu-types/src/features.rs
@@ -1058,7 +1058,8 @@ bitflags_array! {
         ///
         /// This is a native only feature.
         const SHADER_INT64 = 1 << 37;
-        /// Allows compute and fragment shaders to use the subgroup operation built-ins
+        /// Allows compute and fragment shaders to use the subgroup operation
+        /// built-ins and perform subgroup operations (except barriers).
         ///
         /// Supported Platforms:
         /// - Vulkan
@@ -1067,14 +1068,17 @@ bitflags_array! {
         ///
         /// This is a native only feature.
         const SUBGROUP = 1 << 38;
-        /// Allows vertex shaders to use the subgroup operation built-ins
+        /// Allows vertex shaders to use the subgroup operation built-ins and
+        /// perform subgroup operations (except barriers).
         ///
         /// Supported Platforms:
         /// - Vulkan
         ///
         /// This is a native only feature.
         const SUBGROUP_VERTEX = 1 << 39;
-        /// Allows shaders to use the subgroup barrier
+        /// Allows compute shaders to use the subgroup barrier.
+        ///
+        /// Requires [`Features::SUBGROUP`]. Without it, enables nothing.
         ///
         /// Supported Platforms:
         /// - Vulkan