feat(wgsl-in): parse diagnostic attrs. on fns

CHANGELOG.md

@@ -85,12 +85,12 @@ fn main(@builtin(position) p : vec4f) -> @location(0) vec4f {
 
 There are some limitations to keep in mind with this new functionality:
 
-- We do not yet support `diagnostic(…)` rules in attribute position (i.e., `@diagnostic(…) fn my_func { … }`). This is being tracked in <https://github.com/gfx-rs/wgpu/issues/5320>. We expect that rules in `fn` attribute position will be relaxed shortly (see <https://github.com/gfx-rs/wgpu/pull/6353>), but the prioritization for statement positions is unclear. If you are blocked by not being able to parse `diagnostic(…)` rules in statement positions, please let us know in that issue, so we can determine how to prioritize it!
+- We support `@diagnostic(…)` rules as `fn` attributes, but prioritization for rules in statement positions (i.e., `if (…) @diagnostic(…) { … }` is unclear. If you are blocked by not being able to parse `diagnostic(…)` rules in statement positions, please let us know in <https://github.com/gfx-rs/wgpu/issues/5320>, so we can determine how to prioritize it!
 - Standard WGSL specifies `error`, `warning`, `info`, and `off` severity levels. These are all technically usable now! A caveat, though: warning- and info-level are only emitted to `stderr` via the `log` façade, rather than being reported through a `Result::Err` in Naga or the `CompilationInfo` interface in `wgpu{,-core}`. This will require breaking changes in Naga to fix, and is being tracked by <https://github.com/gfx-rs/wgpu/issues/6458>.
 - Not all lints can be controlled with `diagnostic(…)` rules. In fact, only the `derivative_uniformity` triggering rule exists in the WGSL standard. That said, Naga contributors are excited to see how this level of control unlocks a new ecosystem of configurable diagnostics.
 - Finally, `diagnostic(…)` rules are not yet emitted in WGSL output. This means that `wgsl-in` → `wgsl-out` is currently a lossy process. We felt that it was important to unblock users who needed `diagnostic(…)` rules (i.e., <https://github.com/gfx-rs/wgpu/issues/3135>) before we took significant effort to fix this (tracked in <https://github.com/gfx-rs/wgpu/issues/6496>).
 
-By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148](https://github.com/gfx-rs/wgpu/pull/6148), [#6533](https://github.com/gfx-rs/wgpu/pull/6533).
+By @ErichDonGubler in [#6456](https://github.com/gfx-rs/wgpu/pull/6456), [#6148](https://github.com/gfx-rs/wgpu/pull/6148), [#6533](https://github.com/gfx-rs/wgpu/pull/6533), [#6353](https://github.com/gfx-rs/wgpu/pull/6353).
 
 ### New Features
 

naga/src/front/glsl/functions.rs

@@ -1066,6 +1066,7 @@ impl Frontend {
             expressions,
             named_expressions: crate::NamedExpressions::default(),
             body,
+            diagnostic_filter_leaf: None,
         };
 
         'outer: for decl in declaration.overloads.iter_mut() {

naga/src/front/spv/function.rs

@@ -60,6 +60,7 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
                 ),
                 named_expressions: crate::NamedExpressions::default(),
                 body: crate::Block::new(),
+                diagnostic_filter_leaf: None,
             }
         };
 
@@ -311,6 +312,7 @@ impl<I: Iterator<Item = u32>> super::Frontend<I> {
             expressions: Arena::new(),
             named_expressions: crate::NamedExpressions::default(),
             body: crate::Block::new(),
+            diagnostic_filter_leaf: None,
         };
 
         // 1. copy the inputs from arguments to privates

naga/src/front/wgsl/lower/mod.rs

@@ -1284,6 +1284,7 @@ impl<'source, 'temp> Lowerer<'source, 'temp> {
             expressions,
             named_expressions: crate::NamedExpressions::default(),
             body: crate::Block::default(),
+            diagnostic_filter_leaf: f.diagnostic_filter_leaf,
         };
 
         let mut typifier = Typifier::default();

naga/src/front/wgsl/parse/ast.rs

@@ -133,6 +133,7 @@ pub struct Function<'a> {
     pub arguments: Vec<FunctionArgument<'a>>,
     pub result: Option<FunctionResult<'a>>,
     pub body: Block<'a>,
+    pub diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
 }
 
 #[derive(Debug)]

naga/src/front/wgsl/parse/mod.rs

