[naga] Allow unreachable statements (#7718)

Allow unreachable statements after return/break/continue/discard.

Fixes #7536

Author: andyleiserson

cts_runner/test.lst

@@ -17,3 +17,4 @@ webgpu:api,operation,rendering,color_target_state:blend_constant,setting:*
 webgpu:api,operation,rendering,depth:*
 webgpu:api,operation,rendering,draw:*
 webgpu:api,operation,uncapturederror:*
+webgpu:shader,execution,flow_control,return:*

naga/src/valid/function.rs

@@ -114,8 +114,6 @@ pub enum FunctionError {
         name: String,
         space: crate::AddressSpace,
     },
-    #[error("There are instructions after `return`/`break`/`continue`")]
-    InstructionsAfterReturn,
     #[error("The `break` is used outside of a `loop` or `switch` context")]
     BreakOutsideOfLoopOrSwitch,
     #[error("The `continue` is used outside of a `loop` context")]
@@ -236,7 +234,6 @@ bitflags::bitflags! {
 
 struct BlockInfo {
     stages: super::ShaderStages,
-    finished: bool,
 }
 
 struct BlockContext<'a> {
@@ -768,13 +765,8 @@ impl super::Validator {
         context: &BlockContext,
     ) -> Result<BlockInfo, WithSpan<FunctionError>> {
         use crate::{AddressSpace, Statement as S, TypeInner as Ti};
-        let mut finished = false;
         let mut stages = super::ShaderStages::all();
         for (statement, &span) in statements.span_iter() {
-            if finished {
-                return Err(FunctionError::InstructionsAfterReturn
-                    .with_span_static(span, "instructions after return"));
-            }
             match *statement {
                 S::Emit(ref range) => {
                     for handle in range.clone() {
@@ -823,7 +815,6 @@ impl super::Validator {
                 S::Block(ref block) => {
                     let info = self.validate_block(block, context)?;
                     stages &= info.stages;
-                    finished = info.finished;
                 }
                 S::If {
                     condition,
@@ -966,14 +957,12 @@ impl super::Validator {
                         return Err(FunctionError::BreakOutsideOfLoopOrSwitch
                             .with_span_static(span, "invalid break"));
                     }
-                    finished = true;
                 }
                 S::Continue => {
                     if !context.abilities.contains(ControlFlowAbility::CONTINUE) {
                         return Err(FunctionError::ContinueOutsideOfLoop
                             .with_span_static(span, "invalid continue"));
                     }
-                    finished = true;
                 }
                 S::Return { value } => {
                     if !context.abilities.contains(ControlFlowAbility::RETURN) {
@@ -1013,11 +1002,9 @@ impl super::Validator {
                             .with_span_static(span, "invalid return"));
                         }
                     }
-                    finished = true;
                 }
                 S::Kill => {
                     stages &= super::ShaderStages::FRAGMENT;
-                    finished = true;
                 }
                 S::Barrier(barrier) => {
                     stages &= super::ShaderStages::COMPUTE;
@@ -1635,7 +1622,7 @@ impl super::Validator {
                 }
             }
         }
-        Ok(BlockInfo { stages, finished })
+        Ok(BlockInfo { stages })
     }
 
     fn validate_block(

naga/tests/naga/validation.rs

@@ -923,3 +923,34 @@ fn main() {
         naga::valid::GlobalUse::QUERY
     );
 }
+
+#[cfg(feature = "wgsl-in")]
+#[test]
+fn global_use_unreachable() {
+    // We should allow statements after `return`, and such statements should
+    // still contribute to global usage. (Unreachable statements should not
+    // contribute to uniformity analysis, but there are multiple issues with
+    // the current implementation of uniformity analysis, see #4369.)
+
+    let source = "
+@group(0) @binding(0)
+var<storage, read_write> global: u32;
+
+@compute @workgroup_size(64)
+fn main() {
+    var used: u32;
+    return;
+    used = global;
+}
+    ";
+
+    let module = naga::front::wgsl::parse_str(source).expect("module should parse");
+    let mut validator = valid::Validator::new(Default::default(), valid::Capabilities::all());
+    let info = validator.validate(&module).unwrap();
+
+    let global = module.global_variables.iter().next().unwrap().0;
+    assert_eq!(
+        info.get_entry_point(0)[global],
+        naga::valid::GlobalUse::READ,
+    );
+}

naga/tests/naga/wgsl_errors.rs

@@ -1979,37 +1979,6 @@ fn invalid_local_vars() {
     }
 }
 
-#[test]
-fn dead_code() {
-    check_validation! {
-        "
-        fn dead_code_after_if(condition: bool) -> i32 {
-            if (condition) {
-                return 1;
-            } else {
-                return 2;
-            }
-            return 3;
-        }
-        ":
-        Ok(_)
-    }
-    check_validation! {
-        "
-        fn dead_code_after_block() -> i32 {
-            {
-                return 1;
-            }
-            return 2;
-        }
-        ":
-        Err(naga::valid::ValidationError::Function {
-            source: naga::valid::FunctionError::InstructionsAfterReturn,
-            ..
-        })
-    }
-}
-
 #[test]
 fn invalid_runtime_sized_arrays() {
     // You can't have structs whose last member is an unsized struct. An unsized