[naga] Track globals used directly by a named expression (gfx-rs#7540)

Author: andyleiserson

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

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

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)
     }
 

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<storage, read_write> 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<storage, read_write> global: array<f32>;
+
+@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<storage, read_write> global: array<f32>;
+
+@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<storage, read_write> 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
+    );
+}