[ir] Split the Terminator into two nodes.

This CL moves to having a `FunctionTerminator` and a `RootTerminator` so
we can assert if the IR is in a function depending on the terminator
seen.

Bug: tint:1929
Change-Id: Ie9e3aed71b7cf3b91439efbcca20885ec2cabe24
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131281
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Reviewed-by: James Price <jrprice@google.com>

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",

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

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<Value*, 2> args;
 };

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<Block>();
 }
 
-Terminator* Builder::CreateTerminator() {
-    return ir.flow_nodes.Create<Terminator>();
+RootTerminator* Builder::CreateRootTerminator() {
+    return ir.flow_nodes.Create<RootTerminator>();
+}
+
+FunctionTerminator* Builder::CreateFunctionTerminator() {
+    return ir.flow_nodes.Create<FunctionTerminator>();
 }
 
 Function* Builder::CreateFunction() {
     auto* ir_func = ir.flow_nodes.Create<Function>();
     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);

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

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"

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

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<private, u32, read_write>) = 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<private, u32, read_write>) = var private read_write
 store %1(ref<private, u32, read_write>), 2u
-ret
+
+
 
 )");
 }
@@ -2117,7 +2119,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
     EXPECT_EQ(Disassemble(m), R"(%fn0 = block
 %1(ref<private, vec3<f32>, read_write>) = var private read_write
 store %1(ref<private, vec3<f32>, read_write>), vec3<f32> 0.0f
-ret
+
+
 
 )");
 }
@@ -2135,7 +2138,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
     EXPECT_EQ(Disassemble(m), R"(%fn0 = block
 %1(ref<private, f32, read_write>) = var private read_write
 store %1(ref<private, f32, read_write>), 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<private, i32, read_write>) = var private read_write
 store %1(ref<private, i32, read_write>), 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<private, f32, read_write>) = var private read_write
 store %1(ref<private, f32, read_write>), 1.0f
-ret
+
+
 
 %fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
   %fn2 = block

src/tint/ir/debug.cc

@@ -18,10 +18,10 @@
 #include <unordered_set>
 
 #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
             });
     };

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<Terminator>()) {
+            if (b->branch.target->Is<FunctionTerminator>()) {
                 Indent() << "ret";
+            } else if (b->branch.target->Is<RootTerminator>()) {
+                // 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<Terminator>()) {
+            if (!b->branch.target->Is<FunctionTerminator>()) {
                 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;
         });
 }

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<Function, FlowNode> {
     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);

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

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<FunctionTerminator, FlowNode> {
+  public:
+    /// Constructor
+    FunctionTerminator();
+    ~FunctionTerminator() override;
+};
+
+}  // namespace tint::ir
+
+#endif  // SRC_TINT_IR_FUNCTION_TERMINATOR_H_

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

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<Terminator, FlowNode> {
+class RootTerminator : public utils::Castable<RootTerminator, FlowNode> {
   public:
     /// Constructor
-    Terminator();
-    ~Terminator() override;
+    RootTerminator();
+    ~RootTerminator() override;
 };
 
 }  // namespace tint::ir
 
-#endif  // SRC_TINT_IR_TERMINATOR_H_
+#endif  // SRC_TINT_IR_ROOT_TERMINATOR_H_