diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index a805fe12a6..2f5bc0e62e 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -1154,6 +1154,8 @@ libtint_source_set("libtint_ir_src") { "ir/flow_node.h", "ir/function.cc", "ir/function.h", + "ir/function_terminator.cc", + "ir/function_terminator.h", "ir/if.cc", "ir/if.h", "ir/instruction.cc", @@ -1162,12 +1164,12 @@ libtint_source_set("libtint_ir_src") { "ir/loop.h", "ir/module.cc", "ir/module.h", + "ir/root_terminator.cc", + "ir/root_terminator.h", "ir/store.cc", "ir/store.h", "ir/switch.cc", "ir/switch.h", - "ir/terminator.cc", - "ir/terminator.h", "ir/unary.cc", "ir/unary.h", "ir/user_call.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index 2e1ded863f..87333b3476 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -730,6 +730,8 @@ if(${TINT_BUILD_IR}) ir/flow_node.h ir/function.cc ir/function.h + ir/function_terminator.cc + ir/function_terminator.h ir/if.cc ir/if.h ir/instruction.cc @@ -738,12 +740,12 @@ if(${TINT_BUILD_IR}) ir/loop.h ir/module.cc ir/module.h + ir/root_terminator.cc + ir/root_terminator.h ir/store.cc ir/store.h ir/switch.cc ir/switch.h - ir/terminator.cc - ir/terminator.h ir/unary.cc ir/unary.h ir/user_call.cc diff --git a/src/tint/ir/branch.h b/src/tint/ir/branch.h index 4fc8d7c1f6..667f131f87 100644 --- a/src/tint/ir/branch.h +++ b/src/tint/ir/branch.h @@ -26,7 +26,7 @@ struct Branch { FlowNode* target = nullptr; /// The arguments provided for that branch. These arguments could be the - /// return value in the case of a branch to the terminator, or they could + /// return value in the case of a branch to the function terminator, or they could /// be the basic block arguments passed into the block. utils::Vector args; }; diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc index 7090cb6e39..8badd31809 100644 --- a/src/tint/ir/builder.cc +++ b/src/tint/ir/builder.cc @@ -29,8 +29,8 @@ ir::Block* Builder::CreateRootBlockIfNeeded() { ir.root_block = CreateBlock(); // Everything in the module scope must have been const-eval's, so everything will go into a - // single block. So, we can create the terminator for the root-block now. - ir.root_block->branch.target = CreateTerminator(); + // single block. So, we can create the root terminator for the root-block now. + ir.root_block->branch.target = CreateRootTerminator(); } return ir.root_block; } @@ -39,14 +39,18 @@ Block* Builder::CreateBlock() { return ir.flow_nodes.Create(); } -Terminator* Builder::CreateTerminator() { - return ir.flow_nodes.Create(); +RootTerminator* Builder::CreateRootTerminator() { + return ir.flow_nodes.Create(); +} + +FunctionTerminator* Builder::CreateFunctionTerminator() { + return ir.flow_nodes.Create(); } Function* Builder::CreateFunction() { auto* ir_func = ir.flow_nodes.Create(); ir_func->start_target = CreateBlock(); - ir_func->end_target = CreateTerminator(); + ir_func->end_target = CreateFunctionTerminator(); // Function is always branching into the start target ir_func->start_target->inbound_branches.Push(ir_func); diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h index 618679bf6e..24c20d0c63 100644 --- a/src/tint/ir/builder.h +++ b/src/tint/ir/builder.h @@ -26,12 +26,13 @@ #include "src/tint/ir/convert.h" #include "src/tint/ir/discard.h" #include "src/tint/ir/function.h" +#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/loop.h" #include "src/tint/ir/module.h" +#include "src/tint/ir/root_terminator.h" #include "src/tint/ir/store.h" #include "src/tint/ir/switch.h" -#include "src/tint/ir/terminator.h" #include "src/tint/ir/unary.h" #include "src/tint/ir/user_call.h" #include "src/tint/ir/value.h" @@ -59,8 +60,11 @@ class Builder { /// @returns a new block flow node Block* CreateBlock(); - /// @returns a new terminator flow node - Terminator* CreateTerminator(); + /// @returns a new root terminator flow node + RootTerminator* CreateRootTerminator(); + + /// @returns a new function terminator flow node + FunctionTerminator* CreateFunctionTerminator(); /// Creates a function flow node /// @returns the flow node diff --git a/src/tint/ir/builder_impl.cc b/src/tint/ir/builder_impl.cc index 6ee8170224..56849d2cea 100644 --- a/src/tint/ir/builder_impl.cc +++ b/src/tint/ir/builder_impl.cc @@ -61,7 +61,6 @@ #include "src/tint/ir/module.h" #include "src/tint/ir/store.h" #include "src/tint/ir/switch.h" -#include "src/tint/ir/terminator.h" #include "src/tint/ir/value.h" #include "src/tint/program.h" #include "src/tint/sem/builtin.h" diff --git a/src/tint/ir/builder_impl.h b/src/tint/ir/builder_impl.h index fd80a3e34d..f6bed120bd 100644 --- a/src/tint/ir/builder_impl.h +++ b/src/tint/ir/builder_impl.h @@ -60,14 +60,6 @@ class UnaryOpExpression; class WhileStatement; class Variable; } // namespace tint::ast -namespace tint::ir { -class Block; -class If; -class Function; -class Loop; -class Switch; -class Terminator; -} // namespace tint::ir namespace tint::sem { class Builtin; } // namespace tint::sem diff --git a/src/tint/ir/builder_impl_test.cc b/src/tint/ir/builder_impl_test.cc index 9fc8edf49f..493bbfb8ca 100644 --- a/src/tint/ir/builder_impl_test.cc +++ b/src/tint/ir/builder_impl_test.cc @@ -1567,7 +1567,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) { EXPECT_EQ(Disassemble(m), R"(%fn0 = block %1(ref) = var private read_write -ret + + )"); } @@ -1583,7 +1584,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) { EXPECT_EQ(Disassemble(m), R"(%fn0 = block %1(ref) = var private read_write store %1(ref), 2u -ret + + )"); } @@ -2117,7 +2119,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) { EXPECT_EQ(Disassemble(m), R"(%fn0 = block %1(ref, read_write>) = var private read_write store %1(ref, read_write>), vec3 0.0f -ret + + )"); } @@ -2135,7 +2138,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) { EXPECT_EQ(Disassemble(m), R"(%fn0 = block %1(ref) = var private read_write store %1(ref), 1.0f -ret + + %fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn2 = block @@ -2159,7 +2163,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) { EXPECT_EQ(Disassemble(m), R"(%fn0 = block %1(ref) = var private read_write store %1(ref), 1i -ret + + %fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn2 = block @@ -2199,7 +2204,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) { EXPECT_EQ(Disassemble(m), R"(%fn0 = block %1(ref) = var private read_write store %1(ref), 1.0f -ret + + %fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)] %fn2 = block diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc index f69805e5de..655c45507e 100644 --- a/src/tint/ir/debug.cc +++ b/src/tint/ir/debug.cc @@ -18,10 +18,10 @@ #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/switch.h" -#include "src/tint/ir/terminator.h" #include "src/tint/switch.h" #include "src/tint/utils/string_stream.h" @@ -136,7 +136,7 @@ std::string Debug::AsDotGraph(const Module* mod) { Graph(l->continuing.target); Graph(l->merge.target); }, - [&](const ir::Terminator*) { + [&](const ir::FunctionTerminator*) { // Already done }); }; diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc index acadde3d94..b3deeafaa5 100644 --- a/src/tint/ir/disassembler.cc +++ b/src/tint/ir/disassembler.cc @@ -15,10 +15,11 @@ #include "src/tint/ir/disassembler.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/root_terminator.h" #include "src/tint/ir/switch.h" -#include "src/tint/ir/terminator.h" #include "src/tint/switch.h" #include "src/tint/type/type.h" #include "src/tint/utils/scoped_assignment.h" @@ -136,8 +137,10 @@ void Disassembler::Walk(const FlowNode* node) { Indent() << "%fn" << GetIdForNode(b) << " = block" << std::endl; EmitBlockInstructions(b); - if (b->branch.target->Is()) { + if (b->branch.target->Is()) { Indent() << "ret"; + } else if (b->branch.target->Is()) { + // Nothing to do } else { Indent() << "branch " << "%fn" << GetIdForNode(b->branch.target); @@ -153,7 +156,7 @@ void Disassembler::Walk(const FlowNode* node) { } out_ << std::endl; - if (!b->branch.target->Is()) { + if (!b->branch.target->Is()) { out_ << std::endl; } @@ -272,10 +275,12 @@ void Disassembler::Walk(const FlowNode* node) { Walk(l->merge.target); } }, - [&](const ir::Terminator*) { - if (in_function_) { - Indent() << "func_end" << std::endl; - } + [&](const ir::FunctionTerminator*) { + TINT_ASSERT(IR, in_function_); + Indent() << "func_end" << std::endl << std::endl; + }, + [&](const ir::RootTerminator*) { + TINT_ASSERT(IR, !in_function_); out_ << std::endl; }); } diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h index 7393b94b73..e4aa4edd90 100644 --- a/src/tint/ir/function.h +++ b/src/tint/ir/function.h @@ -24,7 +24,7 @@ // Forward declarations namespace tint::ir { class Block; -class Terminator; +class FunctionTerminator; } // namespace tint::ir namespace tint::ir { @@ -84,7 +84,7 @@ class Function : public utils::Castable { Block* start_target = nullptr; /// The end target is the end of the function. It is used as the branch target if a return is /// encountered in the function. - Terminator* end_target = nullptr; + FunctionTerminator* end_target = nullptr; }; utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value); diff --git a/src/tint/ir/function_terminator.cc b/src/tint/ir/function_terminator.cc new file mode 100644 index 0000000000..59893117ee --- /dev/null +++ b/src/tint/ir/function_terminator.cc @@ -0,0 +1,25 @@ +// 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/ir/function_terminator.h" + +TINT_INSTANTIATE_TYPEINFO(tint::ir::FunctionTerminator); + +namespace tint::ir { + +FunctionTerminator::FunctionTerminator() : Base() {} + +FunctionTerminator::~FunctionTerminator() = default; + +} // namespace tint::ir diff --git a/src/tint/ir/function_terminator.h b/src/tint/ir/function_terminator.h new file mode 100644 index 0000000000..668eab6cd4 --- /dev/null +++ b/src/tint/ir/function_terminator.h @@ -0,0 +1,33 @@ +// 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/flow_node.h" + +namespace tint::ir { + +/// Flow node used as the end of a function. Must only be used as the `end_target` in a function +/// flow node. There are no instructions and no branches from this node. +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/terminator.cc b/src/tint/ir/root_terminator.cc similarity index 77% rename from src/tint/ir/terminator.cc rename to src/tint/ir/root_terminator.cc index 6d7678d841..bfccf46a21 100644 --- a/src/tint/ir/terminator.cc +++ b/src/tint/ir/root_terminator.cc @@ -12,14 +12,14 @@ // See the License for the specific language governing permissions and // limitations under the License. -#include "src/tint/ir/terminator.h" +#include "src/tint/ir/root_terminator.h" -TINT_INSTANTIATE_TYPEINFO(tint::ir::Terminator); +TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator); namespace tint::ir { -Terminator::Terminator() : Base() {} +RootTerminator::RootTerminator() : Base() {} -Terminator::~Terminator() = default; +RootTerminator::~RootTerminator() = default; } // namespace tint::ir diff --git a/src/tint/ir/terminator.h b/src/tint/ir/root_terminator.h similarity index 78% rename from src/tint/ir/terminator.h rename to src/tint/ir/root_terminator.h index 35fbcee1f0..361aa6d464 100644 --- a/src/tint/ir/terminator.h +++ b/src/tint/ir/root_terminator.h @@ -12,8 +12,8 @@ // See the License for the specific language governing permissions and // limitations under the License. -#ifndef SRC_TINT_IR_TERMINATOR_H_ -#define SRC_TINT_IR_TERMINATOR_H_ +#ifndef SRC_TINT_IR_ROOT_TERMINATOR_H_ +#define SRC_TINT_IR_ROOT_TERMINATOR_H_ #include "src/tint/ir/flow_node.h" @@ -21,13 +21,13 @@ namespace tint::ir { /// Flow node used as the end of a function. Must only be used as the `end_target` in a function /// flow node. There are no instructions and no branches from this node. -class Terminator : public utils::Castable { +class RootTerminator : public utils::Castable { public: /// Constructor - Terminator(); - ~Terminator() override; + RootTerminator(); + ~RootTerminator() override; }; } // namespace tint::ir -#endif // SRC_TINT_IR_TERMINATOR_H_ +#endif // SRC_TINT_IR_ROOT_TERMINATOR_H_