From 09178f2400aa13cf9d5d94ce1529d8fd9e9d0bca Mon Sep 17 00:00:00 2001 From: Andy Leiserson Date: Tue, 15 Apr 2025 12:01:19 -0700 Subject: [PATCH] [naga] Track globals used directly by a named expression (#7540) --- CHANGELOG.md | 6 ++ naga/src/front/wgsl/lower/mod.rs | 3 + naga/src/valid/analyzer.rs | 22 ++++++- naga/tests/naga/validation.rs | 100 +++++++++++++++++++++++++++++++ 4 files changed, 128 insertions(+), 3 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index c376a118cd..90d0f16217 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -44,6 +44,12 @@ Bottom level categories: #### Naga +### Bux Fixes + +#### Naga + +Naga now infers the correct binding layout when a resource appears only in an assignment to `_`. By @andyleiserson in [#7540](https://github.com/gfx-rs/wgpu/pull/7540). + - Add polyfills for `dot4U8Packed` and `dot4I8Packed` for all backends. By @robamler in [#7494](https://github.com/gfx-rs/wgpu/pull/7494). ### Changes diff --git a/naga/src/front/wgsl/lower/mod.rs b/naga/src/front/wgsl/lower/mod.rs index 53f7d0dc84..03e4a6fc78 100644 --- a/naga/src/front/wgsl/lower/mod.rs +++ b/naga/src/front/wgsl/lower/mod.rs @@ -1988,6 +1988,9 @@ impl<'source, 'temp> Lowerer<'source, 'temp> { return Ok(()); } ast::StatementKind::Phony(expr) => { + // Remembered the RHS of the phony assignment as a named expression. This + // is important (1) to preserve the RHS for validation, (2) to track any + // referenced globals. let mut emitter = Emitter::default(); emitter.start(&ctx.function.expressions); diff --git a/naga/src/valid/analyzer.rs b/naga/src/valid/analyzer.rs index 335a5b8a4e..cf7fec2060 100644 --- a/naga/src/valid/analyzer.rs +++ b/naga/src/valid/analyzer.rs @@ -158,8 +158,11 @@ pub struct ExpressionInfo { /// originates. Otherwise, this is `None`. pub uniformity: Uniformity, - /// The number of statements and other expressions using this - /// expression's value. + /// The number of direct references to this expression in statements and + /// other expressions. + /// + /// This is a _local_ reference count only, it may be non-zero for + /// expressions that are ultimately unused. pub ref_count: usize, /// The global variable into which this expression produces a pointer. @@ -362,7 +365,7 @@ impl FunctionInfo { ) -> NonUniformResult { let info = &mut self.expressions[expr.index()]; info.ref_count += 1; - // mark the used global as read + // Record usage if this expression may access a global if let Some(global) = info.assignable_global { self.global_uses[global.index()] |= global_use; } @@ -1220,6 +1223,19 @@ impl ModuleInfo { info.uniformity = uniformity.result; info.may_kill = uniformity.exit.contains(ExitFlags::MAY_KILL); + // If there are any globals referenced directly by a named expression, + // ensure they are marked as used even if they are not referenced + // anywhere else. An important case where this matters is phony + // assignments used to include a global in the shader's resource + // interface. https://www.w3.org/TR/WGSL/#phony-assignment-section + for &handle in fun.named_expressions.keys() { + if let Some(global) = info[handle].assignable_global { + if info.global_uses[global.index()].is_empty() { + info.global_uses[global.index()] = GlobalUse::QUERY; + } + } + } + Ok(info) } diff --git a/naga/tests/naga/validation.rs b/naga/tests/naga/validation.rs index a72d43f2d1..4b813d384e 100644 --- a/naga/tests/naga/validation.rs +++ b/naga/tests/naga/validation.rs @@ -823,3 +823,103 @@ fn arity_check() { assert!(validate(Mf::Pow, &[1]).is_ok()); assert!(validate(Mf::Pow, &[3]).is_err()); } + +#[cfg(feature = "wgsl-in")] +#[test] +fn global_use_scalar() { + let source = " +@group(0) @binding(0) +var global: u32; + +@compute @workgroup_size(64) +fn main() { + let used = &global; +} + "; + + let module = naga::front::wgsl::parse_str(source).expect("module should parse"); + let info = valid::Validator::new(Default::default(), valid::Capabilities::all()) + .validate(&module) + .unwrap(); + + let global = module.global_variables.iter().next().unwrap().0; + assert_eq!( + info.get_entry_point(0)[global], + naga::valid::GlobalUse::QUERY + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn global_use_array() { + let source = " +@group(0) @binding(0) +var global: array; + +@compute @workgroup_size(64) +fn main() { + let used = &global; +} + "; + + let module = naga::front::wgsl::parse_str(source).expect("module should parse"); + let info = valid::Validator::new(Default::default(), valid::Capabilities::all()) + .validate(&module) + .unwrap(); + + let global = module.global_variables.iter().next().unwrap().0; + assert_eq!( + info.get_entry_point(0)[global], + naga::valid::GlobalUse::QUERY + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn global_use_array_index() { + let source = " +@group(0) @binding(0) +var global: array; + +@compute @workgroup_size(64) +fn main() { + let used = &global[0]; +} + "; + + let module = naga::front::wgsl::parse_str(source).expect("module should parse"); + let info = valid::Validator::new(Default::default(), valid::Capabilities::all()) + .validate(&module) + .unwrap(); + + let global = module.global_variables.iter().next().unwrap().0; + assert_eq!( + info.get_entry_point(0)[global], + naga::valid::GlobalUse::QUERY + ); +} + +#[cfg(feature = "wgsl-in")] +#[test] +fn global_use_phony() { + let source = " +@group(0) @binding(0) +var global: u32; + +@compute @workgroup_size(64) +fn main() { + _ = &global; +} + "; + + let module = naga::front::wgsl::parse_str(source).expect("module should parse"); + let info = valid::Validator::new(Default::default(), valid::Capabilities::all()) + .validate(&module) + .unwrap(); + + let global = module.global_variables.iter().next().unwrap().0; + assert_eq!( + info.get_entry_point(0)[global], + naga::valid::GlobalUse::QUERY + ); +}