From 68a8b094d21beff983e668cdbda7dfe0a6d96512 Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Fri, 26 May 2023 04:31:50 +0000 Subject: [PATCH] [ir] Add explicit Return instructions. This Cl adds a `ret` instruction into the IR. The `FunctionTerminator` block has been removed. Bug: tint:1718 Change-Id: Ie5fcdbfa8983b4c960773494b0c58793bd9ef503 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134461 Reviewed-by: James Price Commit-Queue: Dan Sinclair Kokoro: Kokoro Reviewed-by: Ben Clayton --- src/tint/BUILD.gn | 4 +- src/tint/CMakeLists.txt | 4 +- src/tint/ir/branch.cc | 11 +- src/tint/ir/branch.h | 9 +- src/tint/ir/builder.cc | 9 +- src/tint/ir/builder.h | 11 +- src/tint/ir/debug.cc | 38 +-- src/tint/ir/disassembler.cc | 22 +- src/tint/ir/from_program.cc | 20 +- src/tint/ir/from_program_binary_test.cc | 296 ++++++++---------- src/tint/ir/from_program_builtin_test.cc | 3 +- src/tint/ir/from_program_call_test.cc | 29 +- src/tint/ir/from_program_function_test.cc | 27 +- src/tint/ir/from_program_materialize_test.cc | 3 +- src/tint/ir/from_program_store_test.cc | 3 +- src/tint/ir/from_program_test.cc | 169 +++------- src/tint/ir/from_program_unary_test.cc | 36 +-- src/tint/ir/from_program_var_test.cc | 6 +- src/tint/ir/function.h | 7 - src/tint/ir/function_terminator.h | 33 -- .../ir/{function_terminator.cc => return.cc} | 15 +- src/tint/ir/return.h | 46 +++ src/tint/ir/to_program.cc | 35 +-- .../ir/transform/add_empty_entry_point.cc | 2 +- .../transform/add_empty_entry_point_test.cc | 8 +- src/tint/transform/manager_test.cc | 4 +- src/tint/writer/spirv/ir/generator_impl_ir.cc | 24 +- .../spirv/ir/generator_impl_ir_binary_test.cc | 14 +- .../ir/generator_impl_ir_function_test.cc | 30 +- .../spirv/ir/generator_impl_ir_if_test.cc | 10 +- .../spirv/ir/generator_impl_ir_var_test.cc | 14 +- 31 files changed, 404 insertions(+), 538 deletions(-) delete mode 100644 src/tint/ir/function_terminator.h rename src/tint/ir/{function_terminator.cc => return.cc} (65%) create mode 100644 src/tint/ir/return.h diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 11e16b2271..a45c6121df 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -1236,8 +1236,6 @@ if (tint_build_ir) { "ir/function.h", "ir/function_param.cc", "ir/function_param.h", - "ir/function_terminator.cc", - "ir/function_terminator.h", "ir/if.cc", "ir/if.h", "ir/instruction.cc", @@ -1248,6 +1246,8 @@ if (tint_build_ir) { "ir/loop.h", "ir/module.cc", "ir/module.h", + "ir/return.cc", + "ir/return.h", "ir/root_terminator.cc", "ir/root_terminator.h", "ir/store.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index 7dc64e46b2..a3d7719971 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -746,8 +746,6 @@ if(${TINT_BUILD_IR}) ir/function.h ir/function_param.cc ir/function_param.h - ir/function_terminator.cc - ir/function_terminator.h ir/if.cc ir/if.h ir/instruction.cc @@ -758,6 +756,8 @@ if(${TINT_BUILD_IR}) ir/loop.h ir/module.cc ir/module.h + ir/return.cc + ir/return.h ir/root_terminator.cc ir/root_terminator.h ir/store.cc diff --git a/src/tint/ir/branch.cc b/src/tint/ir/branch.cc index 3648b61e6d..0918962b5c 100644 --- a/src/tint/ir/branch.cc +++ b/src/tint/ir/branch.cc @@ -22,14 +22,19 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Branch); namespace tint::ir { -Branch::Branch(Block* to, utils::VectorRef args) : to_(to), args_(std::move(args)) { - TINT_ASSERT(IR, to_); - to_->AddInboundBranch(this); +Branch::Branch(utils::VectorRef args) : args_(std::move(args)) { for (auto* arg : args) { arg->AddUsage(this); } } +Branch::Branch(Block* to, utils::VectorRef args) : Branch(args) { + to_ = to; + + TINT_ASSERT(IR, to_); + to_->AddInboundBranch(this); +} + Branch::~Branch() = default; } // namespace tint::ir diff --git a/src/tint/ir/branch.h b/src/tint/ir/branch.h index fe08b979b7..aa3d1a0b4a 100644 --- a/src/tint/ir/branch.h +++ b/src/tint/ir/branch.h @@ -26,7 +26,7 @@ class Block; namespace tint::ir { -/// A branch instruction. A branch is a walk terminating jump. +/// A branch instruction. class Branch : public utils::Castable { public: /// Constructor @@ -41,8 +41,13 @@ class Branch : public utils::Castable { /// @returns the branch arguments utils::VectorRef Args() const { return args_; } + protected: + /// Constructor + /// @param args the branch arguments + explicit Branch(utils::VectorRef args); + private: - Block* to_; + Block* to_ = nullptr; utils::Vector args_; }; diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc index d64cd07bcd..28289a7ae8 100644 --- a/src/tint/ir/builder.cc +++ b/src/tint/ir/builder.cc @@ -41,10 +41,6 @@ RootTerminator* Builder::CreateRootTerminator() { return ir.blocks.Create(); } -FunctionTerminator* Builder::CreateFunctionTerminator() { - return ir.blocks.Create(); -} - Function* Builder::CreateFunction(std::string_view name, const type::Type* return_type, Function::PipelineStage stage, @@ -53,7 +49,6 @@ Function* Builder::CreateFunction(std::string_view name, auto* ir_func = ir.values.Create(return_type, stage, wg_size); ir_func->SetStartTarget(CreateBlock()); - ir_func->SetEndTarget(CreateFunctionTerminator()); ir.SetName(ir_func, name); return ir_func; } @@ -214,6 +209,10 @@ ir::Branch* Builder::Branch(Block* to, utils::VectorRef args) { return ir.values.Create(to, args); } +ir::Return* Builder::Return(Function* func, utils::VectorRef args) { + return ir.values.Create(func, args); +} + ir::BlockParam* Builder::BlockParam(const type::Type* type) { return ir.values.Create(type); } diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h index 1c4f3b2181..8c1818cadb 100644 --- a/src/tint/ir/builder.h +++ b/src/tint/ir/builder.h @@ -28,11 +28,11 @@ #include "src/tint/ir/discard.h" #include "src/tint/ir/function.h" #include "src/tint/ir/function_param.h" -#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/load.h" #include "src/tint/ir/loop.h" #include "src/tint/ir/module.h" +#include "src/tint/ir/return.h" #include "src/tint/ir/root_terminator.h" #include "src/tint/ir/store.h" #include "src/tint/ir/switch.h" @@ -65,9 +65,6 @@ class Builder { /// @returns a new root terminator flow node RootTerminator* CreateRootTerminator(); - /// @returns a new function terminator flow node - FunctionTerminator* CreateFunctionTerminator(); - /// Creates a function flow node /// @param name the function name /// @param return_type the function return type @@ -333,6 +330,12 @@ class Builder { /// @returns the instruction ir::Var* Declare(const type::Type* type); + /// Creates a return instruction + /// @param func the function being returned + /// @param args the return arguments + /// @returns the instruction + ir::Return* Return(Function* func, utils::VectorRef args = {}); + /// Creates a branch declaration /// @param to the node being branched too /// @param args the branch arguments diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc index ba674ab7ae..b9f6b608b3 100644 --- a/src/tint/ir/debug.cc +++ b/src/tint/ir/debug.cc @@ -18,9 +18,9 @@ #include #include "src/tint/ir/block.h" -#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/loop.h" +#include "src/tint/ir/return.h" #include "src/tint/ir/switch.h" #include "src/tint/switch.h" #include "src/tint/utils/string_stream.h" @@ -54,25 +54,26 @@ std::string Debug::AsDotGraph(const Module* mod) { } visited.insert(blk); - tint::Switch( - blk, - [&](const ir::FunctionTerminator*) { - // Already done - }, - [&](const ir::Block* b) { - if (block_to_name.count(b) == 0) { - out << name_for(b) << R"( [label="block"])" << std::endl; - } - out << name_for(b) << " -> " << name_for(b->Branch()->To()); + tint::Switch(blk, // + [&](const ir::Block* b) { + if (block_to_name.count(b) == 0) { + out << name_for(b) << R"( [label="block"])" << std::endl; + } + out << name_for(b) << " -> " << name_for(b->Branch()->To()); - // Dashed lines to merge blocks - if (merge_blocks.count(b->Branch()->To()) != 0) { - out << " [style=dashed]"; - } + // Dashed lines to merge blocks + if (merge_blocks.count(b->Branch()->To()) != 0) { + out << " [style=dashed]"; + } - out << std::endl; - Graph(b->Branch()->To()); - }); + out << std::endl; + + if (b->Branch()->Is()) { + return; + } else { + Graph(b->Branch()->To()); + } + }); }; out << "digraph G {" << std::endl; @@ -81,7 +82,6 @@ std::string Debug::AsDotGraph(const Module* mod) { out << "subgraph cluster_" << mod->NameOf(func).Name() << " {" << std::endl; out << R"(label=")" << mod->NameOf(func).Name() << R"(")" << std::endl; out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl; - out << name_for(func->EndTarget()) << R"( [label="end"])" << std::endl; Graph(func->StartTarget()); out << "}" << std::endl; } diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc index 03e21a2711..1525a046d8 100644 --- a/src/tint/ir/disassembler.cc +++ b/src/tint/ir/disassembler.cc @@ -25,10 +25,10 @@ #include "src/tint/ir/construct.h" #include "src/tint/ir/convert.h" #include "src/tint/ir/discard.h" -#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/load.h" #include "src/tint/ir/loop.h" +#include "src/tint/ir/return.h" #include "src/tint/ir/root_terminator.h" #include "src/tint/ir/store.h" #include "src/tint/ir/switch.h" @@ -108,11 +108,6 @@ void Disassembler::Walk(const Block* blk) { tint::Switch( blk, - [&](const ir::FunctionTerminator* t) { - TINT_ASSERT(IR, in_function_); - Indent() << "%b" << IdOf(t) << " = func_terminator" << std::endl; - in_function_ = false; - }, [&](const ir::RootTerminator* t) { TINT_ASSERT(IR, !in_function_); Indent() << "%b" << IdOf(t) << " = root_terminator" << std::endl << std::endl; @@ -142,7 +137,7 @@ void Disassembler::Walk(const Block* blk) { } Indent() << "}" << std::endl; - if (!b->Branch()->To()->Is()) { + if (!b->Branch()->Is()) { out_ << std::endl; } }); @@ -186,7 +181,6 @@ void Disassembler::EmitFunction(const Function* func) { { ScopedIndent si(indent_size_); Walk(func->StartTarget()); - Walk(func->EndTarget()); } Indent() << "}" << std::endl; } @@ -415,11 +409,13 @@ void Disassembler::EmitSwitch(const Switch* s) { void Disassembler::EmitBranch(const Branch* b) { std::string suffix = ""; - out_ << "br %b" << IdOf(b->To()); - if (b->To()->Is()) { - suffix = "return"; - } else if (b->To()->Is()) { - suffix = "root_end"; + if (b->Is()) { + out_ << "ret"; + } else { + out_ << "br %b" << IdOf(b->To()); + if (b->To()->Is()) { + suffix = "root_end"; + } } if (!b->Args().IsEmpty()) { diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc index fa00b6bf80..fde6fca472 100644 --- a/src/tint/ir/from_program.cc +++ b/src/tint/ir/from_program.cc @@ -166,6 +166,21 @@ class Impl { diagnostics_.add_error(tint::diag::System::IR, err, s); } + void SetBranch(Branch* br) { + TINT_ASSERT(IR, current_flow_block_); + TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget()); + + current_flow_block_->Instructions().Push(br); + current_flow_block_ = nullptr; + } + + void SetBranchIfNeeded(Branch* br) { + if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) { + return; + } + SetBranch(br); + } + void BranchTo(Block* node, utils::VectorRef args = {}) { TINT_ASSERT(IR, current_flow_block_); TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget()); @@ -337,7 +352,7 @@ class Impl { // If the branch target has already been set then a `return` was called. Only set in // the case where `return` wasn't called. - BranchToIfNeeded(current_function_->EndTarget()); + SetBranchIfNeeded(builder_.Return(current_function_)); } TINT_ASSERT(IR, control_stack_.IsEmpty()); @@ -716,8 +731,7 @@ class Impl { } ret_value.Push(ret.Get()); } - - BranchTo(current_function_->EndTarget(), std::move(ret_value)); + SetBranch(builder_.Return(current_function_, std::move(ret_value))); } void EmitBreak(const ast::BreakStatement*) { diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc index cdd62274c4..1a81f56598 100644 --- a/src/tint/ir/from_program_binary_test.cc +++ b/src/tint/ir/from_program_binary_test.cc @@ -36,17 +36,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Add) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = add %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -72,9 +70,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) { %3:u32 = load %v1 %4:u32 = add %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -100,9 +97,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) { %3:u32 = load %v1 %4:u32 = add %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -117,17 +113,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Subtract) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = sub %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -153,9 +147,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) { %3:i32 = load %v1 %4:i32 = sub %3, 1i store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -181,9 +174,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) { %3:u32 = load %v1 %4:u32 = sub %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -198,17 +190,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Multiply) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = mul %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -234,9 +224,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) { %3:u32 = load %v1 %4:u32 = mul %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -251,17 +240,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Div) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = div %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -287,9 +274,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) { %3:u32 = load %v1 %4:u32 = div %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -304,17 +290,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Modulo) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = mod %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -340,9 +324,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) { %3:u32 = load %v1 %4:u32 = mod %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -357,17 +340,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_And) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = and %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -393,9 +374,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) { %3:bool = load %v1 %4:bool = and %3, false store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -410,17 +390,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Or) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = or %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -446,9 +424,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) { %3:bool = load %v1 %4:bool = or %3, false store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -463,17 +440,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Xor) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = xor %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -499,9 +474,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) { %3:u32 = load %v1 %4:u32 = xor %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -516,40 +490,39 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %b1 { %b1 = block { - br %b2 true # return + ret true } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:bool = call %my_func - if %3 [t: %b4, f: %b5, m: %b6] + if %3 [t: %b3, f: %b4, m: %b5] # True block - %b4 = block { - br %b6 false + %b3 = block { + br %b5 false } # False block - %b5 = block { - br %b6 %3 + %b4 = block { + br %b5 %3 } # Merge block - %b6 = block (%4:bool) { - if %4:bool [t: %b7, f: %b8, m: %b9] + %b5 = block (%4:bool) { + if %4:bool [t: %b6, f: %b7, m: %b8] # True block - %b7 = block { - br %b9 + %b6 = block { + br %b8 } # False block - %b8 = block { - br %b9 + %b7 = block { + br %b8 } # Merge block - %b9 = block { - br %b10 # return + %b8 = block { + ret } } @@ -557,7 +530,6 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) { } - %b10 = func_terminator } )"); } @@ -572,40 +544,39 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %b1 { %b1 = block { - br %b2 true # return + ret true } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:bool = call %my_func - if %3 [t: %b4, f: %b5, m: %b6] + if %3 [t: %b3, f: %b4, m: %b5] # True block - %b4 = block { - br %b6 %3 + %b3 = block { + br %b5 %3 } # False block - %b5 = block { - br %b6 true + %b4 = block { + br %b5 true } # Merge block - %b6 = block (%4:bool) { - if %4:bool [t: %b7, f: %b8, m: %b9] + %b5 = block (%4:bool) { + if %4:bool [t: %b6, f: %b7, m: %b8] # True block - %b7 = block { - br %b9 + %b6 = block { + br %b8 } # False block - %b8 = block { - br %b9 + %b7 = block { + br %b8 } # Merge block - %b9 = block { - br %b10 # return + %b8 = block { + ret } } @@ -613,7 +584,6 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) { } - %b10 = func_terminator } )"); } @@ -628,17 +598,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Equal) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:bool = eq %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -653,17 +621,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_NotEqual) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:bool = neq %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -678,17 +644,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThan) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:bool = lt %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -703,17 +667,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThan) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:bool = gt %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -728,17 +690,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThanEqual) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:bool = lte %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -753,17 +713,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThanEqual) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:bool = gte %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -778,17 +736,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftLeft) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = shiftl %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -814,9 +770,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) { %3:u32 = load %v1 %4:u32 = shiftl %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -831,17 +786,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftRight) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 0u # return + ret 0u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = shiftr %3, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -867,9 +820,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) { %3:u32 = load %v1 %4:u32 = shiftr %3, 1u store %v1, %4 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -886,38 +838,36 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 -> %b1 { %b1 = block { - br %b2 0.0f # return + ret 0.0f } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:f32 = call %my_func %4:bool = lt %3, 2.0f - if %4 [t: %b4, f: %b5, m: %b6] + if %4 [t: %b3, f: %b4, m: %b5] # True block - %b4 = block { + %b3 = block { %5:f32 = call %my_func %6:f32 = call %my_func %7:f32 = mul 2.29999995231628417969f, %6 %8:f32 = div %5, %7 %9:bool = gt 2.5f, %8 - br %b6 %9 + br %b5 %9 } # False block - %b5 = block { - br %b6 %4 + %b4 = block { + br %b5 %4 } # Merge block - %b6 = block (%tint_symbol:bool) { - br %b7 # return + %b5 = block (%tint_symbol:bool) { + ret } } - %b7 = func_terminator } )"); } @@ -933,16 +883,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:bool):bool -> %b1 { %b1 = block { - br %b2 true # return + ret true } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %tint_symbol:bool = call %my_func, false - br %b4 # return + ret } - %b4 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_builtin_test.cc b/src/tint/ir/from_program_builtin_test.cc index 37505ef3cb..3993fbbbf8 100644 --- a/src/tint/ir/from_program_builtin_test.cc +++ b/src/tint/ir/from_program_builtin_test.cc @@ -46,9 +46,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) { %b3 = block { %3:f32 = load %i %tint_symbol:f32 = asin %3 - br %b4 # return + ret } - %b4 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc index efd90ccf60..43a9709be4 100644 --- a/src/tint/ir/from_program_call_test.cc +++ b/src/tint/ir/from_program_call_test.cc @@ -37,17 +37,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 -> %b1 { %b1 = block { - br %b2 0.0f # return + ret 0.0f } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:f32 = call %my_func %tint_symbol:f32 = bitcast %3 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -65,9 +63,8 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Discard) { EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():void [@fragment] -> %b1 { %b1 = block { discard - br %b2 # return + ret } - %b2 = func_terminator } )"); } @@ -82,16 +79,14 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:f32):void -> %b1 { %b1 = block { - br %b2 # return + ret } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %4:void = call %my_func, 6.0f - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -116,9 +111,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) { %b3 = block { %3:i32 = load %i %tint_symbol:f32 = convert i32, %3 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -161,9 +155,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) { %b3 = block { %3:f32 = load %i %tint_symbol:vec3 = construct 2.0f, 3.0f, %3 - br %b4 # return + ret } - %b4 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_function_test.cc b/src/tint/ir/from_program_function_test.cc index d5ae1068f3..1be388a42f 100644 --- a/src/tint/ir/from_program_function_test.cc +++ b/src/tint/ir/from_program_function_test.cc @@ -36,9 +36,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4 [@vertex ra: @position] -> %b1 { %b1 = block { - br %b2 vec4 0.0f # return + ret vec4 0.0f } - %b2 = func_terminator } )"); } @@ -52,9 +51,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Fragment) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():void [@fragment] -> %b1 { %b1 = block { - br %b2 # return + ret } - %b2 = func_terminator } )"); } @@ -69,9 +67,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Compute) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():void [@compute @workgroup_size(8, 4, 2)] -> %b1 { %b1 = block { - br %b2 # return + ret } - %b2 = func_terminator } )"); } @@ -85,9 +82,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Return) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec3 -> %b1 { %b1 = block { - br %b2 vec3 0.0f # return + ret vec3 0.0f } - %b2 = func_terminator } )"); } @@ -102,9 +98,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4 [@vertex ra: @position] -> %b1 { %b1 = block { - br %b2 vec4 1.0f, 2.0f, 3.0f, 4.0f # return + ret vec4 1.0f, 2.0f, 3.0f, 4.0f } - %b2 = func_terminator } )"); } @@ -120,9 +115,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4 [@vertex ra: @position @invariant] -> %b1 { %b1 = block { - br %b2 vec4 1.0f, 2.0f, 3.0f, 4.0f # return + ret vec4 1.0f, 2.0f, 3.0f, 4.0f } - %b2 = func_terminator } )"); } @@ -137,9 +131,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():vec4 [@fragment ra: @location(1)] -> %b1 { %b1 = block { - br %b2 vec4 1.0f, 2.0f, 3.0f, 4.0f # return + ret vec4 1.0f, 2.0f, 3.0f, 4.0f } - %b2 = func_terminator } )"); } @@ -154,9 +147,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnFragDepth) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():f32 [@fragment ra: @frag_depth] -> %b1 { %b1 = block { - br %b2 1.0f # return + ret 1.0f } - %b2 = func_terminator } )"); } @@ -171,9 +163,8 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnSampleMask) { EXPECT_EQ(Disassemble(m.Get()), R"(%test = func():u32 [@fragment ra: @sample_mask] -> %b1 { %b1 = block { - br %b2 1u # return + ret 1u } - %b2 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_materialize_test.cc b/src/tint/ir/from_program_materialize_test.cc index 1fdc5ab564..126257978f 100644 --- a/src/tint/ir/from_program_materialize_test.cc +++ b/src/tint/ir/from_program_materialize_test.cc @@ -36,9 +36,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) { EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():f32 -> %b1 { %b1 = block { - br %b2 2.0f # return + ret 2.0f } - %b2 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc index bebd300d92..2d3f3d730c 100644 --- a/src/tint/ir/from_program_store_test.cc +++ b/src/tint/ir/from_program_store_test.cc @@ -46,9 +46,8 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { %b3 = block { store %a, 4u - br %b4 # return + ret } - %b4 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc index 89fa102d0b..406bde588f 100644 --- a/src/tint/ir/from_program_test.cc +++ b/src/tint/ir/from_program_test.cc @@ -19,7 +19,6 @@ #include "src/tint/ast/int_literal_expression.h" #include "src/tint/constant/scalar.h" #include "src/tint/ir/block.h" -#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/loop.h" #include "src/tint/ir/switch.h" @@ -64,17 +63,13 @@ TEST_F(IR_BuilderImplTest, Func) { auto* f = m->functions[0]; ASSERT_NE(f->StartTarget(), nullptr); - ASSERT_NE(f->EndTarget(), nullptr); - - EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length()); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():void -> %b1 { %b1 = block { - br %b2 # return + ret } - %b2 = func_terminator } )"); } @@ -89,17 +84,13 @@ TEST_F(IR_BuilderImplTest, Func_WithParam) { auto* f = m->functions[0]; ASSERT_NE(f->StartTarget(), nullptr); - ASSERT_NE(f->EndTarget(), nullptr); - - EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length()); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32):u32 -> %b1 { %b1 = block { - br %b2 %a # return + ret %a } - %b2 = func_terminator } )"); } @@ -115,17 +106,13 @@ TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) { auto* f = m->functions[0]; ASSERT_NE(f->StartTarget(), nullptr); - ASSERT_NE(f->EndTarget(), nullptr); - - EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length()); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32, %b:i32, %c:bool):void -> %b1 { %b1 = block { - br %b2 # return + ret } - %b2 = func_terminator } )"); } @@ -151,12 +138,10 @@ TEST_F(IR_BuilderImplTest, IfStatement) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -174,12 +159,11 @@ TEST_F(IR_BuilderImplTest, IfStatement) { # Merge block %b4 = block { - br %b5 # return + ret } } - %b5 = func_terminator } )"); } @@ -195,12 +179,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -208,7 +190,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) { if true [t: %b2, f: %b3, m: %b4] # True block %b2 = block { - br %b5 # return + ret } # False block %b3 = block { @@ -217,12 +199,11 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) { # Merge block %b4 = block { - br %b5 # return + ret } } - %b5 = func_terminator } )"); } @@ -238,12 +219,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -256,16 +235,15 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) { # False block %b3 = block { - br %b5 # return + ret } # Merge block %b4 = block { - br %b5 # return + ret } } - %b5 = func_terminator } )"); } @@ -281,12 +259,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -294,16 +270,15 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) { if true [t: %b2, f: %b3] # True block %b2 = block { - br %b4 # return + ret } # False block %b3 = block { - br %b4 # return + ret } } - %b4 = func_terminator } )"); } @@ -354,12 +329,11 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) { # Merge block %b4 = block { - br %b8 # return + ret } } - %b8 = func_terminator } )"); } @@ -375,12 +349,10 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -397,12 +369,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) { # Merge block %b4 = block { - br %b5 # return + ret } } - %b5 = func_terminator } )"); } @@ -421,7 +392,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) { auto* if_flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); @@ -429,7 +399,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) { EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -462,12 +431,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) { # Merge block %b4 = block { - br %b8 # return + ret } } - %b8 = func_terminator } )"); } @@ -485,7 +453,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) { auto* break_if_flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); @@ -493,7 +460,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) { EXPECT_EQ(1u, break_if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, break_if_flow->False()->InboundBranches().Length()); EXPECT_EQ(2u, break_if_flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -526,12 +492,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) { # Merge block %b4 = block { - br %b8 # return + ret } } - %b8 = func_terminator } )"); } @@ -577,12 +542,11 @@ TEST_F(IR_BuilderImplTest, Loop_Continuing_Body_Scope) { # Merge block %b4 = block { - br %b8 # return + ret } } - %b8 = func_terminator } )"); } @@ -600,7 +564,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { auto* if_flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); @@ -608,7 +571,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -618,7 +580,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { if true [t: %b4, f: %b5, m: %b6] # True block %b4 = block { - br %b7 # return + ret } # False block %b5 = block { @@ -641,7 +603,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { } - %b7 = func_terminator } )"); } @@ -657,19 +618,17 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) { auto* loop_flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { %b1 = block { loop [s: %b2, c: %b3] %b2 = block { - br %b4 # return + ret } # Continuing block %b3 = block { @@ -679,7 +638,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) { } - %b4 = func_terminator } )"); } @@ -704,35 +662,33 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) { auto* loop_flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(3u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { %b1 = block { loop [s: %b2, c: %b3, m: %b4] %b2 = block { - br %b5 # return + ret } # Continuing block %b3 = block { - if true [t: %b6, f: %b7, m: %b8] + if true [t: %b5, f: %b6, m: %b7] # True block - %b6 = block { + %b5 = block { br %b4 } # False block - %b7 = block { - br %b8 + %b6 = block { + br %b7 } # Merge block - %b8 = block { + %b7 = block { br %b2 } @@ -741,19 +697,19 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) { # Merge block %b4 = block { - if true [t: %b9, f: %b10, m: %b11] + if true [t: %b8, f: %b9, m: %b10] # True block - %b9 = block { - br %b5 # return + %b8 = block { + ret } # False block - %b10 = block { - br %b11 + %b9 = block { + br %b10 } # Merge block - %b11 = block { - br %b5 # return + %b10 = block { + ret } } @@ -761,7 +717,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) { } - %b5 = func_terminator } )"); } @@ -779,7 +734,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) { auto* if_flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); @@ -787,7 +741,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) { EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -815,12 +768,11 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) { # Merge block %b4 = block { - br %b7 # return + ret } } - %b7 = func_terminator } )"); } @@ -966,12 +918,11 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) { # Merge block %b4 = block { - br %b26 # return + ret } } - %b26 = func_terminator } )"); } @@ -991,9 +942,7 @@ TEST_F(IR_BuilderImplTest, While) { auto* if_flow = flow->Start()->Branch()->As(); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); @@ -1032,12 +981,11 @@ TEST_F(IR_BuilderImplTest, While) { # Merge block %b4 = block { - br %b8 # return + ret } } - %b8 = func_terminator } )"); } @@ -1057,9 +1005,7 @@ TEST_F(IR_BuilderImplTest, While_Return) { auto* if_flow = flow->Start()->Branch()->As(); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; - EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); @@ -1085,7 +1031,7 @@ TEST_F(IR_BuilderImplTest, While_Return) { # Merge block %b7 = block { - br %b8 # return + ret } } @@ -1097,12 +1043,11 @@ TEST_F(IR_BuilderImplTest, While_Return) { # Merge block %b4 = block { - br %b8 # return + ret } } - %b8 = func_terminator } )"); } @@ -1135,9 +1080,7 @@ TEST_F(IR_BuilderImplTest, DISABLED_For) { auto* if_flow = flow->Start()->Branch()->As(); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); @@ -1159,12 +1102,10 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -1181,12 +1122,11 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) { # Merge block %b4 = block { - br %b5 # return + ret } } - %b5 = func_terminator } )"); } @@ -1205,7 +1145,6 @@ TEST_F(IR_BuilderImplTest, Switch) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(3u, cases.Length()); @@ -1227,7 +1166,6 @@ TEST_F(IR_BuilderImplTest, Switch) { EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[2].Start()->InboundBranches().Length()); EXPECT_EQ(4u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -1250,12 +1188,11 @@ TEST_F(IR_BuilderImplTest, Switch) { # Merge block %b5 = block { - br %b6 # return + ret } } - %b6 = func_terminator } )"); } @@ -1275,7 +1212,6 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(1u, cases.Length()); @@ -1292,7 +1228,6 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) { EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -1305,12 +1240,11 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) { # Merge block %b3 = block { - br %b4 # return + ret } } - %b4 = func_terminator } )"); } @@ -1326,7 +1260,6 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(1u, cases.Length()); @@ -1335,7 +1268,6 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) { EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -1348,12 +1280,11 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) { # Merge block %b3 = block { - br %b4 # return + ret } } - %b4 = func_terminator } )"); } @@ -1371,7 +1302,6 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(2u, cases.Length()); @@ -1387,7 +1317,6 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) { EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length()); // This is 1 because the if is dead-code eliminated and the return doesn't happen. - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -1405,12 +1334,11 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) { # Merge block %b4 = block { - br %b5 # return + ret } } - %b5 = func_terminator } )"); } @@ -1431,7 +1359,6 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) { auto* flow = FindSingleValue(m); ASSERT_EQ(1u, m.functions.Length()); - auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(2u, cases.Length()); @@ -1446,7 +1373,6 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) { EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); - EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m), R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { @@ -1454,16 +1380,15 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) { switch 1i [c: (0i, %b2), c: (default, %b3)] # Case block %b2 = block { - br %b4 # return + ret } # Case block %b3 = block { - br %b4 # return + ret } } - %b4 = func_terminator } )"); } @@ -1478,16 +1403,14 @@ TEST_F(IR_BuilderImplTest, Emit_Phony) { EXPECT_EQ(Disassemble(m.Get()), R"(%b = func():i32 -> %b1 { %b1 = block { - br %b2 1i # return + ret 1i } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:i32 = call %b - br %b4 # return + ret } - %b4 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc index bbdb773f0f..774be8cb51 100644 --- a/src/tint/ir/from_program_unary_test.cc +++ b/src/tint/ir/from_program_unary_test.cc @@ -36,17 +36,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Not) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %b1 { %b1 = block { - br %b2 false # return + ret false } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:bool = call %my_func %tint_symbol:bool = eq %3, false - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -61,17 +59,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Complement) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %b1 { %b1 = block { - br %b2 1u # return + ret 1u } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:u32 = call %my_func %tint_symbol:u32 = complement %3 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -86,17 +82,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Negation) { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():i32 -> %b1 { %b1 = block { - br %b2 1i # return + ret 1i } - %b2 = func_terminator } -%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { - %b3 = block { +%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 { + %b2 = block { %3:i32 = call %my_func %tint_symbol:i32 = negation %3 - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -120,9 +114,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { %b3 = block { - br %b4 # return + ret } - %b4 = func_terminator } )"); } @@ -149,9 +142,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 { %b3 = block { store %v3, 42i - br %b4 # return + ret } - %b4 = func_terminator } )"); } diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc index 1f793ef931..29377e4f33 100644 --- a/src/tint/ir/from_program_var_test.cc +++ b/src/tint/ir/from_program_var_test.cc @@ -72,9 +72,8 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { %b1 = block { %a:ptr = var - br %b2 # return + ret } - %b2 = func_terminator } )"); } @@ -91,9 +90,8 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { %b1 = block { %a:ptr = var, 2u - br %b2 # return + ret } - %b2 = func_terminator } )"); } diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h index f0ab80375d..172afdc554 100644 --- a/src/tint/ir/function.h +++ b/src/tint/ir/function.h @@ -117,12 +117,6 @@ class Function : public utils::Castable { /// @returns the function start target Block* StartTarget() const { return start_target_; } - /// Sets the end target for the function - /// @param target the end target - void SetEndTarget(FunctionTerminator* target) { end_target_ = target; } - /// @returns the function end target - FunctionTerminator* EndTarget() const { return end_target_; } - private: const type::Type* return_type_; PipelineStage pipeline_stage_; @@ -134,7 +128,6 @@ class Function : public utils::Castable { utils::Vector params_; Block* start_target_ = nullptr; - FunctionTerminator* end_target_ = nullptr; }; utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value); diff --git a/src/tint/ir/function_terminator.h b/src/tint/ir/function_terminator.h deleted file mode 100644 index 42aa01e437..0000000000 --- a/src/tint/ir/function_terminator.h +++ /dev/null @@ -1,33 +0,0 @@ -// 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_IR_FUNCTION_TERMINATOR_H_ -#define SRC_TINT_IR_FUNCTION_TERMINATOR_H_ - -#include "src/tint/ir/block.h" - -namespace tint::ir { - -/// Block used as the end of a function. Must only be used as the `end_target` in a function. There -/// are no instructions in this block. -class FunctionTerminator : public utils::Castable { - public: - /// Constructor - FunctionTerminator(); - ~FunctionTerminator() override; -}; - -} // namespace tint::ir - -#endif // SRC_TINT_IR_FUNCTION_TERMINATOR_H_ diff --git a/src/tint/ir/function_terminator.cc b/src/tint/ir/return.cc similarity index 65% rename from src/tint/ir/function_terminator.cc rename to src/tint/ir/return.cc index 59893117ee..bf7fca7207 100644 --- a/src/tint/ir/function_terminator.cc +++ b/src/tint/ir/return.cc @@ -1,4 +1,4 @@ -// Copyright 2022 The Tint Authors. +// Copyright 2023 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. @@ -12,14 +12,19 @@ // See the License for the specific language governing permissions and // limitations under the License. -#include "src/tint/ir/function_terminator.h" +#include "src/tint/ir/return.h" -TINT_INSTANTIATE_TYPEINFO(tint::ir::FunctionTerminator); +#include "src/tint/ir/function.h" + +TINT_INSTANTIATE_TYPEINFO(tint::ir::Return); namespace tint::ir { -FunctionTerminator::FunctionTerminator() : Base() {} +Return::Return(Function* func, utils::VectorRef args) : Base(args), func_(func) { + TINT_ASSERT(IR, func_); + func_->AddUsage(this); +} -FunctionTerminator::~FunctionTerminator() = default; +Return::~Return() = default; } // namespace tint::ir diff --git a/src/tint/ir/return.h b/src/tint/ir/return.h new file mode 100644 index 0000000000..4e4bfba7bb --- /dev/null +++ b/src/tint/ir/return.h @@ -0,0 +1,46 @@ +// Copyright 2023 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_IR_RETURN_H_ +#define SRC_TINT_IR_RETURN_H_ + +#include "src/tint/ir/branch.h" +#include "src/tint/utils/castable.h" + +// Forward declarations +namespace tint::ir { +class Function; +} // namespace tint::ir + +namespace tint::ir { + +/// A return instruction. +class Return : public utils::Castable { + public: + /// Constructor + /// @param func the function being returned + /// @param args the branch arguments + explicit Return(Function* func, utils::VectorRef args = {}); + ~Return() override; + + /// @returns the function being returned + const Function* Func() const { return func_; } + + private: + Function* func_ = nullptr; +}; + +} // namespace tint::ir + +#endif // SRC_TINT_IR_RETURN_H_ diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc index c9c4f3a2cd..af4a957a8e 100644 --- a/src/tint/ir/to_program.cc +++ b/src/tint/ir/to_program.cc @@ -20,11 +20,11 @@ #include "src/tint/ir/block.h" #include "src/tint/ir/call.h" #include "src/tint/ir/constant.h" -#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/instruction.h" #include "src/tint/ir/load.h" #include "src/tint/ir/module.h" +#include "src/tint/ir/return.h" #include "src/tint/ir/store.h" #include "src/tint/ir/switch.h" #include "src/tint/ir/user_call.h" @@ -126,8 +126,6 @@ class State { Status status = tint::Switch( block, - [&](const ir::FunctionTerminator*) { return kStop; }, - [&](const ir::Block* blk) { for (auto* inst : blk->Instructions()) { auto stmt = Stmt(inst); @@ -240,9 +238,9 @@ class State { return b.Switch(cond, std::move(cases)); } - utils::Result FunctionTerminator(const ir::Branch* branch) { - if (branch->Args().IsEmpty()) { - // Branch to function terminator has no arguments. + utils::Result Return(const ir::Return* ret) { + if (ret->Args().IsEmpty()) { + // Return has no arguments. // If this block is nested withing some control flow, then we must // emit a 'return' statement, otherwise we've just naturally reached // the end of the function where the 'return' is redundant. @@ -252,16 +250,14 @@ class State { return nullptr; } - // Branch to function terminator has arguments - this is the return - // value. - if (branch->Args().Length() != 1) { - TINT_ICE(IR, b.Diagnostics()) << "expected 1 value for function " - "terminator (return value), got " - << branch->Args().Length(); + // Return has arguments - this is the return value. + if (ret->Args().Length() != 1) { + TINT_ICE(IR, b.Diagnostics()) + << "expected 1 value for return, got " << ret->Args().Length(); return utils::Failure; } - auto* val = Expr(branch->Args().Front()); + auto* val = Expr(ret->Args().Front()); if (TINT_UNLIKELY(!val)) { return utils::Failure; } @@ -275,10 +271,8 @@ class State { return true; } if (auto* br = node->Instructions().Front()->As()) { - return br->To() == stop_at; + return !br->Is() && br->To() == stop_at; } - // TODO(dsinclair): This should possibly walk over Jump instructions that - // just jump to empty blocks if we want to be comprehensive. return false; } @@ -291,12 +285,9 @@ class State { [&](const ir::Store* i) { return Store(i); }, // [&](const ir::If* if_) { return If(if_); }, [&](const ir::Switch* switch_) { return Switch(switch_); }, - [&](const ir::Branch* branch) { - if (branch->To()->Is()) { - return utils::Result{FunctionTerminator(branch)}; - } - return utils::Result{nullptr}; - }, + [&](const ir::Return* ret) { return Return(ret); }, + // TODO(dsinclair): Remove when branch is only a parent ... + [&](const ir::Branch*) { return utils::Result{nullptr}; }, [&](Default) { UNHANDLED_CASE(inst); return utils::Failure; diff --git a/src/tint/ir/transform/add_empty_entry_point.cc b/src/tint/ir/transform/add_empty_entry_point.cc index 770ba072f7..9cb58e2a06 100644 --- a/src/tint/ir/transform/add_empty_entry_point.cc +++ b/src/tint/ir/transform/add_empty_entry_point.cc @@ -37,7 +37,7 @@ void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const { ir::Builder builder(*ir); auto* ep = builder.CreateFunction("unused_entry_point", ir->Types().void_(), Function::PipelineStage::kCompute, std::array{1u, 1u, 1u}); - ep->StartTarget()->SetInstructions(utils::Vector{builder.Branch(ep->EndTarget())}); + ep->StartTarget()->SetInstructions(utils::Vector{builder.Return(ep)}); ir->functions.Push(ep); } diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc index 72fde36473..ca918f2da5 100644 --- a/src/tint/ir/transform/add_empty_entry_point_test.cc +++ b/src/tint/ir/transform/add_empty_entry_point_test.cc @@ -27,9 +27,8 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) { auto* expect = R"( %unused_entry_point = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { %b1 = block { - br %b2 # return + ret } - %b2 = func_terminator } )"; @@ -40,15 +39,14 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) { TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) { auto* ep = b.CreateFunction("main", mod.Types().void_(), Function::PipelineStage::kFragment); - ep->StartTarget()->SetInstructions(utils::Vector{b.Branch(ep->EndTarget())}); + ep->StartTarget()->SetInstructions(utils::Vector{b.Return(ep)}); mod.functions.Push(ep); auto* expect = R"( %main = func():void [@fragment] -> %b1 { %b1 = block { - br %b2 # return + ret } - %b2 = func_terminator } )"; diff --git a/src/tint/transform/manager_test.cc b/src/tint/transform/manager_test.cc index 10a72aff95..86c5961161 100644 --- a/src/tint/transform/manager_test.cc +++ b/src/tint/transform/manager_test.cc @@ -51,7 +51,7 @@ class IR_AddFunction final : public ir::transform::Transform { void Run(ir::Module* mod, const DataMap&, DataMap&) const override { ir::Builder builder(*mod); auto* func = builder.CreateFunction("ir_func", mod->Types().Get()); - func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{builder.Return(func)}); mod->functions.Push(func); } }; @@ -68,7 +68,7 @@ ir::Module MakeIR() { ir::Module mod; ir::Builder builder(mod); auto* func = builder.CreateFunction("main", mod.Types().Get()); - func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{builder.Return(func)}); builder.ir.functions.Push(func); return mod; } diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc index bc599263f6..29c89e6619 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc @@ -19,10 +19,10 @@ #include "spirv/unified1/spirv.h" #include "src/tint/ir/binary.h" #include "src/tint/ir/block.h" -#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/load.h" #include "src/tint/ir/module.h" +#include "src/tint/ir/return.h" #include "src/tint/ir/store.h" #include "src/tint/ir/transform/add_empty_entry_point.h" #include "src/tint/ir/user_call.h" @@ -354,18 +354,20 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) { } void GeneratorImplIr::EmitBranch(const ir::Branch* b) { + if (b->Is()) { + if (!b->Args().IsEmpty()) { + TINT_ASSERT(Writer, b->Args().Length() == 1u); + OperandList operands; + operands.push_back(Value(b->Args()[0])); + current_function_.push_inst(spv::Op::OpReturnValue, operands); + } else { + current_function_.push_inst(spv::Op::OpReturn, {}); + } + return; + } + Switch( b->To(), - [&](const ir::FunctionTerminator*) { - if (!b->Args().IsEmpty()) { - TINT_ASSERT(Writer, b->Args().Length() == 1u); - OperandList operands; - operands.push_back(Value(b->Args()[0])); - current_function_.push_inst(spv::Op::OpReturnValue, operands); - } else { - current_function_.push_inst(spv::Op::OpReturn, {}); - } - }, [&](const ir::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); }, [&](Default) { // A block may not have an outward branch (e.g. an unreachable merge diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc index 71a2a36bf2..cfe3dc4eea 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc @@ -125,7 +125,7 @@ TEST_P(Arithmetic, Scalar) { func->StartTarget()->SetInstructions( utils::Vector{b.CreateBinary(params.kind, MakeScalarType(params.type), MakeScalarValue(params.type), MakeScalarValue(params.type)), - b.Branch(func->EndTarget())}); + b.Return(func)}); generator_.EmitFunction(func); EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst)); @@ -138,7 +138,7 @@ TEST_P(Arithmetic, Vector) { utils::Vector{b.CreateBinary(params.kind, MakeVectorType(params.type), MakeVectorValue(params.type), MakeVectorValue(params.type)), - b.Branch(func->EndTarget())}); + b.Return(func)}); generator_.EmitFunction(func); EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst)); @@ -172,7 +172,7 @@ TEST_P(Bitwise, Scalar) { func->StartTarget()->SetInstructions( utils::Vector{b.CreateBinary(params.kind, MakeScalarType(params.type), MakeScalarValue(params.type), MakeScalarValue(params.type)), - b.Branch(func->EndTarget())}); + b.Return(func)}); generator_.EmitFunction(func); EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst)); @@ -185,7 +185,7 @@ TEST_P(Bitwise, Vector) { utils::Vector{b.CreateBinary(params.kind, MakeVectorType(params.type), MakeVectorValue(params.type), MakeVectorValue(params.type)), - b.Branch(func->EndTarget())}); + b.Return(func)}); generator_.EmitFunction(func); EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst)); @@ -211,7 +211,7 @@ TEST_P(Comparison, Scalar) { func->StartTarget()->SetInstructions( utils::Vector{b.CreateBinary(params.kind, mod.Types().bool_(), MakeScalarValue(params.type), MakeScalarValue(params.type)), - b.Branch(func->EndTarget())}); + b.Return(func)}); generator_.EmitFunction(func); EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst)); @@ -224,7 +224,7 @@ TEST_P(Comparison, Vector) { utils::Vector{b.CreateBinary(params.kind, mod.Types().vec2(mod.Types().bool_()), MakeVectorValue(params.type), MakeVectorValue(params.type)), - b.Branch(func->EndTarget())}); + b.Return(func)}); generator_.EmitFunction(func); EXPECT_THAT(DumpModule(generator_.Module()), ::testing::HasSubstr(params.spirv_inst)); @@ -279,7 +279,7 @@ TEST_F(SpvGeneratorImplTest, Binary_Chain) { auto* func = b.CreateFunction("foo", mod.Types().void_()); auto* a = b.Subtract(mod.Types().i32(), b.Constant(1_i), b.Constant(2_i)); func->StartTarget()->SetInstructions( - utils::Vector{a, b.Add(mod.Types().i32(), a, a), b.Branch(func->EndTarget())}); + utils::Vector{a, b.Add(mod.Types().i32(), a, a), b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc index 604835d404..6412bc65e3 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc @@ -19,7 +19,7 @@ namespace { TEST_F(SpvGeneratorImplTest, Function_Empty) { auto* func = b.CreateFunction("foo", mod.Types().void_()); - func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -35,7 +35,7 @@ OpFunctionEnd // Test that we do not emit the same function type more than once. TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) { auto* func = b.CreateFunction("foo", mod.Types().void_()); - func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)}); generator_.EmitFunction(func); generator_.EmitFunction(func); @@ -48,7 +48,7 @@ TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) { TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Compute) { auto* func = b.CreateFunction("main", mod.Types().void_(), ir::Function::PipelineStage::kCompute, {{32, 4, 1}}); - func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main" @@ -66,7 +66,7 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Fragment) { auto* func = b.CreateFunction("main", mod.Types().void_(), ir::Function::PipelineStage::kFragment); - func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Fragment %1 "main" @@ -84,7 +84,7 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Vertex) { auto* func = b.CreateFunction("main", mod.Types().void_(), ir::Function::PipelineStage::kVertex); - func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Vertex %1 "main" @@ -101,15 +101,15 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Multiple) { auto* f1 = b.CreateFunction("main1", mod.Types().void_(), ir::Function::PipelineStage::kCompute, {{32, 4, 1}}); - f1->StartTarget()->SetInstructions(utils::Vector{b.Branch(f1->EndTarget())}); + f1->StartTarget()->SetInstructions(utils::Vector{b.Return(f1)}); auto* f2 = b.CreateFunction("main2", mod.Types().void_(), ir::Function::PipelineStage::kCompute, {{8, 2, 16}}); - f2->StartTarget()->SetInstructions(utils::Vector{b.Branch(f2->EndTarget())}); + f2->StartTarget()->SetInstructions(utils::Vector{b.Return(f2)}); auto* f3 = b.CreateFunction("main3", mod.Types().void_(), ir::Function::PipelineStage::kFragment); - f3->StartTarget()->SetInstructions(utils::Vector{b.Branch(f3->EndTarget())}); + f3->StartTarget()->SetInstructions(utils::Vector{b.Return(f3)}); generator_.EmitFunction(f1); generator_.EmitFunction(f2); @@ -143,7 +143,7 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_ReturnValue) { auto* func = b.CreateFunction("foo", mod.Types().i32()); func->StartTarget()->SetInstructions( - utils::Vector{b.Branch(func->EndTarget(), utils::Vector{b.Constant(i32(42))})}); + utils::Vector{b.Return(func, utils::Vector{b.Constant(i32(42))})}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -165,7 +165,7 @@ TEST_F(SpvGeneratorImplTest, Function_Parameters) { auto* func = b.CreateFunction("foo", i32); func->SetParams(utils::Vector{x, y}); func->StartTarget()->SetInstructions( - utils::Vector{result, b.Branch(func->EndTarget(), utils::Vector{result})}); + utils::Vector{result, b.Return(func, utils::Vector{result})}); mod.SetName(x, "x"); mod.SetName(y, "y"); @@ -193,12 +193,12 @@ TEST_F(SpvGeneratorImplTest, Function_Call) { auto* foo = b.CreateFunction("foo", i32_ty); foo->SetParams(utils::Vector{x, y}); foo->StartTarget()->SetInstructions( - utils::Vector{result, b.Branch(foo->EndTarget(), utils::Vector{result})}); + utils::Vector{result, b.Return(foo, utils::Vector{result})}); auto* bar = b.CreateFunction("bar", mod.Types().void_()); bar->StartTarget()->SetInstructions(utils::Vector{ b.UserCall(i32_ty, foo, utils::Vector{b.Constant(i32(2)), b.Constant(i32(3))}), - b.Branch(bar->EndTarget())}); + b.Return(bar)}); generator_.EmitFunction(foo); generator_.EmitFunction(bar); @@ -227,11 +227,11 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_Call_Void) { auto* foo = b.CreateFunction("foo", mod.Types().void_()); - foo->StartTarget()->SetInstructions(utils::Vector{b.Branch(foo->EndTarget())}); + foo->StartTarget()->SetInstructions(utils::Vector{b.Return(foo)}); auto* bar = b.CreateFunction("bar", mod.Types().void_()); - bar->StartTarget()->SetInstructions(utils::Vector{ - b.UserCall(mod.Types().void_(), foo, utils::Empty), b.Branch(bar->EndTarget())}); + bar->StartTarget()->SetInstructions( + utils::Vector{b.UserCall(mod.Types().void_(), foo, utils::Empty), b.Return(bar)}); generator_.EmitFunction(foo); generator_.EmitFunction(bar); diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc index ff7defc320..b526c91c32 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc @@ -25,7 +25,7 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) { auto* i = b.CreateIf(b.Constant(true)); i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); - i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); func->StartTarget()->SetInstructions(utils::Vector{i}); @@ -50,7 +50,7 @@ TEST_F(SpvGeneratorImplTest, If_FalseEmpty) { auto* i = b.CreateIf(b.Constant(true)); i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); - i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); auto* true_block = i->True(); true_block->SetInstructions(utils::Vector{ @@ -84,7 +84,7 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty) { auto* i = b.CreateIf(b.Constant(true)); i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); - i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); auto* false_block = i->False(); false_block->SetInstructions(utils::Vector{ @@ -117,8 +117,8 @@ TEST_F(SpvGeneratorImplTest, If_BothBranchesReturn) { auto* func = b.CreateFunction("foo", mod.Types().void_()); auto* i = b.CreateIf(b.Constant(true)); - i->True()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); - i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + i->True()->SetInstructions(utils::Vector{b.Return(func)}); + i->False()->SetInstructions(utils::Vector{b.Return(func)}); func->StartTarget()->SetInstructions(utils::Vector{i}); diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc index 4a42bdc664..d66df98885 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc @@ -25,7 +25,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_NoInit) { auto* ty = mod.Types().Get(mod.Types().i32(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); - func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty), b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty), b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -49,7 +49,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_WithInit) { auto* v = b.Declare(ty); v->SetInitializer(b.Constant(42_i)); - func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{v, b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -73,7 +73,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_Name) { auto* ty = mod.Types().Get(mod.Types().i32(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); auto* v = b.Declare(ty); - func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{v, b.Return(func)}); mod.SetName(v, "myvar"); generator_.EmitFunction(func); @@ -101,8 +101,8 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) { auto* i = b.CreateIf(b.Constant(true)); i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())}); - i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); - i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + i->False()->SetInstructions(utils::Vector{b.Return(func)}); + i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); func->StartTarget()->SetInstructions(utils::Vector{i}); @@ -138,7 +138,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_Load) { auto* ty = mod.Types().Get(store_ty, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); auto* v = b.Declare(ty); - func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v), b.Branch(func->EndTarget())}); + func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v), b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -162,7 +162,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_Store) { builtin::Access::kReadWrite); auto* v = b.Declare(ty); func->StartTarget()->SetInstructions( - utils::Vector{v, b.Store(v, b.Constant(42_i)), b.Branch(func->EndTarget())}); + utils::Vector{v, b.Store(v, b.Constant(42_i)), b.Return(func)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"