diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index d2c69e1312..92dfe06266 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -511,6 +511,8 @@ libtint_source_set("libtint_core_all_src") { "transform/localize_struct_array_assignment.h", "transform/manager.cc", "transform/manager.h", + "transform/merge_return.cc", + "transform/merge_return.h", "transform/module_scope_var_to_entry_point_param.cc", "transform/module_scope_var_to_entry_point_param.h", "transform/multiplanar_external_texture.cc", @@ -1224,6 +1226,7 @@ if (tint_build_unittests) { "transform/first_index_offset_test.cc", "transform/for_loop_to_loop_test.cc", "transform/localize_struct_array_assignment_test.cc", + "transform/merge_return_test.cc", "transform/module_scope_var_to_entry_point_param_test.cc", "transform/multiplanar_external_texture_test.cc", "transform/num_workgroups_from_uniform_test.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index 2399ee12ad..d4c13ca85f 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -444,6 +444,8 @@ set(TINT_LIB_SRCS transform/localize_struct_array_assignment.h transform/manager.cc transform/manager.h + transform/merge_return.cc + transform/merge_return.h transform/module_scope_var_to_entry_point_param.cc transform/module_scope_var_to_entry_point_param.h transform/multiplanar_external_texture.cc @@ -1162,6 +1164,7 @@ if(TINT_BUILD_TESTS) transform/for_loop_to_loop_test.cc transform/expand_compound_assignment.cc transform/localize_struct_array_assignment_test.cc + transform/merge_return.cc transform/module_scope_var_to_entry_point_param_test.cc transform/multiplanar_external_texture_test.cc transform/num_workgroups_from_uniform_test.cc diff --git a/src/tint/transform/merge_return.cc b/src/tint/transform/merge_return.cc new file mode 100644 index 0000000000..aec6b6d764 --- /dev/null +++ b/src/tint/transform/merge_return.cc @@ -0,0 +1,239 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/transform/merge_return.h" + +#include + +#include "src/tint/program_builder.h" +#include "src/tint/sem/statement.h" +#include "src/tint/utils/scoped_assignment.h" + +TINT_INSTANTIATE_TYPEINFO(tint::transform::MergeReturn); + +using namespace tint::number_suffixes; // NOLINT + +namespace tint::transform { + +namespace { + +/// Returns `true` if `stmt` has the behavior `behavior`. +bool HasBehavior(const Program* program, const ast::Statement* stmt, sem::Behavior behavior) { + return program->Sem().Get(stmt)->Behaviors().Contains(behavior); +} + +/// Returns `true` if `func` needs to be transformed. +bool NeedsTransform(const Program* program, const ast::Function* func) { + // Entry points and intrinsic declarations never need transforming. + if (func->IsEntryPoint() || func->body == nullptr) { + return false; + } + + // Avoid transforming functions that only have a single exit point. + // TODO(jrprice): Alternatively, use the uniformity analysis to decide which + // functions need to be transformed. + for (auto* s : func->body->statements) { + // Find the first statement that has contains the Return behavior. + if (HasBehavior(program, s, sem::Behavior::kReturn)) { + // If this statement is itself a return, it will be the only exit point, + // so no need to apply the transform to the function. + if (s->Is()) { + return false; + } else { + // Apply the transform in all other cases. + return true; + } + } + } + return false; +} + +} // namespace + +MergeReturn::MergeReturn() = default; + +MergeReturn::~MergeReturn() = default; + +bool MergeReturn::ShouldRun(const Program* program, const DataMap&) const { + for (auto* func : program->AST().Functions()) { + if (NeedsTransform(program, func)) { + return true; + } + } + return false; +} + +namespace { + +/// Internal class used to during the transform. +class State { + private: + /// The clone context. + CloneContext& ctx; + + /// The program builder. + ProgramBuilder& b; + + /// The function. + const ast::Function* function; + + /// The symbol for the return flag variable. + Symbol flag; + + /// The symbol for the return value variable. + Symbol retval; + + /// Tracks whether we are currently inside a loop or switch statement. + bool is_in_loop_or_switch = false; + + public: + /// Constructor + /// @param context the clone context + State(CloneContext& context, const ast::Function* func) + : ctx(context), b(*ctx.dst), function(func) {} + + /// Process a statement (recursively). + void ProcessStatement(const ast::Statement* stmt) { + if (stmt == nullptr || !HasBehavior(ctx.src, stmt, sem::Behavior::kReturn)) { + return; + } + + Switch( + stmt, [&](const ast::BlockStatement* block) { ProcessBlock(block); }, + [&](const ast::CaseStatement* c) { ProcessStatement(c->body); }, + [&](const ast::ForLoopStatement* f) { + TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true); + ProcessStatement(f->body); + }, + [&](const ast::IfStatement* i) { + ProcessStatement(i->body); + ProcessStatement(i->else_statement); + }, + [&](const ast::LoopStatement* l) { + TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true); + ProcessStatement(l->body); + }, + [&](const ast::ReturnStatement* r) { + utils::Vector stmts; + // Set the return flag to signal that we have hit a return. + stmts.Push(b.Assign(b.Expr(flag), true)); + if (r->value) { + // Set the return value if necessary. + stmts.Push(b.Assign(b.Expr(retval), ctx.Clone(r->value))); + } + if (is_in_loop_or_switch) { + // If we are in a loop or switch statement, break out of it. + stmts.Push(b.Break()); + } + ctx.Replace(r, b.Block(std::move(stmts))); + }, + [&](const ast::SwitchStatement* s) { + TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true); + for (auto* c : s->body) { + ProcessStatement(c); + } + }, + [&](const ast::WhileStatement* w) { + TINT_SCOPED_ASSIGNMENT(is_in_loop_or_switch, true); + ProcessStatement(w->body); + }, + [&](Default) { TINT_ICE(Transform, b.Diagnostics()) << "unhandled statement type"; }); + } + + void ProcessBlock(const ast::BlockStatement* block) { + // We will rebuild the contents of the block statement. + // We may introduce conditionals around statements that follow a statement with the + // `Return` behavior, so build a stack of statement lists that represent the new + // (potentially nested) conditional blocks. + utils::Vector, 8> new_stmts({{}}); + + // Insert variables for the return flag and return value at the top of the function. + if (block == function->body) { + flag = b.Symbols().New("tint_return_flag"); + new_stmts[0].Push(b.Decl(b.Var(flag, b.ty.bool_()))); + + if (!function->return_type->Is()) { + retval = b.Symbols().New("tint_return_value"); + new_stmts[0].Push(b.Decl(b.Var(retval, ctx.Clone(function->return_type)))); + } + } + + for (auto* s : block->statements) { + // Process the statement and add it to the current block. + ProcessStatement(s); + new_stmts.Back().Push(ctx.Clone(s)); + + // Check if the statement is or contains a return statement. + // We need to make sure any statements that follow this one do not get executed if the + // return flag has been set. + if (HasBehavior(ctx.src, s, sem::Behavior::kReturn)) { + if (is_in_loop_or_switch) { + // We're in a loop/switch, and so we would have inserted a `break`. + // If we've just come out of a loop/switch statement, we need to `break` again. + if (s->IsAnyOf()) { + // If the loop only has the 'Return' behavior, we can just unconditionally + // break. Otherwise check the return flag. + if (HasBehavior(ctx.src, s, sem::Behavior::kNext)) { + new_stmts.Back().Push( + b.If(b.Expr(flag), b.Block(utils::Vector{b.Break()}))); + } else { + new_stmts.Back().Push(b.Break()); + } + } + } else { + // Create a new list for any subsequent statements, which we will wrap in a + // conditional block. + new_stmts.Push({}); + } + } + } + + // Descend the stack of new block statements, wrapping them in conditionals. + while (new_stmts.Length() > 1) { + const ast::IfStatement* i = nullptr; + if (new_stmts.Back().Length() > 0) { + i = b.If(b.Not(b.Expr(flag)), b.Block(new_stmts.Back())); + } + new_stmts.Pop(); + if (i) { + new_stmts.Back().Push(i); + } + } + + // Insert the final return statement at the end of the function body. + if (block == function->body && retval.IsValid()) { + new_stmts[0].Push(b.Return(b.Expr(retval))); + } + + ctx.Replace(block, b.Block(new_stmts[0])); + } +}; + +} // namespace + +void MergeReturn::Run(CloneContext& ctx, const DataMap&, DataMap&) const { + for (auto* func : ctx.src->AST().Functions()) { + if (!NeedsTransform(ctx.src, func)) { + continue; + } + + State state(ctx, func); + state.ProcessStatement(func->body); + } + + ctx.Clone(); +} + +} // namespace tint::transform diff --git a/src/tint/transform/merge_return.h b/src/tint/transform/merge_return.h new file mode 100644 index 0000000000..1334a5cdc5 --- /dev/null +++ b/src/tint/transform/merge_return.h @@ -0,0 +1,47 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#ifndef SRC_TINT_TRANSFORM_MERGE_RETURN_H_ +#define SRC_TINT_TRANSFORM_MERGE_RETURN_H_ + +#include "src/tint/transform/transform.h" + +namespace tint::transform { + +/// Merge return statements into a single return at the end of the function. +class MergeReturn final : public Castable { + public: + /// Constructor + MergeReturn(); + /// Destructor + ~MergeReturn() override; + + /// @param program the program to inspect + /// @param data optional extra transform-specific input data + /// @returns true if this transform should be run for the given program + bool ShouldRun(const Program* program, const DataMap& data = {}) const override; + + protected: + /// Runs the transform using the CloneContext built for transforming a + /// program. Run() is responsible for calling Clone() on the CloneContext. + /// @param ctx the CloneContext primed with the input program and + /// ProgramBuilder + /// @param inputs optional extra transform-specific input data + /// @param outputs optional extra transform-specific output data + void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) const override; +}; + +} // namespace tint::transform + +#endif // SRC_TINT_TRANSFORM_MERGE_RETURN_H_ diff --git a/src/tint/transform/merge_return_test.cc b/src/tint/transform/merge_return_test.cc new file mode 100644 index 0000000000..fb2377dff2 --- /dev/null +++ b/src/tint/transform/merge_return_test.cc @@ -0,0 +1,862 @@ +// Copyright 2022 The Tint Authors. +// +// Licensed under the Apache License, Version 2.0 (the "License"); +// you may not use this file except in compliance with the License. +// You may obtain a copy of the License at +// +// http://www.apache.org/licenses/LICENSE-2.0 +// +// Unless required by applicable law or agreed to in writing, software +// distributed under the License is distributed on an "AS IS" BASIS, +// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +// See the License for the specific language governing permissions and +// limitations under the License. + +#include "src/tint/transform/merge_return.h" + +#include + +#include "src/tint/transform/test_helper.h" + +namespace tint::transform { +namespace { + +using MergeReturnTest = TransformTest; + +TEST_F(MergeReturnTest, ShouldRunEmptyModule) { + auto* src = R"()"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(MergeReturnTest, ShouldRunOnlyEntryPoints) { + auto* src = R"( +@fragment +fn foo() {} + +@compute @workgroup_size(1) +fn bar() {} +)"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(MergeReturnTest, ShouldRunUserFuncNoReturn) { + auto* src = R"( +fn foo() {} + +@compute @workgroup_size(1) +fn bar() { + foo(); +} +)"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(MergeReturnTest, ShouldRunUserFuncSingleReturn) { + auto* src = R"( +fn foo() -> i32 { + return 42; +} + +@compute @workgroup_size(1) +fn bar() { + foo(); +} +)"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(MergeReturnTest, ShouldRunUserFuncMultipleReturn) { + auto* src = R"( +fn foo(x : i32) -> i32 { + if (x == 0) { + return 42; + } + return 0; +} + +@compute @workgroup_size(1) +fn bar() { + foo(7); +} +)"; + + EXPECT_TRUE(ShouldRun(src)); +} + +TEST_F(MergeReturnTest, EmptyModule) { + auto* src = R"()"; + + auto* expect = src; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, SingleExitPoint) { + auto* src = R"( +fn foo(x : i32) -> i32 { + var retval = 0; + if ((x == 0)) { + retval = 42; + } else { + retval = 7; + } + return retval; +} +)"; + + auto* expect = src; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, Basic) { + auto* src = R"( +var non_uniform_global : i32; + +fn foo() { + if (non_uniform_global == 0) { + return; + } else { + return; + } +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +fn foo() { + var tint_return_flag : bool; + if ((non_uniform_global == 0)) { + { + tint_return_flag = true; + } + } else { + { + tint_return_flag = true; + } + } +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, BasicRetval) { + auto* src = R"( +var non_uniform_global : i32; + +fn foo() -> i32 { + if (non_uniform_global == 0) { + return 1; + } else { + return 2; + } +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + if ((non_uniform_global == 0)) { + { + tint_return_flag = true; + tint_return_value = 1; + } + } else { + { + tint_return_flag = true; + tint_return_value = 2; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, CodeAfterIf) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + if (non_uniform_global == 0) { + return 1; + } + bar = 42; + return 2; +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + if ((non_uniform_global == 0)) { + { + tint_return_flag = true; + tint_return_value = 1; + } + } + if (!(tint_return_flag)) { + bar = 42; + { + tint_return_flag = true; + tint_return_value = 2; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, SequenceOfConditionalReturns) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + if (non_uniform_global == 0) { + return 1; + } + bar = 42; + if (non_uniform_global == 1) { + return 2; + } + bar = 43; + if (non_uniform_global == 2) { + return 3; + } + bar = 44; + if (non_uniform_global == 3) { + return 4; + } + bar = 45; + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + if ((non_uniform_global == 0)) { + { + tint_return_flag = true; + tint_return_value = 1; + } + } + if (!(tint_return_flag)) { + bar = 42; + if ((non_uniform_global == 1)) { + { + tint_return_flag = true; + tint_return_value = 2; + } + } + if (!(tint_return_flag)) { + bar = 43; + if ((non_uniform_global == 2)) { + { + tint_return_flag = true; + tint_return_value = 3; + } + } + if (!(tint_return_flag)) { + bar = 44; + if ((non_uniform_global == 3)) { + { + tint_return_flag = true; + tint_return_value = 4; + } + } + if (!(tint_return_flag)) { + bar = 45; + { + tint_return_flag = true; + tint_return_value = 0; + } + } + } + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, NestedConditionals) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + if (non_uniform_global == 0) { + if (true) { + return 1; + } else { + return 2; + } + } else { + if (true) { + bar = 42; + } else if (false) { + return 3; + } else if (true) { + return 4; + } + return 5; + } +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + if ((non_uniform_global == 0)) { + if (true) { + { + tint_return_flag = true; + tint_return_value = 1; + } + } else { + { + tint_return_flag = true; + tint_return_value = 2; + } + } + } else { + if (true) { + bar = 42; + } else if (false) { + { + tint_return_flag = true; + tint_return_value = 3; + } + } else if (true) { + { + tint_return_flag = true; + tint_return_value = 4; + } + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 5; + } + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, Loop) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var i = 0; + loop { + bar = i; + if (non_uniform_global == 0) { + return non_uniform_global; + } + if (i >= 10) { + break; + } + } + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + var i = 0; + loop { + bar = i; + if ((non_uniform_global == 0)) { + { + tint_return_flag = true; + tint_return_value = non_uniform_global; + break; + } + } + if ((i >= 10)) { + break; + } + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 0; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, NestedLoop) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var i = 0; + loop { + loop { + bar = i; + if (non_uniform_global == 0) { + return non_uniform_global; + } + if (i >= 10) { + break; + } + } + bar = 10; + } + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + var i = 0; + loop { + loop { + bar = i; + if ((non_uniform_global == 0)) { + { + tint_return_flag = true; + tint_return_value = non_uniform_global; + break; + } + } + if ((i >= 10)) { + break; + } + } + if (tint_return_flag) { + break; + } + bar = 10; + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 0; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, NestedLoopAlwaysReturns) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var i = 0; + loop { + bar = i; + loop { + return non_uniform_global; + } + if (i == 10) { + break; + } + } + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + var i = 0; + loop { + bar = i; + loop { + { + tint_return_flag = true; + tint_return_value = non_uniform_global; + break; + } + } + break; + if ((i == 10)) { + break; + } + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 0; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, ForLoop) { + auto* src = R"( +var non_uniform_global : array; + +fn foo() -> i32 { + for (var i = 0; i < 10; i++) { + if (non_uniform_global[i] == 0) { + return i; + } + } + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : array; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + for(var i = 0; (i < 10); i++) { + if ((non_uniform_global[i] == 0)) { + { + tint_return_flag = true; + tint_return_value = i; + break; + } + } + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 0; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, WhileLoop) { + auto* src = R"( +var non_uniform_global : array; + +fn foo() -> i32 { + var i = 0; + while (i < 10) { + if (non_uniform_global[i] == 0) { + return i; + } + i++; + } + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : array; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + var i = 0; + while((i < 10)) { + if ((non_uniform_global[i] == 0)) { + { + tint_return_flag = true; + tint_return_value = i; + break; + } + } + i++; + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 0; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, Switch) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + switch (non_uniform_global) { + case 0 { + return 0; + } + case 1 { + if (false) { + return 1; + } + bar = 6; + fallthrough; + } + case 2 { + bar = 5; + } + default { + return 42; + } + } + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + switch(non_uniform_global) { + case 0: { + { + tint_return_flag = true; + tint_return_value = 0; + break; + } + } + case 1: { + if (false) { + { + tint_return_flag = true; + tint_return_value = 1; + break; + } + } + bar = 6; + fallthrough; + } + case 2: { + bar = 5; + } + default: { + { + tint_return_flag = true; + tint_return_value = 42; + break; + } + } + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 0; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(MergeReturnTest, ComplexNesting) { + auto* src = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + for (var i = 0; i < 10; i++) { + bar = i; + + switch (non_uniform_global) { + case 0 { + return 0; + } + case 1 { + if (false) { + loop { + if (i == non_uniform_global) { + return 1; + } + continuing { + bar++; + break if (i == 7); + } + } + } + bar = 0; + } + default { + return 42; + } + } + } + return 0; +} +)"; + + auto* expect = R"( +var non_uniform_global : i32; + +var bar : i32; + +fn foo() -> i32 { + var tint_return_flag : bool; + var tint_return_value : i32; + for(var i = 0; (i < 10); i++) { + bar = i; + switch(non_uniform_global) { + case 0: { + { + tint_return_flag = true; + tint_return_value = 0; + break; + } + } + case 1: { + if (false) { + loop { + if ((i == non_uniform_global)) { + { + tint_return_flag = true; + tint_return_value = 1; + break; + } + } + + continuing { + bar++; + break if (i == 7); + } + } + if (tint_return_flag) { + break; + } + } + bar = 0; + } + default: { + { + tint_return_flag = true; + tint_return_value = 42; + break; + } + } + } + if (tint_return_flag) { + break; + } + } + if (!(tint_return_flag)) { + { + tint_return_flag = true; + tint_return_value = 0; + } + } + return tint_return_value; +} +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +} // namespace +} // namespace tint::transform