feat(wgsl-in): parse diagnostic(…) attributes 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/diagnostic_filter.rs

@@ -86,6 +86,18 @@ pub struct DiagnosticFilter {
     pub triggering_rule: FilterableTriggeringRule,
 }
 
+/// Determines whether [`DiagnosticFilterMap::add`] should consider full duplicates a conflict.
+///
+/// In WGSL, directive position does not consider this case a conflict, while attribute position
+/// does.
+#[cfg(feature = "wgsl-in")]
+pub(crate) enum ShouldConflictOnFullDuplicate {
+    /// Use this for attributes in WGSL.
+    Yes,
+    /// Use this for directives in WGSL.
+    No,
+}
+
 /// A map of diagnostic filters to their severity and first occurrence's span.
 ///
 /// Intended for front ends' first step into storing parsed [`DiagnosticFilter`]s.
@@ -104,6 +116,7 @@ impl DiagnosticFilterMap {
         &mut self,
         diagnostic_filter: DiagnosticFilter,
         span: Span,
+        should_conflict_on_full_duplicate: ShouldConflictOnFullDuplicate,
     ) -> Result<(), ConflictingDiagnosticRuleError> {
         use indexmap::map::Entry;
 
@@ -119,7 +132,11 @@ impl DiagnosticFilterMap {
             }
             Entry::Occupied(entry) => {
                 let &(first_severity, first_span) = entry.get();
-                if first_severity != new_severity {
+                let should_conflict_on_full_duplicate = match should_conflict_on_full_duplicate {
+                    ShouldConflictOnFullDuplicate::Yes => true,
+                    ShouldConflictOnFullDuplicate::No => false,
+                };
+                if first_severity != new_severity || should_conflict_on_full_duplicate {
                     return Err(ConflictingDiagnosticRuleError {
                         triggering_rule,
                         triggering_rule_spans: [first_span, span],

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/error.rs

@@ -1047,13 +1047,14 @@ impl<'a> Error<'a> {
                         (first_span, "first rule".into()),
                         (second_span, "second rule".into()),
                     ],
-                    notes: vec![concat!(
-                        "multiple `diagnostic(…)` rules with the same rule name ",
-                        "conflict unless the severity is the same; ",
-                        "delete the rule you don't want, or ",
-                        "ensure that all severities with the same rule name match"
-                    )
-                    .into()],
+                    notes: vec![
+                        concat!(
+                            "Multiple `diagnostic(…)` rules with the same rule name ",
+                            "conflict unless it is a directive and the severity is the same.",
+                        )
+                        .into(),
+                        "You should delete the rule you don't want.".into(),
+                    ],
                 }
             }
             Error::DiagnosticAttributeNotYetImplementedAtParseSite {

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

@@ -1,5 +1,6 @@
 use crate::diagnostic_filter::{
     self, DiagnosticFilter, DiagnosticFilterMap, DiagnosticFilterNode, FilterableTriggeringRule,
+    ShouldConflictOnFullDuplicate,
 };
 use crate::front::wgsl::error::{Error, ExpectedToken};
 use crate::front::wgsl::parse::directive::enable_extension::{
@@ -2167,7 +2168,7 @@ impl Parser {
             if let Some(DirectiveKind::Diagnostic) = DirectiveKind::from_ident(name) {
                 if let Some(filter) = self.diagnostic_filter(lexer)? {
                     let span = self.peek_rule_span(lexer);
-                    diagnostic_filters.add(filter, span)?;
+                    diagnostic_filters.add(filter, span, ShouldConflictOnFullDuplicate::Yes)?;
                 }
             } else {
                 return Err(Error::Unexpected(
@@ -2218,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>> {
@@ -2290,6 +2292,7 @@ impl Parser {
             arguments,
             result,
             body,
+            diagnostic_filter_leaf,
         };
 
         // done
@@ -2369,7 +2372,7 @@ impl Parser {
             if let Some(DirectiveKind::Diagnostic) = DirectiveKind::from_ident(name) {
                 if let Some(filter) = self.diagnostic_filter(lexer)? {
                     let span = self.peek_rule_span(lexer);
-                    diagnostic_filters.add(filter, span)?;
+                    diagnostic_filters.add(filter, span, ShouldConflictOnFullDuplicate::Yes)?;
                 }
                 continue;
             }
@@ -2525,13 +2528,13 @@ impl Parser {
                 Some(ast::GlobalDeclKind::Var(var))
             }
             (Token::Word("fn"), _) => {
-                if !diagnostic_filters.is_empty() {
-                    return Err(Error::DiagnosticAttributeNotYetImplementedAtParseSite {
-                        site_name_plural: "functions",
-                        spans: diagnostic_filters.spans().collect(),
-                    });
-                }
-                let function = self.function_decl(lexer, out, &mut dependencies)?;
+                let diagnostic_filter_leaf = Self::write_diagnostic_filters(
+                    &mut out.diagnostic_filters,
+                    diagnostic_filters,
+                    out.diagnostic_filter_leaf,
+                );
+                let function =
+                    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() {
@@ -2602,7 +2605,11 @@ impl Parser {
                     DirectiveKind::Diagnostic => {
                         if let Some(diagnostic_filter) = self.diagnostic_filter(&mut lexer)? {
                             let span = self.peek_rule_span(&lexer);
-                            diagnostic_filters.add(diagnostic_filter, span)?;
+                            diagnostic_filters.add(
+                                diagnostic_filter,
+                                span,
+                                ShouldConflictOnFullDuplicate::No,
+                            )?;
                         }
                         lexer.expect(Token::Separator(';'))?;
                     }

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,
             ),
         ),
     ],