archive-gh-me/wgpu
1
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-12-08 21:26:17 +00:00
Code Issues Projects Releases Wiki Activity

Clarify the effect of the SUBGROUP features and capabilities (#8203)

Browse Source
This commit is contained in:
Andy Leiserson 2025-09-10 17:53:56 -07:00 committed by GitHub
parent 3758b08be9
commit d31d944ed5
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
4 changed files with 212 additions and 9 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

6
CHANGELOG.md
View File

@ -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

29
naga/src/valid/mod.rs
View File

@ -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;

176
naga/tests/naga/wgsl_errors.rs
View File

@ -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! {

10
wgpu-types/src/features.rs
View File

@ -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