diff --git a/CHANGELOG.md b/CHANGELOG.md index c781d3f604..c9eccafcda 100644 --- a/CHANGELOG.md +++ b/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 diff --git a/naga/src/back/continue_forward.rs b/naga/src/back/continue_forward.rs new file mode 100644 index 0000000000..cecb93a837 --- /dev/null +++ b/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 , 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 , 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, + + /// 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, + }, + /// 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, + }, +} + +/// 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, +} + +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> { + 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 + } + } +} diff --git a/naga/src/back/glsl/mod.rs b/naga/src/back/glsl/mod.rs index bc2d2a90d8..7ad1f3c597 100644 --- a/naga/src/back/glsl/mod.rs +++ b/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, /// 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 { - 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:")?, - } + // Some GLSL consumers may not handle switches with a single + // body correctly: See wgpu#4514. Write such switch statements + // as a `do {} while(false);` loop instead. + // + // 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. - let write_block_braces = !(case.fall_through && case.body.is_empty()); - if write_block_braces { - writeln!(self.out, " {{")?; - } else { - writeln!(self.out)?; + // Body + if let Some(case) = cases.last() { + for sta in case.body.iter() { + self.write_stmt(sta, ctx, l2)?; + } } - - for sta in case.body.iter() { - self.write_stmt(sta, ctx, l2.next())?; + // 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:")?, + } - if !case.fall_through && case.body.last().map_or(true, |s| !s.is_terminator()) { - writeln!(self.out, "{}break;", l2.next())?; - } + 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}}}")?; + if write_block_braces { + writeln!(self.out, "{l2}}}")?; + } } - } - writeln!(self.out, "{level}}}")? + 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}")?; - writeln!(self.out, "continue;")? + // Sometimes we must render a `Continue` statement as a `break`. + // 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 } => { diff --git a/naga/src/back/hlsl/mod.rs b/naga/src/back/hlsl/mod.rs index 49ff07ebf2..d28b387bf7 100644 --- a/naga/src/back/hlsl/mod.rs +++ b/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. diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index afa12cccab..982bf0cfea 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/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, + 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`). + // 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 - 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`). - // 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}}}")? + self.write_switch(module, func_ctx, level, selector, cases)?; } Statement::RayQuery { .. } => unreachable!(), Statement::SubgroupBallot { result, predicate } => { diff --git a/naga/src/back/mod.rs b/naga/src/back/mod.rs index cd9496e3ff..43d88a437d 100644 --- a/naga/src/back/mod.rs +++ b/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. diff --git a/naga/tests/in/control-flow.wgsl b/naga/tests/in/control-flow.wgsl index 5a0ef1cbbf..a25c899a44 100644 --- a/naga/tests/in/control-flow.wgsl +++ b/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 + } +} diff --git a/naga/tests/out/glsl/control-flow.main.Compute.glsl b/naga/tests/out/glsl/control-flow.main.Compute.glsl index b877f9cb69..391fca84f4 100644 --- a/naga/tests/out/glsl/control-flow.main.Compute.glsl +++ b/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) { - default: { - break; - } - } + do { + 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) { - default: { - pos = 1; - break; - } - } + do { + pos = 1; + } while(false); int _e4 = pos; switch(_e4) { case 1: { diff --git a/naga/tests/out/hlsl/control-flow.hlsl b/naga/tests/out/hlsl/control-flow.hlsl index 1e253add21..2438858a8a 100644 --- a/naga/tests/out/hlsl/control-flow.hlsl +++ b/naga/tests/out/hlsl/control-flow.hlsl @@ -1,10 +1,8 @@ void switch_default_break(int i) { - switch(i) { - default: { - break; - } - } + do { + 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) { - default: { - pos = 1; - break; - } - } + do { + pos = 1; + } while(false); int _e4 = pos; switch(_e4) { case 1: { diff --git a/naga/tests/out/msl/control-flow.msl b/naga/tests/out/msl/control-flow.msl index 0d0e082e41..11771693aa 100644 --- a/naga/tests/out/msl/control-flow.msl +++ b/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_( diff --git a/naga/tests/out/spv/control-flow.spvasm b/naga/tests/out/spv/control-flow.spvasm index 2fc9337cfe..f3c3644b4f 100644 --- a/naga/tests/out/spv/control-flow.spvasm +++ b/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 -OpExecutionMode %36 LocalSize 1 1 1 -OpDecorate %33 BuiltIn GlobalInvocationId +OpEntryPoint GLCompute %104 "main" %101 +OpExecutionMode %104 LocalSize 1 1 1 +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 -%33 = OpVariable %34 Input -%37 = OpConstant %5 1 -%38 = OpConstant %5 2 -%39 = OpConstant %5 3 -%40 = OpConstant %5 4 -%41 = OpConstant %4 0 -%43 = OpTypePointer Function %5 -%44 = OpConstantNull %5 -%46 = OpConstant %4 2 -%47 = OpConstant %4 1 -%48 = OpConstant %4 72 -%49 = OpConstant %4 264 +%37 = OpTypeFunction %2 %5 %5 %5 +%73 = OpTypeFunction %2 %5 %5 %5 %5 +%74 = OpConstant %5 1 +%75 = OpConstant %5 2 +%77 = OpTypePointer Function %5 +%102 = OpTypePointer Input %3 +%101 = OpVariable %102 Input +%105 = OpConstant %5 3 +%106 = OpConstant %5 4 +%107 = OpConstant %4 0 +%109 = OpConstantNull %5 +%111 = OpConstant %4 2 +%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 -%35 = OpLoad %3 %33 -OpBranch %45 +OpBranch %38 +%38 = OpLabel +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 -OpControlBarrier %46 %46 %49 -OpSelectionMerge %50 None -OpSwitch %37 %51 -%51 = OpLabel -OpStore %42 %37 +OpSelectionMerge %47 None +OpSwitch %34 %49 1 %48 +%48 = OpLabel +OpBranch %42 +%49 = OpLabel OpBranch %50 %50 = OpLabel -%52 = OpLoad %5 %42 -OpSelectionMerge %53 None -OpSwitch %52 %58 1 %54 2 %55 3 %56 4 %56 5 %57 6 %58 -%54 = OpLabel -OpStore %42 %16 -OpBranch %53 +OpLoopMerge %51 %53 None +OpBranch %52 +%52 = OpLabel +OpSelectionMerge %54 None +OpSwitch %35 %56 1 %55 %55 = OpLabel -OpStore %42 %37 OpBranch %53 %56 = OpLabel -OpStore %42 %38 -OpBranch %53 -%57 = OpLabel -OpStore %42 %39 -OpBranch %53 -%58 = OpLabel -OpStore %42 %40 +OpBranch %54 +%54 = OpLabel OpBranch %53 %53 = OpLabel -OpSelectionMerge %59 None -OpSwitch %41 %61 0 %60 -%60 = OpLabel -OpBranch %59 -%61 = OpLabel +OpBranch %50 +%51 = OpLabel +OpBranch %47 +%47 = 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 -OpBranch %63 -%65 = OpLabel -OpStore %42 %37 -OpReturn +OpSelectionMerge %65 None +OpSwitch %35 %66 %66 = OpLabel -OpStore %42 %38 +OpBranch %62 +%65 = OpLabel +OpBranch %63 +%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 -%68 = OpLabel -OpStore %42 %39 +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 -%63 = OpLabel +%131 = OpLabel +OpStore %108 %75 +OpReturn +%132 = OpLabel +OpReturn +%133 = OpLabel +OpStore %108 %105 +OpReturn +%128 = OpLabel OpReturn OpFunctionEnd \ No newline at end of file diff --git a/naga/tests/out/wgsl/control-flow.wgsl b/naga/tests/out/wgsl/control-flow.wgsl index dcc3f90365..ad071af58a 100644 --- a/naga/tests/out/wgsl/control-flow.wgsl +++ b/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) { var pos: i32; diff --git a/naga/xtask/src/validate.rs b/naga/xtask/src/validate.rs index d90ee8d84a..fa330f0a96 100644 --- a/naga/xtask/src/validate.rs +++ b/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<()> { diff --git a/tests/src/init.rs b/tests/src/init.rs index 3a11b3abe3..140bb202fc 100644 --- a/tests/src/init.rs +++ b/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) { - let instance = initialize_instance(); +pub async fn initialize_adapter( + adapter_index: usize, + force_fxc: bool, +) -> (Instance, Adapter, Option) { + let instance = initialize_instance(force_fxc); #[allow(unused_variables)] let surface: Option; let surface_guard: Option; diff --git a/tests/src/params.rs b/tests/src/params.rs index 2f54e65bbb..e5d50a4859 100644 --- a/tests/src/params.rs +++ b/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, @@ -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); diff --git a/tests/src/run.rs b/tests/src/run.rs index 82ddb93399..82c1d34e69 100644 --- a/tests/src/run.rs +++ b/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(); diff --git a/tests/tests/create_surface_error.rs b/tests/tests/create_surface_error.rs index 87aeb15726..e3b48cb757 100644 --- a/tests/tests/create_surface_error.rs +++ b/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(); diff --git a/tests/tests/device.rs b/tests/tests/device.rs index f932faa2f1..ae463cca46 100644 --- a/tests/tests/device.rs +++ b/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( diff --git a/tests/tests/regression/issue_4485.rs b/tests/tests/regression/issue_4485.rs new file mode 100644 index 0000000000..101712fe02 --- /dev/null +++ b/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 . +/// +/// 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; +} diff --git a/tests/tests/regression/issue_4485.wgsl b/tests/tests/regression/issue_4485.wgsl new file mode 100644 index 0000000000..e72ed6d1ea --- /dev/null +++ b/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 { + let x = i32(vertex_index) / 2; + let y = i32(vertex_index) & 1; + return vec4( + f32(x) * 4.0 - 1.0, + 1.0 - f32(y) * 4.0, + 0.0, 1.0 + ); +} + + +@fragment +fn fs_main() -> @location(0) vec4 { + 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(x, y, z, w); +} diff --git a/tests/tests/regression/issue_4514.rs b/tests/tests/regression/issue_4514.rs new file mode 100644 index 0000000000..f447f879bf --- /dev/null +++ b/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 . +/// +/// 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; +} diff --git a/tests/tests/regression/issue_4514.wgsl b/tests/tests/regression/issue_4514.wgsl new file mode 100644 index 0000000000..d4bd2f80c0 --- /dev/null +++ b/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 { + let x = i32(vertex_index) / 2; + let y = i32(vertex_index) & 1; + return vec4( + 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) -> @location(0) vec4 { + 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(x, y, z, w); +} diff --git a/tests/tests/root.rs b/tests/tests/root.rs index 384cfcf78f..df0dce5fed 100644 --- a/tests/tests/root.rs +++ b/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; }