[naga hlsl-out glsl-out] Work around backend loop/switch bugs.

Introduce a new module, `naga:🔙:continue_forward`, containing
shared code for rendering Naga `Continue` statements as backend
`break` statements and assignments to introduced `bool` locals.
See the module's documentation for details.

- [hlsl-out] Transform degenerate single body switches into `do-while`
  loops. Properly render `Continue` statements enclosed by
  `Switch` statements enclosed by `Loop` statements.

- [glsl-out] Transform degenerate single body switches into `do-while`
  loops.

Improve `naga xtask validate spv` error message.

Fixes #4485.
Fixes #4514.

CHANGELOG.md

@@ -296,6 +296,11 @@ This release fixes the validation errors whenever a surface is used with the vul
 
 -  Fix regression on OpenGL (EGL) where non-sRGB still used sRGB [#5642](https://github.com/gfx-rs/wgpu/pull/5642)
 
+#### Naga
+
+- Work around shader consumers that have bugs handling `switch` statements with a single body for all cases. These are now written as `do {} while(false);` loops in hlsl-out and glsl-out. By @Imberflur in [#5654](https://github.com/gfx-rs/wgpu/pull/5654)
+- In hlsl-out, defer `continue` statements in switches by setting a flag and breaking from the switch. This allows such constructs to work with FXC which does not support `continue` within a switch. By @Imberflur in [#5654](https://github.com/gfx-rs/wgpu/pull/5654)
+
 ## v0.20.0 (2024-04-28)
 
 ### Major Changes

naga/src/back/continue_forward.rs

@@ -0,0 +1,311 @@
+//! Workarounds for platform bugs and limitations in switches and loops.
+//!
+//! In these docs, we use CamelCase links for Naga IR concepts, and ordinary
+//! `code` formatting for HLSL or GLSL concepts.
+//!
+//! ## Avoiding `continue` within `switch`
+//!
+//! As described in <https://github.com/gfx-rs/wgpu/issues/4485>, the FXC HLSL
+//! compiler doesn't allow `continue` statements within `switch` statements, but
+//! Naga IR does. We work around this by introducing synthetic boolean local
+//! variables and branches.
+//!
+//! Specifically:
+//!
+//! - We generate code for [`Continue`] statements within [`SwitchCase`]s that
+//!   sets an introduced `bool` local to `true` and does a `break`, jumping to
+//!   immediately after the generated `switch`.
+//!
+//! - When generating code for a [`Switch`] statement, we conservatively assume
+//!   it might contain such a [`Continue`] statement, so:
+//!
+//!   - If it's the outermost such [`Switch`] within a [`Loop`], we declare the
+//!     `bool` local ahead of the switch, initialized to `false`. Immediately
+//!     after the `switch`, we check the local and do a `continue` if it's set.
+//!
+//!   - If the [`Switch`] is nested within other [`Switch`]es, then after the
+//!     generated `switch`, we check the local (which we know was declared
+//!     before the surrounding `switch`) and do a `break` if it's set.
+//!
+//!   - As an optimization, we only generate the check of the local if a
+//!     [`Continue`] statement is encountered within the [`Switch`]. This may
+//!     help drivers more easily identify that the `bool` is unused.
+//!
+//! So while we "weaken" the [`Continue`] statement by rendering it as a `break`
+//! statement, we also place checks immediately at the locations to which those
+//! `break` statements will jump, until we can be sure we've reached the
+//! intended target of the original [`Continue`].
+//!
+//! In the case of nested [`Loop`] and [`Switch`] statements, there may be
+//! multiple introduced `bool` locals in scope, but there's no problem knowing
+//! which one to operate on. At any point, there is at most one [`Loop`]
+//! statement that could be targeted by a [`Continue`] statement, so the correct
+//! `bool` local to set and test is always the one introduced for the innermost
+//! enclosing [`Loop`]'s outermost [`Switch`].
+//!
+//! # Avoiding single body `switch` statements
+//!
+//! As described in <https://github.com/gfx-rs/wgpu/issues/4514>, some language
+//! front ends miscompile `switch` statements where all cases branch to the same
+//! body. Our HLSL and GLSL backends render [`Switch`] statements with a single
+//! [`SwitchCase`] as `do {} while(false);` loops.
+//!
+//! However, this rewriting introduces a new loop that could "capture"
+//! `continue` statements in its body. To avoid doing so, we apply the
+//! [`Continue`]-to-`break` transformation described above.
+//!
+//! [`Continue`]: crate::Statement::Continue
+//! [`Loop`]: crate::Statement::Loop
+//! [`Switch`]: crate::Statement::Switch
+//! [`SwitchCase`]: crate::SwitchCase
+
+use crate::proc::Namer;
+use std::rc::Rc;
+
+/// A summary of the code surrounding a statement.
+enum Nesting {
+    /// Currently nested in at least one [`Loop`] statement.
+    ///
+    /// [`Continue`] should apply to the innermost loop.
+    ///
+    /// When this entry is on the top of the stack:
+    ///
+    /// * When entering an inner [`Loop`] statement, push a [`Loop`][nl] state
+    ///   onto the stack.
+    ///
+    /// * When entering a nested [`Switch`] statement, push a [`Switch`][ns]
+    ///   state onto the stack with a new variable name. Before the generated
+    ///   `switch`, introduce a `bool` local with that name, initialized to
+    ///   `false`.
+    ///
+    /// When exiting the [`Loop`] for which this entry was pushed, pop it from
+    /// the stack.
+    ///
+    /// [`Continue`]: crate::Statement::Continue
+    /// [`Loop`]: crate::Statement::Loop
+    /// [`Switch`]: crate::Statement::Switch
+    /// [ns]: Nesting::Switch
+    /// [nl]: Nesting::Loop
+    Loop,
+
+    /// Currently nested in at least one [`Switch`] that may need to forward
+    /// [`Continue`]s.
+    ///
+    /// This includes [`Switch`]es rendered as `do {} while(false)` loops, but
+    /// doesn't need to include regular [`Switch`]es in backends that can
+    /// support `continue` within switches.
+    ///
+    /// [`Continue`] should be forwarded to the innermost surrounding [`Loop`].
+    ///
+    /// When this entry is on the top of the stack:
+    ///
+    /// * When entering a nested [`Loop`], push a [`Loop`][nl] state onto the
+    ///   stack.
+    ///
+    /// * When entering a nested [`Switch`], push a [`Switch`][ns] state onto
+    ///   the stack with a clone of the introduced `bool` variable's name.
+    ///
+    /// * When encountering a [`Continue`] statement, render it as code to set
+    ///   the introduced `bool` local (whose name is held in [`variable`]) to
+    ///   `true`, and then `break`. Set [`continue_encountered`] to `true` to
+    ///   record that the [`Switch`] contains a [`Continue`].
+    ///
+    /// * When exiting this [`Switch`], pop its entry from the stack. If
+    ///   [`continue_encountered`] is set, then we have rendered [`Continue`]
+    ///   statements as `break` statements that jump to this point. Generate
+    ///   code to check `variable`, and if it is `true`:
+    ///
+    ///     * If there is another [`Switch`][ns] left on top of the stack, set
+    ///       its `continue_encountered` flag, and generate a `break`. (Both
+    ///       [`Switch`][ns]es are within the same [`Loop`] and share the same
+    ///       introduced variable, so there's no need to set another flag to
+    ///       continue to exit the `switch`es.)
+    ///
+    ///     * Otherwise, `continue`.
+    ///
+    /// When we exit the [`Switch`] for which this entry was pushed, pop it.
+    ///
+    /// [`Continue`]: crate::Statement::Continue
+    /// [`Loop`]: crate::Statement::Loop
+    /// [`Switch`]: crate::Statement::Switch
+    /// [`variable`]: Nesting::Switch::variable
+    /// [`continue_encountered`]: Nesting::Switch::continue_encountered
+    /// [ns]: Nesting::Switch
+    /// [nl]: Nesting::Loop
+    Switch {
+        variable: Rc<String>,
+
+        /// Set if we've generated code for a [`Continue`] statement with this
+        /// entry on the top of the stack.
+        ///
+        /// If this is still clear when we finish rendering the [`Switch`], then
+        /// we know we don't need to generate branch forwarding code. Omitting
+        /// that may make it easier for drivers to tell that the `bool` we
+        /// introduced ahead of the [`Switch`] is actually unused.
+        ///
+        /// [`Continue`]: crate::Statement::Continue
+        /// [`Switch`]: crate::Statement::Switch
+        continue_encountered: bool,
+    },
+}
+
+/// A micro-IR for code a backend should generate after a [`Switch`].
+///
+/// [`Switch`]: crate::Statement::Switch
+pub(super) enum ExitControlFlow {
+    None,
+    /// Emit `if (continue_variable) { continue; }`
+    Continue {
+        variable: Rc<String>,
+    },
+    /// Emit `if (continue_variable) { break; }`
+    ///
+    /// Used after a [`Switch`] to exit from an enclosing [`Switch`].
+    ///
+    /// After the enclosing switch, its associated check will consult this same
+    /// variable, see that it is set, and exit early.
+    ///
+    /// [`Switch`]: crate::Statement::Switch
+    Break {
+        variable: Rc<String>,
+    },
+}
+
+/// Utility for tracking nesting of loops and switches to orchestrate forwarding
+/// of continue statements inside of a switch to the enclosing loop.
+///
+/// See [module docs](self) for why we need this.
+#[derive(Default)]
+pub(super) struct ContinueCtx {
+    stack: Vec<Nesting>,
+}
+
+impl ContinueCtx {
+    /// Resets internal state.
+    ///
+    /// Use this to reuse memory between writing sessions.
+    pub fn clear(&mut self) {
+        self.stack.clear();
+    }
+
+    /// Updates internal state to record entering a [`Loop`] statement.
+    ///
+    /// [`Loop`]: crate::Statement::Loop
+    pub fn enter_loop(&mut self) {
+        self.stack.push(Nesting::Loop);
+    }
+
+    /// Updates internal state to record exiting a [`Loop`] statement.
+    ///
+    /// [`Loop`]: crate::Statement::Loop
+    pub fn exit_loop(&mut self) {
+        if !matches!(self.stack.pop(), Some(Nesting::Loop)) {
+            unreachable!("ContinueCtx stack out of sync");
+        }
+    }
+
+    /// Updates internal state to record entering a [`Switch`] statement.
+    ///
+    /// Return `Some(variable)` if this [`Switch`] is nested within a [`Loop`],
+    /// and the caller should introcue a new `bool` local variable named
+    /// `variable` above the `switch`, for forwarding [`Continue`] statements.
+    ///
+    /// `variable` is guaranteed not to conflict with any names used by the
+    /// program itself.
+    ///
+    /// [`Continue`]: crate::Statement::Continue
+    /// [`Loop`]: crate::Statement::Loop
+    /// [`Switch`]: crate::Statement::Switch
+    pub fn enter_switch(&mut self, namer: &mut Namer) -> Option<Rc<String>> {
+        match self.stack.last() {
+            // If the stack is empty, we are not in loop, so we don't need to
+            // forward continue statements within this `Switch`. We can leave
+            // the stack empty.
+            None => None,
+            Some(&Nesting::Loop { .. }) => {
+                let variable = Rc::new(namer.call("should_continue"));
+                self.stack.push(Nesting::Switch {
+                    variable: Rc::clone(&variable),
+                    continue_encountered: false,
+                });
+                Some(variable)
+            }
+            Some(&Nesting::Switch { ref variable, .. }) => {
+                self.stack.push(Nesting::Switch {
+                    variable: Rc::clone(variable),
+                    continue_encountered: false,
+                });
+                // We have already declared the variable before some enclosing
+                // `Switch`.
+                None
+            }
+        }
+    }
+
+    /// Update internal state to record leaving a [`Switch`] statement.
+    ///
+    /// Return an [`ExitControlFlow`] value indicating what code should be
+    /// introduced after the generated `switch` to forward continues.
+    ///
+    /// [`Switch`]: crate::Statement::Switch
+    pub fn exit_switch(&mut self) -> ExitControlFlow {
+        match self.stack.pop() {
+            // This doesn't indicate a problem: we don't start pushing entries
+            // for `Switch` statements unless we have an enclosing `Loop`.
+            None => ExitControlFlow::None,
+            Some(Nesting::Loop { .. }) => {
+                unreachable!("Unexpected loop state when exiting switch");
+            }
+            Some(Nesting::Switch {
+                variable,
+                continue_encountered: inner_continue,
+            }) => {
+                if !inner_continue {
+                    // No `Continue` statement was encountered, so we didn't
+                    // introduce any `break`s jumping to this point.
+                    ExitControlFlow::None
+                } else if let Some(&mut Nesting::Switch {
+                    continue_encountered: ref mut outer_continue,
+                    ..
+                }) = self.stack.last_mut()
+                {
+                    // This is nested in another `Switch`. Propagate upwards
+                    // that there is a continue statement present.
+                    *outer_continue = true;
+                    ExitControlFlow::Break { variable }
+                } else {
+                    ExitControlFlow::Continue { variable }
+                }
+            }
+        }
+    }
+
+    /// Determine what to generate for a [`Continue`] statement.
+    ///
+    /// If we can generate an ordinary `continue` statement, return `None`.
+    ///
+    /// Otherwise, we're enclosed by a [`Switch`] that is itself enclosed by a
+    /// [`Loop`]. Return `Some(variable)` to indicate that the [`Continue`]
+    /// should be rendered as setting `variable` to `true`, and then doing a
+    /// `break`.
+    ///
+    /// This also notes that we've encountered a [`Continue`] statement, so that
+    /// we can generate the right code to forward the branch following the
+    /// enclosing `switch`.
+    ///
+    /// [`Continue`]: crate::Statement::Continue
+    /// [`Loop`]: crate::Statement::Loop
+    /// [`Switch`]: crate::Statement::Switch
+    pub fn continue_encountered(&mut self) -> Option<&str> {
+        if let Some(&mut Nesting::Switch {
+            ref variable,
+            ref mut continue_encountered,
+        }) = self.stack.last_mut()
+        {
+            *continue_encountered = true;
+            Some(variable)
+        } else {
+            None
+        }
+    }
+}

naga/src/back/glsl/mod.rs

@@ -545,6 +545,11 @@ pub struct Writer<'a, W> {
     named_expressions: crate::NamedExpressions,
     /// Set of expressions that need to be baked to avoid unnecessary repetition in output
     need_bake_expressions: back::NeedBakeExpressions,
+    /// Information about nesting of loops and switches.
+    ///
+    /// Used for forwarding continue statements in switches that have been
+    /// transformed to `do {} while(false);` loops.
+    continue_ctx: back::continue_forward::ContinueCtx,
     /// How many views to render to, if doing multiview rendering.
     multiview: Option<std::num::NonZeroU32>,
     /// Mapping of varying variables to their location. Needed for reflections.
@@ -619,6 +624,7 @@ impl<'a, W: Write> Writer<'a, W> {
             block_id: IdGenerator::default(),
             named_expressions: Default::default(),
             need_bake_expressions: Default::default(),
+            continue_ctx: back::continue_forward::ContinueCtx::default(),
             varying: Default::default(),
         };
 
@@ -2082,42 +2088,94 @@ impl<'a, W: Write> Writer<'a, W> {
                 selector,
                 ref cases,
             } => {
-                // Start the switch
-                write!(self.out, "{level}")?;
-                write!(self.out, "switch(")?;
-                self.write_expr(selector, ctx)?;
-                writeln!(self.out, ") {{")?;
-
-                // Write all cases
                 let l2 = level.next();
-                for case in cases {
+                // Some GLSL consumers may not handle switches with a single
-                    match case.value {
+                // body correctly: See wgpu#4514. Write such switch statements
-                        crate::SwitchValue::I32(value) => write!(self.out, "{l2}case {value}:")?,
+                // as a `do {} while(false);` loop instead.
-                        crate::SwitchValue::U32(value) => write!(self.out, "{l2}case {value}u:")?,
+                //
-                        crate::SwitchValue::Default => write!(self.out, "{l2}default:")?,
+                // Since doing so may inadvertently capture `continue`
+                // statements in the switch body, we must apply continue
+                // forwarding. See the `naga::back::continue_forward` module
+                // docs for details.
+                let one_body = cases
+                    .iter()
+                    .rev()
+                    .skip(1)
+                    .all(|case| case.fall_through && case.body.is_empty());
+                if one_body {
+                    // Unlike HLSL, in GLSL `continue_ctx` only needs to know
+                    // about [`Switch`] statements that are being rendered as
+                    // `do-while` loops.
+                    if let Some(variable) = self.continue_ctx.enter_switch(&mut self.namer) {
+                        writeln!(self.out, "{level}bool {variable} = false;",)?;
+                    };
+                    writeln!(self.out, "{level}do {{")?;
+                    // Note: Expressions have no side-effects so we don't need to emit selector expression.
+
+                    // Body
+                    if let Some(case) = cases.last() {
+                        for sta in case.body.iter() {
+                            self.write_stmt(sta, ctx, l2)?;
+                        }
+                    }
+                    // End do-while
+                    writeln!(self.out, "{level}}} while(false);")?;
+
+                    // Handle any forwarded continue statements.
+                    use back::continue_forward::ExitControlFlow;
+                    let op = match self.continue_ctx.exit_switch() {
+                        ExitControlFlow::None => None,
+                        ExitControlFlow::Continue { variable } => Some(("continue", variable)),
+                        ExitControlFlow::Break { variable } => Some(("break", variable)),
+                    };
+                    if let Some((control_flow, variable)) = op {
+                        writeln!(self.out, "{level}if ({variable}) {{")?;
+                        writeln!(self.out, "{l2}{control_flow};")?;
+                        writeln!(self.out, "{level}}}")?;
+                    }
+                } else {
+                    // Start the switch
+                    write!(self.out, "{level}")?;
+                    write!(self.out, "switch(")?;
+                    self.write_expr(selector, ctx)?;
+                    writeln!(self.out, ") {{")?;
+
+                    // Write all cases
+                    for case in cases {
+                        match case.value {
+                            crate::SwitchValue::I32(value) => {
+                                write!(self.out, "{l2}case {value}:")?
+                            }
+                            crate::SwitchValue::U32(value) => {
+                                write!(self.out, "{l2}case {value}u:")?
+                            }
+                            crate::SwitchValue::Default => write!(self.out, "{l2}default:")?,
+                        }
+
+                        let write_block_braces = !(case.fall_through && case.body.is_empty());
+                        if write_block_braces {
+                            writeln!(self.out, " {{")?;
+                        } else {
+                            writeln!(self.out)?;
+                        }
+
+                        for sta in case.body.iter() {
+                            self.write_stmt(sta, ctx, l2.next())?;
+                        }
+
+                        if !case.fall_through
+                            && case.body.last().map_or(true, |s| !s.is_terminator())
+                        {
+                            writeln!(self.out, "{}break;", l2.next())?;
+                        }
+
+                        if write_block_braces {
+                            writeln!(self.out, "{l2}}}")?;
+                        }
                     }
 
-                    let write_block_braces = !(case.fall_through && case.body.is_empty());
+                    writeln!(self.out, "{level}}}")?
-                    if write_block_braces {
-                        writeln!(self.out, " {{")?;
-                    } else {
-                        writeln!(self.out)?;
-                    }
-
-                    for sta in case.body.iter() {
-                        self.write_stmt(sta, ctx, l2.next())?;
-                    }
-
-                    if !case.fall_through && case.body.last().map_or(true, |s| !s.is_terminator()) {
-                        writeln!(self.out, "{}break;", l2.next())?;
-                    }
-
-                    if write_block_braces {
-                        writeln!(self.out, "{l2}}}")?;
-                    }
                 }
-
-                writeln!(self.out, "{level}}}")?
             }
             // Loops in naga IR are based on wgsl loops, glsl can emulate the behaviour by using a
             // while true loop and appending the continuing block to the body resulting on:
@@ -2134,6 +2192,7 @@ impl<'a, W: Write> Writer<'a, W> {
                 ref continuing,
                 break_if,
             } => {
+                self.continue_ctx.enter_loop();
                 if !continuing.is_empty() || break_if.is_some() {
                     let gate_name = self.namer.call("loop_init");
                     writeln!(self.out, "{level}bool {gate_name} = true;")?;
@@ -2159,7 +2218,8 @@ impl<'a, W: Write> Writer<'a, W> {
                 for sta in body {
                     self.write_stmt(sta, ctx, level.next())?;
                 }
-                writeln!(self.out, "{level}}}")?
+                writeln!(self.out, "{level}}}")?;
+                self.continue_ctx.exit_loop();
             }
             // Break, continue and return as written as in C
             // `break;`
@@ -2169,8 +2229,14 @@ impl<'a, W: Write> Writer<'a, W> {
             }
             // `continue;`
             Statement::Continue => {
-                write!(self.out, "{level}")?;
+                // Sometimes we must render a `Continue` statement as a `break`.
-                writeln!(self.out, "continue;")?
+                // See the docs for the `back::continue_forward` module.
+                if let Some(variable) = self.continue_ctx.continue_encountered() {
+                    writeln!(self.out, "{level}{variable} = true;",)?;
+                    writeln!(self.out, "{level}break;")?
+                } else {
+                    writeln!(self.out, "{level}continue;")?
+                }
             }
             // `return expr;`, `expr` is optional
             Statement::Return { value } => {

naga/src/back/hlsl/mod.rs

@@ -327,6 +327,7 @@ pub struct Writer<'a, W> {
     /// Set of expressions that have associated temporary variables
     named_expressions: crate::NamedExpressions,
     wrapped: Wrapped,
+    continue_ctx: back::continue_forward::ContinueCtx,
 
     /// A reference to some part of a global variable, lowered to a series of
     /// byte offset calculations.

naga/src/back/hlsl/writer.rs

@@ -104,6 +104,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
             entry_point_io: Vec::new(),
             named_expressions: crate::NamedExpressions::default(),
             wrapped: super::Wrapped::default(),
+            continue_ctx: back::continue_forward::ContinueCtx::default(),
             temp_access_chain: Vec::new(),
             need_bake_expressions: Default::default(),
         }
@@ -122,6 +123,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
         self.entry_point_io.clear();
         self.named_expressions.clear();
         self.wrapped.clear();
+        self.continue_ctx.clear();
         self.need_bake_expressions.clear();
     }
 
@@ -1439,6 +1441,151 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
         self.write_barrier(crate::Barrier::WORK_GROUP, level)
     }
 
+    /// Helper method used to write switches
+    fn write_switch(
+        &mut self,
+        module: &Module,
+        func_ctx: &back::FunctionCtx<'_>,
+        level: back::Level,
+        selector: Handle<crate::Expression>,
+        cases: &[crate::SwitchCase],
+    ) -> BackendResult {
+        // Write all cases
+        let indent_level_1 = level.next();
+        let indent_level_2 = indent_level_1.next();
+
+        // See docs of `back::continue_forward` module.
+        if let Some(variable) = self.continue_ctx.enter_switch(&mut self.namer) {
+            writeln!(self.out, "{level}bool {variable} = false;",)?;
+        };
+
+        // Check if there is only one body, by seeing if all except the last case are fall through
+        // with empty bodies. FXC doesn't handle these switches correctly, so
+        // we generate a `do {} while(false);` loop instead. There must be a default case, so there
+        // is no need to check if one of the cases would have matched.
+        let one_body = cases
+            .iter()
+            .rev()
+            .skip(1)
+            .all(|case| case.fall_through && case.body.is_empty());
+        if one_body {
+            // Start the do-while
+            writeln!(self.out, "{level}do {{")?;
+            // Note: Expressions have no side-effects so we don't need to emit selector expression.
+
+            // Body
+            if let Some(case) = cases.last() {
+                for sta in case.body.iter() {
+                    self.write_stmt(module, sta, func_ctx, indent_level_1)?;
+                }
+            }
+            // End do-while
+            writeln!(self.out, "{level}}} while(false);")?;
+        } else {
+            // Start the switch
+            write!(self.out, "{level}")?;
+            write!(self.out, "switch(")?;
+            self.write_expr(module, selector, func_ctx)?;
+            writeln!(self.out, ") {{")?;
+
+            for (i, case) in cases.iter().enumerate() {
+                match case.value {
+                    crate::SwitchValue::I32(value) => {
+                        write!(self.out, "{indent_level_1}case {value}:")?
+                    }
+                    crate::SwitchValue::U32(value) => {
+                        write!(self.out, "{indent_level_1}case {value}u:")?
+                    }
+                    crate::SwitchValue::Default => write!(self.out, "{indent_level_1}default:")?,
+                }
+
+                // The new block is not only stylistic, it plays a role here:
+                // We might end up having to write the same case body
+                // multiple times due to FXC not supporting fallthrough.
+                // Therefore, some `Expression`s written by `Statement::Emit`
+                // will end up having the same name (`_expr<handle_index>`).
+                // So we need to put each case in its own scope.
+                let write_block_braces = !(case.fall_through && case.body.is_empty());
+                if write_block_braces {
+                    writeln!(self.out, " {{")?;
+                } else {
+                    writeln!(self.out)?;
+                }
+
+                // Although FXC does support a series of case clauses before
+                // a block[^yes], it does not support fallthrough from a
+                // non-empty case block to the next[^no]. If this case has a
+                // non-empty body with a fallthrough, emulate that by
+                // duplicating the bodies of all the cases it would fall
+                // into as extensions of this case's own body. This makes
+                // the HLSL output potentially quadratic in the size of the
+                // Naga IR.
+                //
+                // [^yes]: ```hlsl
+                // case 1:
+                // case 2: do_stuff()
+                // ```
+                // [^no]: ```hlsl
+                // case 1: do_this();
+                // case 2: do_that();
+                // ```
+                if case.fall_through && !case.body.is_empty() {
+                    let curr_len = i + 1;
+                    let end_case_idx = curr_len
+                        + cases
+                            .iter()
+                            .skip(curr_len)
+                            .position(|case| !case.fall_through)
+                            .unwrap();
+                    let indent_level_3 = indent_level_2.next();
+                    for case in &cases[i..=end_case_idx] {
+                        writeln!(self.out, "{indent_level_2}{{")?;
+                        let prev_len = self.named_expressions.len();
+                        for sta in case.body.iter() {
+                            self.write_stmt(module, sta, func_ctx, indent_level_3)?;
+                        }
+                        // Clear all named expressions that were previously inserted by the statements in the block
+                        self.named_expressions.truncate(prev_len);
+                        writeln!(self.out, "{indent_level_2}}}")?;
+                    }
+
+                    let last_case = &cases[end_case_idx];
+                    if last_case.body.last().map_or(true, |s| !s.is_terminator()) {
+                        writeln!(self.out, "{indent_level_2}break;")?;
+                    }
+                } else {
+                    for sta in case.body.iter() {
+                        self.write_stmt(module, sta, func_ctx, indent_level_2)?;
+                    }
+                    if !case.fall_through && case.body.last().map_or(true, |s| !s.is_terminator()) {
+                        writeln!(self.out, "{indent_level_2}break;")?;
+                    }
+                }
+
+                if write_block_braces {
+                    writeln!(self.out, "{indent_level_1}}}")?;
+                }
+            }
+
+            writeln!(self.out, "{level}}}")?;
+        }
+
+        // Handle any forwarded continue statements.
+        use back::continue_forward::ExitControlFlow;
+        let op = match self.continue_ctx.exit_switch() {
+            ExitControlFlow::None => None,
+            ExitControlFlow::Continue { variable } => Some(("continue", variable)),
+            ExitControlFlow::Break { variable } => Some(("break", variable)),
+        };
+        if let Some((control_flow, variable)) = op {
+            writeln!(self.out, "{level}if ({variable}) {{")?;
+            writeln!(self.out, "{indent_level_1}{control_flow};")?;
+            writeln!(self.out, "{level}}}")?;
+        }
+
+        Ok(())
+    }
+
     /// Helper method used to write statements
     ///
     /// # Notes
@@ -1882,6 +2029,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 ref continuing,
                 break_if,
             } => {
+                self.continue_ctx.enter_loop();
                 let l2 = level.next();
                 if !continuing.is_empty() || break_if.is_some() {
                     let gate_name = self.namer.call("loop_init");
@@ -1908,10 +2056,18 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 for sta in body.iter() {
                     self.write_stmt(module, sta, func_ctx, l2)?;
                 }
-                writeln!(self.out, "{level}}}")?
+                writeln!(self.out, "{level}}}")?;
+                self.continue_ctx.exit_loop();
             }
             Statement::Break => writeln!(self.out, "{level}break;")?,
-            Statement::Continue => writeln!(self.out, "{level}continue;")?,
+            Statement::Continue => {
+                if let Some(variable) = self.continue_ctx.continue_encountered() {
+                    writeln!(self.out, "{level}{variable} = true;")?;
+                    writeln!(self.out, "{level}break;")?
+                } else {
+                    writeln!(self.out, "{level}continue;")?
+                }
+            }
             Statement::Barrier(barrier) => {
                 self.write_barrier(barrier, level)?;
             }
@@ -2063,100 +2219,7 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                 selector,
                 ref cases,
             } => {
-                // Start the switch
+                self.write_switch(module, func_ctx, level, selector, cases)?;
-                write!(self.out, "{level}")?;
-                write!(self.out, "switch(")?;
-                self.write_expr(module, selector, func_ctx)?;
-                writeln!(self.out, ") {{")?;
-
-                // Write all cases
-                let indent_level_1 = level.next();
-                let indent_level_2 = indent_level_1.next();
-
-                for (i, case) in cases.iter().enumerate() {
-                    match case.value {
-                        crate::SwitchValue::I32(value) => {
-                            write!(self.out, "{indent_level_1}case {value}:")?
-                        }
-                        crate::SwitchValue::U32(value) => {
-                            write!(self.out, "{indent_level_1}case {value}u:")?
-                        }
-                        crate::SwitchValue::Default => {
-                            write!(self.out, "{indent_level_1}default:")?
-                        }
-                    }
-
-                    // The new block is not only stylistic, it plays a role here:
-                    // We might end up having to write the same case body
-                    // multiple times due to FXC not supporting fallthrough.
-                    // Therefore, some `Expression`s written by `Statement::Emit`
-                    // will end up having the same name (`_expr<handle_index>`).
-                    // So we need to put each case in its own scope.
-                    let write_block_braces = !(case.fall_through && case.body.is_empty());
-                    if write_block_braces {
-                        writeln!(self.out, " {{")?;
-                    } else {
-                        writeln!(self.out)?;
-                    }
-
-                    // Although FXC does support a series of case clauses before
-                    // a block[^yes], it does not support fallthrough from a
-                    // non-empty case block to the next[^no]. If this case has a
-                    // non-empty body with a fallthrough, emulate that by
-                    // duplicating the bodies of all the cases it would fall
-                    // into as extensions of this case's own body. This makes
-                    // the HLSL output potentially quadratic in the size of the
-                    // Naga IR.
-                    //
-                    // [^yes]: ```hlsl
-                    // case 1:
-                    // case 2: do_stuff()
-                    // ```
-                    // [^no]: ```hlsl
-                    // case 1: do_this();
-                    // case 2: do_that();
-                    // ```
-                    if case.fall_through && !case.body.is_empty() {
-                        let curr_len = i + 1;
-                        let end_case_idx = curr_len
-                            + cases
-                                .iter()
-                                .skip(curr_len)
-                                .position(|case| !case.fall_through)
-                                .unwrap();
-                        let indent_level_3 = indent_level_2.next();
-                        for case in &cases[i..=end_case_idx] {
-                            writeln!(self.out, "{indent_level_2}{{")?;
-                            let prev_len = self.named_expressions.len();
-                            for sta in case.body.iter() {
-                                self.write_stmt(module, sta, func_ctx, indent_level_3)?;
-                            }
-                            // Clear all named expressions that were previously inserted by the statements in the block
-                            self.named_expressions.truncate(prev_len);
-                            writeln!(self.out, "{indent_level_2}}}")?;
-                        }
-
-                        let last_case = &cases[end_case_idx];
-                        if last_case.body.last().map_or(true, |s| !s.is_terminator()) {
-                            writeln!(self.out, "{indent_level_2}break;")?;
-                        }
-                    } else {
-                        for sta in case.body.iter() {
-                            self.write_stmt(module, sta, func_ctx, indent_level_2)?;
-                        }
-                        if !case.fall_through
-                            && case.body.last().map_or(true, |s| !s.is_terminator())
-                        {
-                            writeln!(self.out, "{indent_level_2}break;")?;
-                        }
-                    }
-
-                    if write_block_braces {
-                        writeln!(self.out, "{indent_level_1}}}")?;
-                    }
-                }
-
-                writeln!(self.out, "{level}}}")?
             }
             Statement::RayQuery { .. } => unreachable!(),
             Statement::SubgroupBallot { result, predicate } => {

naga/src/back/mod.rs

@@ -19,6 +19,9 @@ pub mod wgsl;
 #[cfg(any(hlsl_out, msl_out, spv_out, glsl_out))]
 pub mod pipeline_constants;
 
+#[cfg(any(feature = "hlsl-out", feature = "glsl-out"))]
+mod continue_forward;
+
 /// Names of vector components.
 pub const COMPONENTS: &[char] = &['x', 'y', 'z', 'w'];
 /// Indent for backends.

naga/tests/in/control-flow.wgsl

@@ -88,3 +88,96 @@ fn loop_switch_continue(x: i32) {
         }
     }
 }
+
+fn loop_switch_continue_nesting(x: i32, y: i32, z: i32) {
+    loop {
+        switch x {
+            case 1: {
+                continue;
+            }
+            case 2: {
+                switch y {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        loop {
+                            switch z {
+                                case 1: {
+                                    continue;
+                                }
+                                default: {}
+                            }
+                        }
+                    }
+                }
+            }
+            default: {}
+        }
+
+
+        // Degenerate switch with continue
+        switch y {
+            default: {
+                continue;
+            }
+        }
+    }
+
+    // In separate loop to avoid spv validation error:
+    // See https://github.com/gfx-rs/wgpu/issues/5658
+    loop {
+        // Nested degenerate switch with continue
+        switch y {
+            case 1, default: {
+                switch z {
+                    default: {
+                        continue;
+                    }
+                }
+            }
+        }
+    }
+}
+
+// Cases with some of the loop nested switches not containing continues.
+// See `continue_forward` module in `naga`.
+fn loop_switch_omit_continue_variable_checks(x: i32, y: i32, z: i32, w: i32) {
+    // switch in loop with no continues, we expect checks after the switch
+    // statement to not be generated
+    var pos: i32 = 0;
+    loop {
+        switch x {
+            case 1: {
+                pos = 1;
+            }
+            default: {}
+        }
+        // check here can be omitted
+    }
+
+    loop {
+        switch x {
+            case 1: {}
+            case 2: {
+                switch y {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        switch z {
+                            case 1: {
+                                pos = 2;
+                            }
+                            default: {}
+                        }
+                        // check here can be omitted
+                    }
+                }
+                // check needs to be generated here
+            }
+            default: {}
+        }
+        // check needs to be generated here
+    }
+}

naga/tests/out/glsl/control-flow.main.Compute.glsl

@@ -7,11 +7,9 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 
 
 void switch_default_break(int i) {
-    switch(i) {
+    do {
-        default: {
+        break;
-            break;
+    } while(false);
-        }
-    }
 }
 
 void switch_case_break() {
@@ -40,6 +38,110 @@ void loop_switch_continue(int x) {
     return;
 }
 
+void loop_switch_continue_nesting(int x_1, int y, int z) {
+    while(true) {
+        switch(x_1) {
+            case 1: {
+                continue;
+            }
+            case 2: {
+                switch(y) {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        while(true) {
+                            switch(z) {
+                                case 1: {
+                                    continue;
+                                }
+                                default: {
+                                    break;
+                                }
+                            }
+                        }
+                        break;
+                    }
+                }
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+        bool should_continue = false;
+        do {
+            should_continue = true;
+            break;
+        } while(false);
+        if (should_continue) {
+            continue;
+        }
+    }
+    while(true) {
+        bool should_continue_1 = false;
+        do {
+            do {
+                should_continue_1 = true;
+                break;
+            } while(false);
+            if (should_continue_1) {
+                break;
+            }
+        } while(false);
+        if (should_continue_1) {
+            continue;
+        }
+    }
+    return;
+}
+
+void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w) {
+    int pos_1 = 0;
+    while(true) {
+        switch(x_2) {
+            case 1: {
+                pos_1 = 1;
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+    }
+    while(true) {
+        switch(x_2) {
+            case 1: {
+                break;
+            }
+            case 2: {
+                switch(y_1) {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        switch(z_1) {
+                            case 1: {
+                                pos_1 = 2;
+                                break;
+                            }
+                            default: {
+                                break;
+                            }
+                        }
+                        break;
+                    }
+                }
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+    }
+    return;
+}
+
 void main() {
     uvec3 global_id = gl_GlobalInvocationID;
     int pos = 0;
@@ -47,12 +149,9 @@ void main() {
     barrier();
     memoryBarrierShared();
     barrier();
-    switch(1) {
+    do {
-        default: {
+        pos = 1;
-            pos = 1;
+    } while(false);
-            break;
-        }
-    }
     int _e4 = pos;
     switch(_e4) {
         case 1: {

naga/tests/out/hlsl/control-flow.hlsl

@@ -1,10 +1,8 @@
 void switch_default_break(int i)
 {
-    switch(i) {
+    do {
-        default: {
+        break;
-            break;
+    } while(false);
-        }
-    }
 }
 
 void switch_case_break()
@@ -23,14 +21,149 @@ void switch_case_break()
 void loop_switch_continue(int x)
 {
     while(true) {
+        bool should_continue = false;
         switch(x) {
             case 1: {
-                continue;
+                should_continue = true;
+                break;
             }
             default: {
                 break;
             }
         }
+        if (should_continue) {
+            continue;
+        }
+    }
+    return;
+}
+
+void loop_switch_continue_nesting(int x_1, int y, int z)
+{
+    while(true) {
+        bool should_continue_1 = false;
+        switch(x_1) {
+            case 1: {
+                should_continue_1 = true;
+                break;
+            }
+            case 2: {
+                switch(y) {
+                    case 1: {
+                        should_continue_1 = true;
+                        break;
+                    }
+                    default: {
+                        while(true) {
+                            bool should_continue_2 = false;
+                            switch(z) {
+                                case 1: {
+                                    should_continue_2 = true;
+                                    break;
+                                }
+                                default: {
+                                    break;
+                                }
+                            }
+                            if (should_continue_2) {
+                                continue;
+                            }
+                        }
+                        break;
+                    }
+                }
+                if (should_continue_1) {
+                    break;
+                }
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+        if (should_continue_1) {
+            continue;
+        }
+        bool should_continue_3 = false;
+        do {
+            should_continue_3 = true;
+            break;
+        } while(false);
+        if (should_continue_3) {
+            continue;
+        }
+    }
+    while(true) {
+        bool should_continue_4 = false;
+        do {
+            do {
+                should_continue_4 = true;
+                break;
+            } while(false);
+            if (should_continue_4) {
+                break;
+            }
+        } while(false);
+        if (should_continue_4) {
+            continue;
+        }
+    }
+    return;
+}
+
+void loop_switch_omit_continue_variable_checks(int x_2, int y_1, int z_1, int w)
+{
+    int pos_1 = 0;
+
+    while(true) {
+        bool should_continue_5 = false;
+        switch(x_2) {
+            case 1: {
+                pos_1 = 1;
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+    }
+    while(true) {
+        bool should_continue_6 = false;
+        switch(x_2) {
+            case 1: {
+                break;
+            }
+            case 2: {
+                switch(y_1) {
+                    case 1: {
+                        should_continue_6 = true;
+                        break;
+                    }
+                    default: {
+                        switch(z_1) {
+                            case 1: {
+                                pos_1 = 2;
+                                break;
+                            }
+                            default: {
+                                break;
+                            }
+                        }
+                        break;
+                    }
+                }
+                if (should_continue_6) {
+                    break;
+                }
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+        if (should_continue_6) {
+            continue;
+        }
     }
     return;
 }
@@ -42,12 +175,9 @@ void main(uint3 global_id : SV_DispatchThreadID)
 
     DeviceMemoryBarrierWithGroupSync();
     GroupMemoryBarrierWithGroupSync();
-    switch(1) {
+    do {
-        default: {
+        pos = 1;
-            pos = 1;
+    } while(false);
-            break;
-        }
-    }
     int _e4 = pos;
     switch(_e4) {
         case 1: {

naga/tests/out/msl/control-flow.msl

@@ -44,6 +44,114 @@ void loop_switch_continue(
     return;
 }
 
+void loop_switch_continue_nesting(
+    int x_1,
+    int y,
+    int z
+) {
+    while(true) {
+        switch(x_1) {
+            case 1: {
+                continue;
+            }
+            case 2: {
+                switch(y) {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        while(true) {
+                            switch(z) {
+                                case 1: {
+                                    continue;
+                                }
+                                default: {
+                                    break;
+                                }
+                            }
+                        }
+                        break;
+                    }
+                }
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+        switch(y) {
+            default: {
+                continue;
+            }
+        }
+    }
+    while(true) {
+        switch(y) {
+            case 1:
+            default: {
+                switch(z) {
+                    default: {
+                        continue;
+                    }
+                }
+                break;
+            }
+        }
+    }
+    return;
+}
+
+void loop_switch_omit_continue_variable_checks(
+    int x_2,
+    int y_1,
+    int z_1,
+    int w
+) {
+    int pos_1 = 0;
+    while(true) {
+        switch(x_2) {
+            case 1: {
+                pos_1 = 1;
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+    }
+    while(true) {
+        switch(x_2) {
+            case 1: {
+                break;
+            }
+            case 2: {
+                switch(y_1) {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        switch(z_1) {
+                            case 1: {
+                                pos_1 = 2;
+                                break;
+                            }
+                            default: {
+                                break;
+                            }
+                        }
+                        break;
+                    }
+                }
+                break;
+            }
+            default: {
+                break;
+            }
+        }
+    }
+    return;
+}
+
 struct main_Input {
 };
 kernel void main_(

naga/tests/out/spv/control-flow.spvasm

@@ -1,13 +1,13 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 69
+; Bound: 134
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %36 "main" %33
+OpEntryPoint GLCompute %104 "main" %101
-OpExecutionMode %36 LocalSize 1 1 1
+OpExecutionMode %104 LocalSize 1 1 1
-OpDecorate %33 BuiltIn GlobalInvocationId
+OpDecorate %101 BuiltIn GlobalInvocationId
 %2 = OpTypeVoid
 %4 = OpTypeInt 32 0
 %3 = OpTypeVector %4 3
@@ -15,19 +15,21 @@ OpDecorate %33 BuiltIn GlobalInvocationId
 %9 = OpTypeFunction %2 %5
 %15 = OpTypeFunction %2
 %16 = OpConstant  %5  0
-%34 = OpTypePointer Input %3
+%37 = OpTypeFunction %2 %5 %5 %5
-%33 = OpVariable  %34  Input
+%73 = OpTypeFunction %2 %5 %5 %5 %5
-%37 = OpConstant  %5  1
+%74 = OpConstant  %5  1
-%38 = OpConstant  %5  2
+%75 = OpConstant  %5  2
-%39 = OpConstant  %5  3
+%77 = OpTypePointer Function %5
-%40 = OpConstant  %5  4
+%102 = OpTypePointer Input %3
-%41 = OpConstant  %4  0
+%101 = OpVariable  %102  Input
-%43 = OpTypePointer Function %5
+%105 = OpConstant  %5  3
-%44 = OpConstantNull  %5
+%106 = OpConstant  %5  4
-%46 = OpConstant  %4  2
+%107 = OpConstant  %4  0
-%47 = OpConstant  %4  1
+%109 = OpConstantNull  %5
-%48 = OpConstant  %4  72
+%111 = OpConstant  %4  2
-%49 = OpConstant  %4  264
+%112 = OpConstant  %4  1
+%113 = OpConstant  %4  72
+%114 = OpConstant  %4  264
 %8 = OpFunction  %2  None %9
 %7 = OpFunctionParameter  %5
 %6 = OpLabel
@@ -76,63 +78,198 @@ OpBranch %25
 %26 = OpLabel
 OpReturn
 OpFunctionEnd
-%36 = OpFunction  %2  None %15
+%36 = OpFunction  %2  None %37
+%33 = OpFunctionParameter  %5
+%34 = OpFunctionParameter  %5
+%35 = OpFunctionParameter  %5
 %32 = OpLabel
-%42 = OpVariable  %43  Function %44
+OpBranch %38
-%35 = OpLoad  %3  %33
+%38 = OpLabel
-OpBranch %45
+OpBranch %39
+%39 = OpLabel
+OpLoopMerge %40 %42 None
+OpBranch %41
+%41 = OpLabel
+OpSelectionMerge %43 None
+OpSwitch %33 %46 1 %44 2 %45
+%44 = OpLabel
+OpBranch %42
 %45 = OpLabel
-OpControlBarrier %46 %47 %48
+OpSelectionMerge %47 None
-OpControlBarrier %46 %46 %49
+OpSwitch %34 %49 1 %48
-OpSelectionMerge %50 None
+%48 = OpLabel
-OpSwitch %37 %51
+OpBranch %42
-%51 = OpLabel
+%49 = OpLabel
-OpStore %42 %37
 OpBranch %50
 %50 = OpLabel
-%52 = OpLoad  %5  %42
+OpLoopMerge %51 %53 None
-OpSelectionMerge %53 None
+OpBranch %52
-OpSwitch %52 %58 1 %54 2 %55 3 %56 4 %56 5 %57 6 %58
+%52 = OpLabel
-%54 = OpLabel
+OpSelectionMerge %54 None
-OpStore %42 %16
+OpSwitch %35 %56 1 %55
-OpBranch %53
 %55 = OpLabel
-OpStore %42 %37
 OpBranch %53
 %56 = OpLabel
-OpStore %42 %38
+OpBranch %54
-OpBranch %53
+%54 = OpLabel
-%57 = OpLabel
-OpStore %42 %39
-OpBranch %53
-%58 = OpLabel
-OpStore %42 %40
 OpBranch %53
 %53 = OpLabel
-OpSelectionMerge %59 None
+OpBranch %50
-OpSwitch %41 %61 0 %60
+%51 = OpLabel
-%60 = OpLabel
+OpBranch %47
-OpBranch %59
+%47 = OpLabel
-%61 = OpLabel
+OpBranch %43
+%46 = OpLabel
+OpBranch %43
+%43 = OpLabel
+OpSelectionMerge %57 None
+OpSwitch %34 %58
+%58 = OpLabel
+OpBranch %42
+%57 = OpLabel
+OpBranch %42
+%42 = OpLabel
+OpBranch %39
+%40 = OpLabel
 OpBranch %59
 %59 = OpLabel
-%62 = OpLoad  %5  %42
+OpLoopMerge %60 %62 None
+OpBranch %61
+%61 = OpLabel
 OpSelectionMerge %63 None
-OpSwitch %62 %68 1 %64 2 %65 3 %66 4 %67
+OpSwitch %34 %64 1 %64
 %64 = OpLabel
-OpStore %42 %16
+OpSelectionMerge %65 None
-OpBranch %63
+OpSwitch %35 %66
-%65 = OpLabel
-OpStore %42 %37
-OpReturn
 %66 = OpLabel
-OpStore %42 %38
+OpBranch %62
-OpReturn
+%65 = OpLabel
-%67 = OpLabel
+OpBranch %63
-OpReturn
-%68 = OpLabel
-OpStore %42 %39
-OpReturn
 %63 = OpLabel
+OpBranch %62
+%62 = OpLabel
+OpBranch %59
+%60 = OpLabel
+OpReturn
+OpFunctionEnd
+%72 = OpFunction  %2  None %73
+%68 = OpFunctionParameter  %5
+%69 = OpFunctionParameter  %5
+%70 = OpFunctionParameter  %5
+%71 = OpFunctionParameter  %5
+%67 = OpLabel
+%76 = OpVariable  %77  Function %16
+OpBranch %78
+%78 = OpLabel
+OpBranch %79
+%79 = OpLabel
+OpLoopMerge %80 %82 None
+OpBranch %81
+%81 = OpLabel
+OpSelectionMerge %83 None
+OpSwitch %68 %85 1 %84
+%84 = OpLabel
+OpStore %76 %74
+OpBranch %83
+%85 = OpLabel
+OpBranch %83
+%83 = OpLabel
+OpBranch %82
+%82 = OpLabel
+OpBranch %79
+%80 = OpLabel
+OpBranch %86
+%86 = OpLabel
+OpLoopMerge %87 %89 None
+OpBranch %88
+%88 = OpLabel
+OpSelectionMerge %90 None
+OpSwitch %68 %93 1 %91 2 %92
+%91 = OpLabel
+OpBranch %90
+%92 = OpLabel
+OpSelectionMerge %94 None
+OpSwitch %69 %96 1 %95
+%95 = OpLabel
+OpBranch %89
+%96 = OpLabel
+OpSelectionMerge %97 None
+OpSwitch %70 %99 1 %98
+%98 = OpLabel
+OpStore %76 %75
+OpBranch %97
+%99 = OpLabel
+OpBranch %97
+%97 = OpLabel
+OpBranch %94
+%94 = OpLabel
+OpBranch %90
+%93 = OpLabel
+OpBranch %90
+%90 = OpLabel
+OpBranch %89
+%89 = OpLabel
+OpBranch %86
+%87 = OpLabel
+OpReturn
+OpFunctionEnd
+%104 = OpFunction  %2  None %15
+%100 = OpLabel
+%108 = OpVariable  %77  Function %109
+%103 = OpLoad  %3  %101
+OpBranch %110
+%110 = OpLabel
+OpControlBarrier %111 %112 %113
+OpControlBarrier %111 %111 %114
+OpSelectionMerge %115 None
+OpSwitch %74 %116
+%116 = OpLabel
+OpStore %108 %74
+OpBranch %115
+%115 = OpLabel
+%117 = OpLoad  %5  %108
+OpSelectionMerge %118 None
+OpSwitch %117 %123 1 %119 2 %120 3 %121 4 %121 5 %122 6 %123
+%119 = OpLabel
+OpStore %108 %16
+OpBranch %118
+%120 = OpLabel
+OpStore %108 %74
+OpBranch %118
+%121 = OpLabel
+OpStore %108 %75
+OpBranch %118
+%122 = OpLabel
+OpStore %108 %105
+OpBranch %118
+%123 = OpLabel
+OpStore %108 %106
+OpBranch %118
+%118 = OpLabel
+OpSelectionMerge %124 None
+OpSwitch %107 %126 0 %125
+%125 = OpLabel
+OpBranch %124
+%126 = OpLabel
+OpBranch %124
+%124 = OpLabel
+%127 = OpLoad  %5  %108
+OpSelectionMerge %128 None
+OpSwitch %127 %133 1 %129 2 %130 3 %131 4 %132
+%129 = OpLabel
+OpStore %108 %16
+OpBranch %128
+%130 = OpLabel
+OpStore %108 %74
+OpReturn
+%131 = OpLabel
+OpStore %108 %75
+OpReturn
+%132 = OpLabel
+OpReturn
+%133 = OpLabel
+OpStore %108 %105
+OpReturn
+%128 = OpLabel
 OpReturn
 OpFunctionEnd

naga/tests/out/wgsl/control-flow.wgsl

@@ -30,6 +30,92 @@ fn loop_switch_continue(x: i32) {
     return;
 }
 
+fn loop_switch_continue_nesting(x_1: i32, y: i32, z: i32) {
+    loop {
+        switch x_1 {
+            case 1: {
+                continue;
+            }
+            case 2: {
+                switch y {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        loop {
+                            switch z {
+                                case 1: {
+                                    continue;
+                                }
+                                default: {
+                                }
+                            }
+                        }
+                    }
+                }
+            }
+            default: {
+            }
+        }
+        switch y {
+            default: {
+                continue;
+            }
+        }
+    }
+    loop {
+        switch y {
+            case 1, default: {
+                switch z {
+                    default: {
+                        continue;
+                    }
+                }
+            }
+        }
+    }
+    return;
+}
+
+fn loop_switch_omit_continue_variable_checks(x_2: i32, y_1: i32, z_1: i32, w: i32) {
+    var pos_1: i32 = 0i;
+
+    loop {
+        switch x_2 {
+            case 1: {
+                pos_1 = 1i;
+            }
+            default: {
+            }
+        }
+    }
+    loop {
+        switch x_2 {
+            case 1: {
+            }
+            case 2: {
+                switch y_1 {
+                    case 1: {
+                        continue;
+                    }
+                    default: {
+                        switch z_1 {
+                            case 1: {
+                                pos_1 = 2i;
+                            }
+                            default: {
+                            }
+                        }
+                    }
+                }
+            }
+            default: {
+            }
+        }
+    }
+    return;
+}
+
 @compute @workgroup_size(1, 1, 1) 
 fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
     var pos: i32;

naga/xtask/src/validate.rs

@@ -208,7 +208,10 @@ fn validate_spirv(path: &Path, spirv_as: &str, spirv_val: &str) -> anyhow::Resul
         buf
     };
     let expected_header_prefix = "; Version: ";
-    let Some(version) = second_line.strip_prefix(expected_header_prefix) else {
+    let Some(version) = second_line
+        .strip_prefix(expected_header_prefix)
+        .map(str::trim)
+    else {
         bail!("no {expected_header_prefix:?} header found in {path:?}");
     };
     let file = open_file(path)?;
@@ -222,7 +225,18 @@ fn validate_spirv(path: &Path, spirv_as: &str, spirv_val: &str) -> anyhow::Resul
     let child = spirv_as_cmd
         .spawn()
         .with_context(|| format!("failed to spawn {spirv_as_cmd:?}"))?;
-    EasyCommand::new(spirv_val, |cmd| cmd.stdin(child.stdout.unwrap())).success()
+    let error_message = || {
+        format!(
+            "Failed to validate {path:?}.
+Note: Labels and line numbers will not match the input file.
+      Use this command to view the corresponding spvasm:
+      '{spirv_as} --target-env spv{version} {} -o - | spirv-dis'\n",
+            path.display(),
+        )
+    };
+    EasyCommand::new(spirv_val, |cmd| cmd.stdin(child.stdout.unwrap()))
+        .success()
+        .with_context(error_message)
 }
 
 fn validate_metal(path: &Path, xcrun: &str) -> anyhow::Result<()> {

tests/src/init.rs

@@ -11,7 +11,7 @@ pub fn init_logger() {
 }
 
 /// Initialize a wgpu instance with the options from the environment.
-pub fn initialize_instance() -> Instance {
+pub fn initialize_instance(force_fxc: bool) -> Instance {
     // We ignore `WGPU_BACKEND` for now, merely using test filtering to only run a single backend's tests.
     //
     // We can potentially work support back into the test runner in the future, but as the adapters are matched up
@@ -27,7 +27,13 @@ pub fn initialize_instance() -> Instance {
     } else {
         Backends::all()
     };
-    let dx12_shader_compiler = wgpu::util::dx12_shader_compiler_from_env().unwrap_or_default();
+    // Some tests need to be able to force demote to FXC, to specifically test workarounds for FXC
+    // behavior.
+    let dx12_shader_compiler = if force_fxc {
+        wgpu::Dx12Compiler::Fxc
+    } else {
+        wgpu::util::dx12_shader_compiler_from_env().unwrap_or_default()
+    };
     let gles_minor_version = wgpu::util::gles_minor_version_from_env().unwrap_or_default();
     Instance::new(wgpu::InstanceDescriptor {
         backends,
@@ -38,8 +44,11 @@ pub fn initialize_instance() -> Instance {
 }
 
 /// Initialize a wgpu adapter, taking the `n`th adapter from the instance.
-pub async fn initialize_adapter(adapter_index: usize) -> (Instance, Adapter, Option<SurfaceGuard>) {
+pub async fn initialize_adapter(
-    let instance = initialize_instance();
+    adapter_index: usize,
+    force_fxc: bool,
+) -> (Instance, Adapter, Option<SurfaceGuard>) {
+    let instance = initialize_instance(force_fxc);
     #[allow(unused_variables)]
     let surface: Option<wgpu::Surface>;
     let surface_guard: Option<SurfaceGuard>;

tests/src/params.rs

@@ -19,6 +19,11 @@ pub struct TestParameters {
     pub required_downlevel_caps: DownlevelCapabilities,
     pub required_limits: Limits,
 
+    /// On Dx12, specifically test against the Fxc compiler.
+    ///
+    /// For testing workarounds to Fxc bugs.
+    pub force_fxc: bool,
+
     /// Conditions under which this test should be skipped.
     pub skips: Vec<FailureCase>,
 
@@ -32,6 +37,7 @@ impl Default for TestParameters {
             required_features: Features::empty(),
             required_downlevel_caps: LOWEST_DOWNLEVEL_PROPERTIES,
             required_limits: Limits::downlevel_webgl2_defaults(),
+            force_fxc: false,
             skips: Vec::new(),
             failures: Vec::new(),
         }
@@ -63,6 +69,11 @@ impl TestParameters {
         self
     }
 
+    pub fn force_fxc(mut self, force_fxc: bool) -> Self {
+        self.force_fxc = force_fxc;
+        self
+    }
+
     /// Mark the test as always failing, but not to be skipped.
     pub fn expect_fail(mut self, when: FailureCase) -> Self {
         self.failures.push(when);

tests/src/run.rs

@@ -42,7 +42,8 @@ pub async fn execute_test(
 
     let _test_guard = isolation::OneTestPerProcessGuard::new();
 
-    let (instance, adapter, _surface_guard) = initialize_adapter(adapter_index).await;
+    let (instance, adapter, _surface_guard) =
+        initialize_adapter(adapter_index, config.params.force_fxc).await;
 
     let adapter_info = adapter.get_info();
     let adapter_downlevel_capabilities = adapter.get_downlevel_capabilities();

tests/tests/create_surface_error.rs

@@ -6,7 +6,7 @@
 #[wasm_bindgen_test::wasm_bindgen_test]
 fn canvas_get_context_returned_null() {
     // Not using the normal testing infrastructure because that goes straight to creating the canvas for us.
-    let instance = wgpu_test::initialize_instance();
+    let instance = wgpu_test::initialize_instance(false);
     // Create canvas
     let canvas = wgpu_test::initialize_html_canvas();
 

tests/tests/device.rs

@@ -107,7 +107,7 @@ static REQUEST_DEVICE_ERROR_MESSAGE_NATIVE: GpuTestConfiguration =
 async fn request_device_error_message() {
     // Not using initialize_test() because that doesn't let us catch the error
     // nor .await anything
-    let (_instance, adapter, _surface_guard) = wgpu_test::initialize_adapter(0).await;
+    let (_instance, adapter, _surface_guard) = wgpu_test::initialize_adapter(0, false).await;
 
     let device_error = adapter
         .request_device(

tests/tests/regression/issue_4485.rs

@@ -0,0 +1,106 @@
+use wgpu_test::{gpu_test, image, GpuTestConfiguration, TestParameters, TestingContext};
+
+/// FXC doesn't accept `continue` inside a switch. Instead we store a flag for whether
+/// the loop should continue that is checked after the switch.
+///
+/// See <https://github.com/gfx-rs/wgpu/issues/4485>.
+///
+/// The shader will fail to compile on Dx12 with FXC without this fix.
+///
+/// This also tests that shaders generated with this fix execute correctly.
+#[gpu_test]
+static CONTINUE_SWITCH: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default().force_fxc(true))
+    .run_async(|ctx| async move { test_impl(&ctx).await });
+
+async fn test_impl(ctx: &TestingContext) {
+    const TEXTURE_HEIGHT: u32 = 2;
+    const TEXTURE_WIDTH: u32 = 2;
+    const BUFFER_SIZE: usize = (TEXTURE_WIDTH * TEXTURE_HEIGHT * 4) as usize;
+
+    let texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
+        label: Some("Offscreen texture"),
+        size: wgpu::Extent3d {
+            width: TEXTURE_WIDTH,
+            height: TEXTURE_HEIGHT,
+            depth_or_array_layers: 1,
+        },
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: wgpu::TextureFormat::Rgba8Unorm,
+        usage: wgpu::TextureUsages::COPY_SRC | wgpu::TextureUsages::RENDER_ATTACHMENT,
+        view_formats: &[],
+    });
+    let texture_view = texture.create_view(&wgpu::TextureViewDescriptor::default());
+
+    let shader = ctx
+        .device
+        .create_shader_module(wgpu::include_wgsl!("issue_4514.wgsl"));
+
+    let pipeline = ctx
+        .device
+        .create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+            label: Some("Pipeline"),
+            layout: None,
+            vertex: wgpu::VertexState {
+                module: &shader,
+                entry_point: "vs_main",
+                compilation_options: Default::default(),
+                buffers: &[],
+            },
+            primitive: wgpu::PrimitiveState::default(),
+            depth_stencil: None,
+            multisample: wgpu::MultisampleState::default(),
+            fragment: Some(wgpu::FragmentState {
+                module: &shader,
+                entry_point: "fs_main",
+                compilation_options: Default::default(),
+                targets: &[Some(wgpu::ColorTargetState {
+                    format: wgpu::TextureFormat::Rgba8Unorm,
+                    blend: None,
+                    write_mask: wgpu::ColorWrites::ALL,
+                })],
+            }),
+            multiview: None,
+            cache: None,
+        });
+
+    let readback_buffer = image::ReadbackBuffers::new(&ctx.device, &texture);
+    {
+        let mut encoder = ctx
+            .device
+            .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+        {
+            let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+                label: Some("Renderpass"),
+                color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                    view: &texture_view,
+                    resolve_target: None,
+                    ops: wgpu::Operations {
+                        // Important: this isn't the color expected below
+                        load: wgpu::LoadOp::Clear(wgpu::Color {
+                            r: 0.0,
+                            g: 0.0,
+                            b: 0.0,
+                            a: 0.0,
+                        }),
+                        store: wgpu::StoreOp::Store,
+                    },
+                })],
+                depth_stencil_attachment: None,
+                timestamp_writes: None,
+                occlusion_query_set: None,
+            });
+            render_pass.set_pipeline(&pipeline);
+            render_pass.draw(0..3, 0..1);
+        }
+        readback_buffer.copy_from(&ctx.device, &mut encoder, &texture);
+        ctx.queue.submit(Some(encoder.finish()));
+    }
+
+    let expected_data = [255; BUFFER_SIZE];
+    readback_buffer
+        .assert_buffer_contents(ctx, &expected_data)
+        .await;
+}

tests/tests/regression/issue_4485.wgsl

@@ -0,0 +1,108 @@
+// meant to be called with 3 vertex indices: 0, 1, 2
+// draws one large triangle over the clip space like this:
+// (the asterisks represent the clip space bounds)
+//-1,1           1,1
+// ---------------------------------
+// |              *              .
+// |              *           .
+// |              *        .
+// |              *      .
+// |              *    .
+// |              * .
+// |***************
+// |            . 1,-1
+// |          .
+// |       .
+// |     .
+// |   .
+// |.
+@vertex
+fn vs_main(@builtin(vertex_index) vertex_index: u32) ->  @builtin(position) vec4<f32> {
+    let x = i32(vertex_index) / 2;
+    let y = i32(vertex_index) & 1;
+    return vec4<f32>(
+        f32(x) * 4.0 - 1.0,
+        1.0 - f32(y) * 4.0,
+        0.0, 1.0
+    );
+}
+
+
+@fragment
+fn fs_main() -> @location(0) vec4<f32> {
+    var x = 0.0;
+    loop {
+        if x != 0.0 { break; }
+        x = 0.5;
+        // Compiled to a do-while in hlsl and glsl,
+        // we want to confirm that continue applies to outer loop.
+        switch 0 {
+            default {
+                x = 1.0;
+                continue;
+            }
+        }
+        x = 0.0;
+    }
+    // expect X == 1.0
+
+    var y = 0.0;
+    loop {
+        if y != 0.0 { break; }
+        y = 0.5;
+        switch 1 {
+            case 0 {
+                continue;
+            }
+            case 1 {}
+        }
+        // test that loop doesn't continue after the switch when the continue case wasn't executed
+        y = 1.0;
+        break;
+    }
+    // expect y == 1.0
+
+    var z = 0.0;
+    loop {
+        if z != 0.0 { break; }
+        switch 0 {
+            case 0 {
+                z = 0.5;
+            }
+            case 1 {
+                z = 0.5;
+            }
+        }
+        // test that loop doesn't continue after the switch that contains no continue statements
+        z = 1.0
+    }
+    // expect z == 1.0
+
+    var w = 0.0;
+    loop {
+        if w != 0.0 { break; }
+        switch 0 {
+            case 0 {
+                loop {
+                    // continue in loop->switch->loop->switch->switch should affect inner loop
+                    switch 1 {
+                        case 0 {}
+                        case 1 {
+                            switch 0 {
+                                default { continue; }
+                            }
+                        }
+                    }
+                    w = 0.5
+                }
+            }
+            case 1 {
+                w = 0.5;
+            }
+        }
+        if w == 0.0 { w = 1.0; }
+    }
+    // expect w == 1.0
+
+    return vec4<f32>(x, y, z, w);
+}

tests/tests/regression/issue_4514.rs

@@ -0,0 +1,106 @@
+use wgpu_test::{gpu_test, image, GpuTestConfiguration, TestParameters, TestingContext};
+
+/// FXC and potentially some glsl consumers have a bug when handling switch statements on a constant
+/// with just a default case. (not sure if the constant part is relevant)
+/// See <https://github.com/gfx-rs/wgpu/issues/4514>.
+///
+/// This test will fail on Dx12 with FXC if this issue is not worked around.
+///
+/// So far no specific buggy glsl consumers have been identified and it isn't known whether the
+/// bug is avoided there.
+#[gpu_test]
+static DEGENERATE_SWITCH: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default().force_fxc(true))
+    .run_async(|ctx| async move { test_impl(&ctx).await });
+
+async fn test_impl(ctx: &TestingContext) {
+    const TEXTURE_HEIGHT: u32 = 2;
+    const TEXTURE_WIDTH: u32 = 2;
+    const BUFFER_SIZE: usize = (TEXTURE_WIDTH * TEXTURE_HEIGHT * 4) as usize;
+
+    let texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
+        label: Some("Offscreen texture"),
+        size: wgpu::Extent3d {
+            width: TEXTURE_WIDTH,
+            height: TEXTURE_HEIGHT,
+            depth_or_array_layers: 1,
+        },
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format: wgpu::TextureFormat::Rgba8Unorm,
+        usage: wgpu::TextureUsages::COPY_SRC | wgpu::TextureUsages::RENDER_ATTACHMENT,
+        view_formats: &[],
+    });
+    let texture_view = texture.create_view(&wgpu::TextureViewDescriptor::default());
+
+    let shader = ctx
+        .device
+        .create_shader_module(wgpu::include_wgsl!("issue_4514.wgsl"));
+
+    let pipeline = ctx
+        .device
+        .create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+            label: Some("Pipeline"),
+            layout: None,
+            vertex: wgpu::VertexState {
+                module: &shader,
+                entry_point: "vs_main",
+                compilation_options: Default::default(),
+                buffers: &[],
+            },
+            primitive: wgpu::PrimitiveState::default(),
+            depth_stencil: None,
+            multisample: wgpu::MultisampleState::default(),
+            fragment: Some(wgpu::FragmentState {
+                module: &shader,
+                entry_point: "fs_main",
+                compilation_options: Default::default(),
+                targets: &[Some(wgpu::ColorTargetState {
+                    format: wgpu::TextureFormat::Rgba8Unorm,
+                    blend: None,
+                    write_mask: wgpu::ColorWrites::ALL,
+                })],
+            }),
+            multiview: None,
+            cache: None,
+        });
+
+    let readback_buffer = image::ReadbackBuffers::new(&ctx.device, &texture);
+    {
+        let mut encoder = ctx
+            .device
+            .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
+        {
+            let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+                label: Some("Renderpass"),
+                color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                    view: &texture_view,
+                    resolve_target: None,
+                    ops: wgpu::Operations {
+                        // Important: this isn't the color expected below
+                        load: wgpu::LoadOp::Clear(wgpu::Color {
+                            r: 0.0,
+                            g: 0.0,
+                            b: 0.0,
+                            a: 0.0,
+                        }),
+                        store: wgpu::StoreOp::Store,
+                    },
+                })],
+                depth_stencil_attachment: None,
+                timestamp_writes: None,
+                occlusion_query_set: None,
+            });
+            render_pass.set_pipeline(&pipeline);
+            render_pass.draw(0..3, 0..1);
+        }
+        readback_buffer.copy_from(&ctx.device, &mut encoder, &texture);
+        ctx.queue.submit(Some(encoder.finish()));
+    }
+
+    let expected_data = [255; BUFFER_SIZE];
+    readback_buffer
+        .assert_buffer_contents(ctx, &expected_data)
+        .await;
+}

tests/tests/regression/issue_4514.wgsl

@@ -0,0 +1,68 @@
+// meant to be called with 3 vertex indices: 0, 1, 2
+// draws one large triangle over the clip space like this:
+// (the asterisks represent the clip space bounds)
+//-1,1           1,1
+// ---------------------------------
+// |              *              .
+// |              *           .
+// |              *        .
+// |              *      .
+// |              *    .
+// |              * .
+// |***************
+// |            . 1,-1
+// |          .
+// |       .
+// |     .
+// |   .
+// |.
+@vertex
+fn vs_main(@builtin(vertex_index) vertex_index: u32) ->  @builtin(position) vec4<f32> {
+    let x = i32(vertex_index) / 2;
+    let y = i32(vertex_index) & 1;
+    return vec4<f32>(
+        f32(x) * 4.0 - 1.0,
+        1.0 - f32(y) * 4.0,
+        0.0, 1.0
+    );
+}
+
+
+@fragment
+fn fs_main(@builtin(position) coord_in: vec4<f32>) -> @location(0) vec4<f32> {
+    var x = 0.0;
+    // Succeeds on FXC without workaround.
+    switch i32(coord_in.x) {
+        default {
+            x = 1.0;
+        }
+    }
+    var y = 0.0;
+    // Fails on FXC without workaround.
+    // (even if we adjust switch above to give different x values based on the input coord)
+    switch i32(x * 30.0) {
+        default {
+            y = 1.0;
+        }
+    }
+    var z = 0.0;
+    // Multiple cases with a single body also fails on FXC without a workaround.
+    switch 0 {
+        case 0, 2, default {
+            z = 1.0;
+        }
+    }
+
+    var w = 0.0;
+    // Succeeds on FXC without workaround.
+    switch 0 {
+        case 0 {
+            w = 1.0;
+        }
+        default {
+            w = 1.0;
+        }
+    }
+
+    return vec4<f32>(x, y, z, w);
+}

tests/tests/root.rs

@@ -3,6 +3,8 @@ mod regression {
     mod issue_3457;
     mod issue_4024;
     mod issue_4122;
+    mod issue_4485;
+    mod issue_4514;
     mod issue_5553;
 }