@@ -2219,6 +2219,7 @@ impl Parser {
     fn function_decl<'a>(
         &mut self,
         lexer: &mut Lexer<'a>,
+        diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
         out: &mut ast::TranslationUnit<'a>,
         dependencies: &mut FastIndexSet<ast::Dependency<'a>>,
     ) -> Result<ast::Function<'a>, Error<'a>> {
@@ -2291,6 +2292,7 @@ impl Parser {
             arguments,
             result,
             body,
+            diagnostic_filter_leaf,
         };
 
         // done
@@ -2526,13 +2528,13 @@ impl Parser {
                 Some(ast::GlobalDeclKind::Var(var))
             }
             (Token::Word("fn"), _) => {
-                if !diagnostic_filters.is_empty() {
+                let diagnostic_filter_leaf = Self::write_diagnostic_filters(
-                    return Err(Error::DiagnosticAttributeNotYetImplementedAtParseSite {
+                    &mut out.diagnostic_filters,
-                        site_name_plural: "functions",
+                    diagnostic_filters,
-                        spans: diagnostic_filters.spans().collect(),
+                    out.diagnostic_filter_leaf,
-                    });
+                );
-                }
+                let function =
-                let function = self.function_decl(lexer, out, &mut dependencies)?;
+                    self.function_decl(lexer, diagnostic_filter_leaf, out, &mut dependencies)?;
                 Some(ast::GlobalDeclKind::Fn(ast::Function {
                     entry_point: if let Some(stage) = stage.value {
                         if stage == ShaderStage::Compute && workgroup_size.value.is_none() {

naga/src/lib.rs

@@ -2121,6 +2121,14 @@ pub struct Function {
     pub named_expressions: NamedExpressions,
     /// Block of instructions comprising the body of the function.
     pub body: Block,
+    /// The leaf of all diagnostic filter rules tree (stored in [`Module::diagnostic_filters`])
+    /// parsed on this function.
+    ///
+    /// In WGSL, this corresponds to `@diagnostic(…)` attributes.
+    ///
+    /// See [`DiagnosticFilterNode`] for details on how the tree is represented and used in
+    /// validation.
+    pub diagnostic_filter_leaf: Option<Handle<DiagnosticFilterNode>>,
 }
 
 /// The main function for a pipeline stage.

naga/src/valid/analyzer.rs

@@ -1150,7 +1150,7 @@ impl ModuleInfo {
             expressions: vec![ExpressionInfo::new(); fun.expressions.len()].into_boxed_slice(),
             sampling: crate::FastHashSet::default(),
             dual_source_blending: false,
-            diagnostic_filter_leaf: module.diagnostic_filter_leaf,
+            diagnostic_filter_leaf: fun.diagnostic_filter_leaf,
         };
         let resolve_context =
             ResolveContext::with_locals(module, &fun.local_variables, &fun.arguments);

naga/src/valid/handles.rs

@@ -122,6 +122,7 @@ impl super::Validator {
                 ref expressions,
                 ref named_expressions,
                 ref body,
+                ref diagnostic_filter_leaf,
             } = function;
 
             for arg in arguments.iter() {
@@ -165,6 +166,10 @@ impl super::Validator {
 
             Self::validate_block_handles(body, expressions, functions)?;
 
+            if let Some(handle) = *diagnostic_filter_leaf {
+                handle.check_valid_for(diagnostic_filters)?;
+            }
+
             Ok(())
         };
 

naga/tests/in/diagnostic-filter.wgsl

@@ -1 +1,6 @@
 diagnostic(off, derivative_uniformity);
+
+fn thing() {}
+
+@diagnostic(warning, derivative_uniformity)
+fn with_diagnostic() {}

naga/tests/out/ir/access.compact.ron

@@ -923,6 +923,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("test_matrix_within_array_within_struct_accesses"),
@@ -1527,6 +1528,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("read_from_private"),
@@ -1560,6 +1562,7 @@
                     value: Some(1),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("test_arr_as_arg"),
@@ -1602,6 +1605,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_through_ptr_fn"),
@@ -1630,6 +1634,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_array_through_ptr_fn"),
@@ -1690,6 +1695,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("fetch_arg_ptr_member"),
@@ -1727,6 +1733,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_to_arg_ptr_member"),
@@ -1763,6 +1770,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("fetch_arg_ptr_array_element"),
@@ -1800,6 +1808,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_to_arg_ptr_array_element"),
@@ -1836,6 +1845,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -2138,6 +2148,7 @@
                         value: Some(52),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
         (
@@ -2329,6 +2340,7 @@
                         value: Some(31),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
         (
@@ -2410,6 +2422,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
         (
@@ -2473,6 +2486,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/access.ron

@@ -923,6 +923,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("test_matrix_within_array_within_struct_accesses"),
@@ -1527,6 +1528,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("read_from_private"),
@@ -1560,6 +1562,7 @@
                     value: Some(1),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("test_arr_as_arg"),
@@ -1602,6 +1605,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_through_ptr_fn"),
@@ -1630,6 +1634,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_array_through_ptr_fn"),
@@ -1690,6 +1695,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("fetch_arg_ptr_member"),
@@ -1727,6 +1733,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_to_arg_ptr_member"),
@@ -1763,6 +1770,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("fetch_arg_ptr_array_element"),
@@ -1800,6 +1808,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("assign_to_arg_ptr_array_element"),
@@ -1836,6 +1845,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -2138,6 +2148,7 @@
                         value: Some(52),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
         (
@@ -2329,6 +2340,7 @@
                         value: Some(31),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
         (
@@ -2410,6 +2422,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
         (
@@ -2473,6 +2486,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/atomic_i_increment.compact.ron

@@ -254,6 +254,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -276,6 +277,7 @@
                         result: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/atomic_i_increment.ron

@@ -279,6 +279,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -301,6 +302,7 @@
                         result: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/collatz.compact.ron

@@ -248,6 +248,7 @@
                     value: Some(23),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -327,6 +328,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/collatz.ron

@@ -248,6 +248,7 @@
                     value: Some(23),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -327,6 +328,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/const_assert.compact.ron

@@ -48,6 +48,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [],

naga/tests/out/ir/const_assert.ron

@@ -48,6 +48,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [],

naga/tests/out/ir/diagnostic-filter.compact.ron

@@ -9,7 +9,36 @@
     overrides: [],
     global_variables: [],
     global_expressions: [],
-    functions: [],
+    functions: [
+        (
+            name: Some("thing"),
+            arguments: [],
+            result: None,
+            local_variables: [],
+            expressions: [],
+            named_expressions: {},
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: Some(0),
+        ),
+        (
+            name: Some("with_diagnostic"),
+            arguments: [],
+            result: None,
+            local_variables: [],
+            expressions: [],
+            named_expressions: {},
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: Some(1),
+        ),
+    ],
     entry_points: [],
     diagnostic_filters: [
         (
@@ -19,6 +48,13 @@
             ),
             parent: None,
         ),
+        (
+            inner: (
+                new_severity: Warning,
+                triggering_rule: DerivativeUniformity,
+            ),
+            parent: Some(0),
+        ),
     ],
     diagnostic_filter_leaf: Some(0),
 )

naga/tests/out/ir/diagnostic-filter.ron

@@ -9,7 +9,36 @@
     overrides: [],
     global_variables: [],
     global_expressions: [],
-    functions: [],
+    functions: [
+        (
+            name: Some("thing"),
+            arguments: [],
+            result: None,
+            local_variables: [],
+            expressions: [],
+            named_expressions: {},
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: Some(0),
+        ),
+        (
+            name: Some("with_diagnostic"),
+            arguments: [],
+            result: None,
+            local_variables: [],
+            expressions: [],
+            named_expressions: {},
+            body: [
+                Return(
+                    value: None,
+                ),
+            ],
+            diagnostic_filter_leaf: Some(1),
+        ),
+    ],
     entry_points: [],
     diagnostic_filters: [
         (
@@ -19,6 +48,13 @@
             ),
             parent: None,
         ),
+        (
+            inner: (
+                new_severity: Warning,
+                triggering_rule: DerivativeUniformity,
+            ),
+            parent: Some(0),
+        ),
     ],
     diagnostic_filter_leaf: Some(0),
 )

naga/tests/out/ir/fetch_depth.compact.ron

@@ -167,6 +167,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -189,6 +190,7 @@
                         result: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/fetch_depth.ron

@@ -237,6 +237,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -259,6 +260,7 @@
                         result: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/index-by-value.compact.ron

@@ -127,6 +127,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("index_let_array"),
@@ -214,6 +215,7 @@
                     value: Some(10),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("index_let_matrix"),
@@ -289,6 +291,7 @@
                     value: Some(10),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -366,6 +369,7 @@
                         value: Some(9),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/index-by-value.ron

@@ -127,6 +127,7 @@
                     value: Some(2),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("index_let_array"),
@@ -214,6 +215,7 @@
                     value: Some(10),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("index_let_matrix"),
@@ -289,6 +291,7 @@
                     value: Some(10),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -366,6 +369,7 @@
                         value: Some(9),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/local-const.compact.ron

@@ -133,6 +133,7 @@
                     end: 5,
                 )),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [],

naga/tests/out/ir/local-const.ron

@@ -133,6 +133,7 @@
                     end: 5,
                 )),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [],

naga/tests/out/ir/overrides-atomicCompareExchangeWeak.compact.ron

@@ -122,6 +122,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/overrides-atomicCompareExchangeWeak.ron

@@ -122,6 +122,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/overrides-ray-query.compact.ron

@@ -253,6 +253,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/overrides-ray-query.ron

@@ -253,6 +253,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/overrides.compact.ron

@@ -193,6 +193,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/overrides.ron

@@ -193,6 +193,7 @@
                         value: None,
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/shadow.compact.ron

@@ -540,6 +540,7 @@
                     value: Some(34),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("fs_main"),
@@ -948,6 +949,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -1023,6 +1025,7 @@
                         value: Some(5),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/shadow.ron

@@ -795,6 +795,7 @@
                     value: Some(70),
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
         (
             name: Some("fs_main"),
@@ -1226,6 +1227,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -1301,6 +1303,7 @@
                         value: Some(5),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/spec-constants.compact.ron

@@ -486,6 +486,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -606,6 +607,7 @@
                         value: Some(14),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],

naga/tests/out/ir/spec-constants.ron

@@ -592,6 +592,7 @@
                     value: None,
                 ),
             ],
+            diagnostic_filter_leaf: None,
         ),
     ],
     entry_points: [
@@ -712,6 +713,7 @@
                         value: Some(14),
                     ),
                 ],
+                diagnostic_filter_leaf: None,
             ),
         ),
     ],