Skip to content

Commit d31d944

Browse files
Clarify the effect of the SUBGROUP features and capabilities (#8203)
1 parent 3758b08 commit d31d944

File tree

4 files changed

+212
-9
lines changed

4 files changed

+212
-9
lines changed

CHANGELOG.md

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -219,6 +219,12 @@ By @cwfitzgerald in [#8162](https://github.com/gfx-rs/wgpu/pull/8162).
219219

220220
- [wgsl-in] Allow a trailing comma in `@blend_src(…)` attributes. By @ErichDonGubler in [#8137](https://github.com/gfx-rs/wgpu/pull/8137).
221221

222+
### Documentation
223+
224+
#### General
225+
226+
- Clarify that subgroup barriers require both the `SUBGROUP` and `SUBGROUP_BARRIER` features / capabilities. By @andyleiserson in TBD.
227+
222228
## v26.0.4 (2025-08-07)
223229

224230
### Bug Fixes

naga/src/valid/mod.rs

Lines changed: 23 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -131,13 +131,26 @@ bitflags::bitflags! {
131131
const CUBE_ARRAY_TEXTURES = 1 << 15;
132132
/// Support for 64-bit signed and unsigned integers.
133133
const SHADER_INT64 = 1 << 16;
134-
/// Support for subgroup operations.
135-
/// Implies support for subgroup operations in both fragment and compute stages,
136-
/// but not necessarily in the vertex stage, which requires [`Capabilities::SUBGROUP_VERTEX_STAGE`].
134+
/// Support for subgroup operations (except barriers) in fragment and compute shaders.
135+
///
136+
/// Subgroup operations in the vertex stage require
137+
/// [`Capabilities::SUBGROUP_VERTEX_STAGE`] in addition to `Capabilities::SUBGROUP`.
138+
/// (But note that `create_validator` automatically sets
139+
/// `Capabilities::SUBGROUP` whenever `Features::SUBGROUP_VERTEX` is
140+
/// available.)
141+
///
142+
/// Subgroup barriers require [`Capabilities::SUBGROUP_BARRIER`] in addition to
143+
/// `Capabilities::SUBGROUP`.
137144
const SUBGROUP = 1 << 17;
138-
/// Support for subgroup barriers.
145+
/// Support for subgroup barriers in compute shaders.
146+
///
147+
/// Requires [`Capabilities::SUBGROUP`]. Without it, enables nothing.
139148
const SUBGROUP_BARRIER = 1 << 18;
140-
/// Support for subgroup operations in the vertex stage.
149+
/// Support for subgroup operations (not including barriers) in the vertex stage.
150+
///
151+
/// Without [`Capabilities::SUBGROUP`], enables nothing. (But note that
152+
/// `create_validator` automatically sets `Capabilities::SUBGROUP`
153+
/// whenever `Features::SUBGROUP_VERTEX` is available.)
141154
const SUBGROUP_VERTEX_STAGE = 1 << 19;
142155
/// Support for [`AtomicFunction::Min`] and [`AtomicFunction::Max`] on
143156
/// 64-bit integers in the [`Storage`] address space, when the return
@@ -206,7 +219,11 @@ bitflags::bitflags! {
206219
#[cfg_attr(feature = "deserialize", derive(serde::Deserialize))]
207220
#[derive(Clone, Copy, Debug, Default, Eq, PartialEq)]
208221
pub struct SubgroupOperationSet: u8 {
209-
/// Elect, Barrier
222+
/// Barriers
223+
// Possibly elections, when that is supported.
224+
// https://github.com/gfx-rs/wgpu/issues/6042#issuecomment-3272603431
225+
// Contrary to what the name "basic" suggests, HLSL/DX12 support the
226+
// other subgroup operations, but do not support subgroup barriers.
210227
const BASIC = 1 << 0;
211228
/// Any, All
212229
const VOTE = 1 << 1;

naga/tests/naga/wgsl_errors.rs

Lines changed: 176 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1166,6 +1166,16 @@ fn validation_error(
11661166
.map_err(|e| e.into_inner()) // TODO(https://github.com/gfx-rs/wgpu/issues/8153): Add tests for spans
11671167
}
11681168

1169+
/// Check that a shader validates successfully.
1170+
///
1171+
/// In a few tests it is useful to check conditions where a validation error
1172+
/// should be absent alongside conditions where it should be present. This
1173+
/// wrapper is less confusing than `validation_error().unwrap()`.
1174+
#[track_caller]
1175+
fn no_validation_error(source: &str, caps: naga::valid::Capabilities) {
1176+
validation_error(source, caps).unwrap();
1177+
}
1178+
11691179
#[test]
11701180
fn int64_capability() {
11711181
check_validation! {
@@ -3585,6 +3595,7 @@ fn issue7165() {
35853595
fn invalid_return_type(a: Struct) -> i32 { return a; }
35863596
";
35873597

3598+
// We need the span for the error, so have to invoke manually.
35883599
let module = naga::front::wgsl::parse_str(shader).unwrap();
35893600
let err = naga::valid::Validator::new(
35903601
naga::valid::ValidationFlags::all(),
@@ -3834,6 +3845,171 @@ fn const_eval_value_errors() {
38343845
assert!(variant("f32(abs(-9223372036854775807 - 1))").is_ok());
38353846
}
38363847

3848+
#[test]
3849+
fn subgroup_capability() {
3850+
// Some of these tests should be `check_extension_validation` tests that
3851+
// also check handling of the enable directive, but that handling is not
3852+
// currently correct. https://github.com/gfx-rs/wgpu/issues/8202
3853+
3854+
// Non-barrier subgroup operations...
3855+
3856+
// ...in fragment and compute shaders require [`Capabilities::SUBGROUP`]`.
3857+
for stage in [naga::ShaderStage::Fragment, naga::ShaderStage::Compute] {
3858+
let stage_attr = match stage {
3859+
naga::ShaderStage::Fragment => "@fragment",
3860+
naga::ShaderStage::Compute => "@compute @workgroup_size(1)",
3861+
_ => unreachable!(),
3862+
};
3863+
check_one_validation! {
3864+
&format!("
3865+
{stage_attr}
3866+
fn main() {{
3867+
subgroupBallot();
3868+
}}
3869+
"),
3870+
Err(naga::valid::ValidationError::EntryPoint {
3871+
stage: err_stage,
3872+
source: naga::valid::EntryPointError::Function(
3873+
naga::valid::FunctionError::MissingCapability(Capabilities::SUBGROUP)
3874+
),
3875+
..
3876+
}) if *err_stage == stage
3877+
}
3878+
}
3879+
3880+
// ...in fragment and compute shaders require *only* [`Capabilities::SUBGROUP`]`.
3881+
for stage in [naga::ShaderStage::Fragment, naga::ShaderStage::Compute] {
3882+
let stage_attr = match stage {
3883+
naga::ShaderStage::Fragment => "@fragment",
3884+
naga::ShaderStage::Compute => "@compute @workgroup_size(1)",
3885+
_ => unreachable!(),
3886+
};
3887+
no_validation_error(
3888+
&format!(
3889+
"
3890+
{stage_attr}
3891+
fn main() {{
3892+
subgroupBallot();
3893+
}}
3894+
"
3895+
),
3896+
Capabilities::SUBGROUP,
3897+
);
3898+
}
3899+
3900+
// ...in vertex shaders require both [`Capabilities::SUBGROUP`] and
3901+
// [`Capabilities::SUBGROUP_VERTEX_STAGE`]`. (But note that
3902+
// `create_validator` automatically sets `Capabilities::SUBGROUP` whenever
3903+
// `Features::SUBGROUP_VERTEX` is available.)
3904+
for cap in [Capabilities::SUBGROUP, Capabilities::SUBGROUP_VERTEX_STAGE] {
3905+
check_validation! {
3906+
"
3907+
@vertex
3908+
fn main() -> @builtin(position) vec4<f32> {{
3909+
subgroupBallot();
3910+
return vec4();
3911+
}}
3912+
":
3913+
Err(_),
3914+
cap
3915+
}
3916+
}
3917+
no_validation_error(
3918+
"
3919+
@vertex
3920+
fn main() -> @builtin(position) vec4<f32> {{
3921+
subgroupBallot();
3922+
return vec4();
3923+
}}
3924+
",
3925+
Capabilities::SUBGROUP | Capabilities::SUBGROUP_VERTEX_STAGE,
3926+
);
3927+
3928+
// Subgroup barriers...
3929+
3930+
// ...require both SUBGROUP and SUBGROUP_BARRIER.
3931+
for cap in [Capabilities::SUBGROUP, Capabilities::SUBGROUP_BARRIER] {
3932+
check_validation! {
3933+
r#"
3934+
@compute @workgroup_size(1)
3935+
fn main() {
3936+
subgroupBarrier();
3937+
}
3938+
"#:
3939+
Err(naga::valid::ValidationError::EntryPoint {
3940+
stage: naga::ShaderStage::Compute,
3941+
source: naga::valid::EntryPointError::Function(
3942+
naga::valid::FunctionError::MissingCapability(required_caps)
3943+
),
3944+
..
3945+
}) if *required_caps == Capabilities::SUBGROUP | Capabilities::SUBGROUP_BARRIER,
3946+
cap
3947+
}
3948+
}
3949+
3950+
// ...are never supported in vertex shaders.
3951+
check_validation! {
3952+
r#"
3953+
@vertex
3954+
fn main() -> @builtin(position) vec4<f32> {
3955+
subgroupBarrier();
3956+
return vec4();
3957+
}
3958+
"#:
3959+
Err(naga::valid::ValidationError::EntryPoint {
3960+
stage: naga::ShaderStage::Vertex,
3961+
source: naga::valid::EntryPointError::ForbiddenStageOperations,
3962+
..
3963+
}),
3964+
Capabilities::SUBGROUP | Capabilities::SUBGROUP_BARRIER | Capabilities::SUBGROUP_VERTEX_STAGE
3965+
}
3966+
3967+
// ...are never supported in fragment shaders.
3968+
check_validation! {
3969+
r#"
3970+
@fragment
3971+
fn main() {
3972+
subgroupBarrier();
3973+
}
3974+
"#:
3975+
Err(naga::valid::ValidationError::EntryPoint {
3976+
stage: naga::ShaderStage::Fragment,
3977+
source: naga::valid::EntryPointError::ForbiddenStageOperations,
3978+
..
3979+
}),
3980+
Capabilities::SUBGROUP | Capabilities::SUBGROUP_BARRIER
3981+
}
3982+
3983+
// The `subgroup_id` built-in...
3984+
3985+
// ...in compute shaders requires [`Capabilities::SUBGROUP`]`.
3986+
check_one_validation! {
3987+
"
3988+
@compute @workgroup_size(1)
3989+
fn main(@builtin(subgroup_id) subgroup_id: u32) {{
3990+
}}
3991+
",
3992+
Err(naga::valid::ValidationError::EntryPoint {
3993+
stage: naga::ShaderStage::Compute,
3994+
source: naga::valid::EntryPointError::Argument(
3995+
_,
3996+
naga::valid::VaryingError::UnsupportedCapability(Capabilities::SUBGROUP)
3997+
),
3998+
..
3999+
})
4000+
}
4001+
4002+
// ...in compute shaders requires *only* [`Capabilities::SUBGROUP`]`.
4003+
no_validation_error(
4004+
"
4005+
@compute @workgroup_size(1)
4006+
fn main(@builtin(subgroup_id) subgroup_id: u32) {{
4007+
}}
4008+
",
4009+
Capabilities::SUBGROUP,
4010+
);
4011+
}
4012+
38374013
#[test]
38384014
fn subgroup_invalid_broadcast() {
38394015
check_validation! {

wgpu-types/src/features.rs

Lines changed: 7 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1058,7 +1058,8 @@ bitflags_array! {
10581058
///
10591059
/// This is a native only feature.
10601060
const SHADER_INT64 = 1 << 37;
1061-
/// Allows compute and fragment shaders to use the subgroup operation built-ins
1061+
/// Allows compute and fragment shaders to use the subgroup operation
1062+
/// built-ins and perform subgroup operations (except barriers).
10621063
///
10631064
/// Supported Platforms:
10641065
/// - Vulkan
@@ -1067,14 +1068,17 @@ bitflags_array! {
10671068
///
10681069
/// This is a native only feature.
10691070
const SUBGROUP = 1 << 38;
1070-
/// Allows vertex shaders to use the subgroup operation built-ins
1071+
/// Allows vertex shaders to use the subgroup operation built-ins and
1072+
/// perform subgroup operations (except barriers).
10711073
///
10721074
/// Supported Platforms:
10731075
/// - Vulkan
10741076
///
10751077
/// This is a native only feature.
10761078
const SUBGROUP_VERTEX = 1 << 39;
1077-
/// Allows shaders to use the subgroup barrier
1079+
/// Allows compute shaders to use the subgroup barrier.
1080+
///
1081+
/// Requires [`Features::SUBGROUP`]. Without it, enables nothing.
10781082
///
10791083
/// Supported Platforms:
10801084
/// - Vulkan

0 commit comments

Comments
 (0)