[ir] Add the Exit instructions.
This CL adds the ExitIf, ExitLoop and ExitSwitch instructions. The dump
to Dot has been removed as the graph is substantially different and it
needs a full re-write if we want to draw the graph.
Bug: tint:1718
Change-Id: I5ff4282abaa7542575d4f8b4b8640a3ed4d5c68f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134464
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 344607d..bff6c37 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1230,12 +1230,16 @@
"ir/continue.h",
"ir/convert.cc",
"ir/convert.h",
- "ir/debug.cc",
- "ir/debug.h",
"ir/disassembler.cc",
"ir/disassembler.h",
"ir/discard.cc",
"ir/discard.h",
+ "ir/exit_if.cc",
+ "ir/exit_if.h",
+ "ir/exit_loop.cc",
+ "ir/exit_loop.h",
+ "ir/exit_switch.cc",
+ "ir/exit_switch.h",
"ir/function.cc",
"ir/function.h",
"ir/function_param.cc",
@@ -1252,8 +1256,6 @@
"ir/module.h",
"ir/return.cc",
"ir/return.h",
- "ir/root_terminator.cc",
- "ir/root_terminator.h",
"ir/store.cc",
"ir/store.h",
"ir/switch.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index ff69240..e5e529b 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -738,12 +738,16 @@
ir/continue.h
ir/convert.cc
ir/convert.h
- ir/debug.cc
- ir/debug.h
ir/disassembler.cc
ir/disassembler.h
ir/discard.cc
ir/discard.h
+ ir/exit_if.cc
+ ir/exit_if.h
+ ir/exit_loop.cc
+ ir/exit_loop.h
+ ir/exit_switch.cc
+ ir/exit_switch.h
ir/from_program.cc
ir/from_program.h
ir/function.cc
@@ -762,8 +766,6 @@
ir/module.h
ir/return.cc
ir/return.h
- ir/root_terminator.cc
- ir/root_terminator.h
ir/store.cc
ir/store.h
ir/switch.cc
diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc
index bfdc41f..5d82479 100644
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@@ -49,7 +49,6 @@
#include "tint/tint.h"
#if TINT_BUILD_IR
-#include "src/tint/ir/debug.h" // nogncheck
#include "src/tint/ir/disassembler.h" // nogncheck
#include "src/tint/ir/from_program.h" // nogncheck
#include "src/tint/ir/module.h" // nogncheck
@@ -110,7 +109,6 @@
#if TINT_BUILD_IR
bool dump_ir = false;
- bool dump_ir_graph = false;
bool use_ir = false;
#endif // TINT_BUILD_IR
@@ -374,8 +372,6 @@
#if TINT_BUILD_IR
} else if (arg == "--dump-ir") {
opts->dump_ir = true;
- } else if (arg == "--dump-ir-graph") {
- opts->dump_ir_graph = true;
} else if (arg == "--use-ir") {
opts->use_ir = true;
#endif // TINT_BUILD_IR
@@ -1072,7 +1068,7 @@
#endif // TINT_BUILD_SYNTAX_TREE_WRITER
#if TINT_BUILD_IR
- if (options.dump_ir || options.dump_ir_graph) {
+ if (options.dump_ir) {
auto result = tint::ir::FromProgram(program.get());
if (!result) {
std::cerr << "Failed to build IR from program: " << result.Failure() << std::endl;
@@ -1082,10 +1078,6 @@
tint::ir::Disassembler d(mod);
std::cout << d.Disassemble() << std::endl;
}
- if (options.dump_ir_graph) {
- auto graph = tint::ir::Debug::AsDotGraph(&mod);
- WriteFile("tint.dot", "w", graph);
- }
}
}
#endif // TINT_BUILD_IR
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 597ff2c..d40002d 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -46,18 +46,6 @@
return instructions_.Back()->As<ir::Branch>();
}
- /// @param target the block to see if we trampoline too
- /// @returns if this block just branches to the provided target.
- bool IsTrampoline(const Block* target) const {
- if (instructions_.Length() != 1) {
- return false;
- }
- if (auto* inst = instructions_.Front()->As<ir::Branch>()) {
- return inst->To() == target;
- }
- return false;
- }
-
/// Sets the instructions in the block
/// @param instructions the instructions to set
void SetInstructions(utils::VectorRef<const Instruction*> instructions) {
diff --git a/src/tint/ir/branch.cc b/src/tint/ir/branch.cc
index 0918962..191831f 100644
--- a/src/tint/ir/branch.cc
+++ b/src/tint/ir/branch.cc
@@ -28,13 +28,6 @@
}
}
-Branch::Branch(Block* to, utils::VectorRef<Value*> 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 aa3d1a0..5c926e9 100644
--- a/src/tint/ir/branch.h
+++ b/src/tint/ir/branch.h
@@ -29,15 +29,8 @@
/// A branch instruction.
class Branch : public utils::Castable<Branch, Instruction> {
public:
- /// Constructor
- /// @param to the block to branch too
- /// @param args the branch arguments
- explicit Branch(Block* to, utils::VectorRef<Value*> args = {});
~Branch() override;
- /// @returns the block being branched too.
- const Block* To() const { return to_; }
-
/// @returns the branch arguments
utils::VectorRef<Value*> Args() const { return args_; }
@@ -47,7 +40,6 @@
explicit Branch(utils::VectorRef<Value*> args);
private:
- Block* to_ = nullptr;
utils::Vector<Value*, 2> args_;
};
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index 4bdb53d..1daf458 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -37,10 +37,6 @@
return ir.blocks.Create<Block>();
}
-RootTerminator* Builder::CreateRootTerminator() {
- return ir.blocks.Create<RootTerminator>();
-}
-
Function* Builder::CreateFunction(std::string_view name,
const type::Type* return_type,
Function::PipelineStage stage,
@@ -205,10 +201,6 @@
return ir.values.Create<ir::Var>(type);
}
-ir::Branch* Builder::Branch(Block* to, utils::VectorRef<Value*> args) {
- return ir.values.Create<ir::Branch>(to, args);
-}
-
ir::Return* Builder::Return(Function* func, utils::VectorRef<Value*> args) {
return ir.values.Create<ir::Return>(func, args);
}
@@ -220,6 +212,17 @@
ir::Continue* Builder::Continue(Loop* loop) {
return ir.values.Create<ir::Continue>(loop);
}
+ir::ExitSwitch* Builder::ExitSwitch(Switch* sw) {
+ return ir.values.Create<ir::ExitSwitch>(sw);
+}
+
+ir::ExitLoop* Builder::ExitLoop(Loop* loop) {
+ return ir.values.Create<ir::ExitLoop>(loop);
+}
+
+ir::ExitIf* Builder::ExitIf(If* i, utils::VectorRef<Value*> args) {
+ return ir.values.Create<ir::ExitIf>(i, args);
+}
ir::BlockParam* Builder::BlockParam(const type::Type* type) {
return ir.values.Create<ir::BlockParam>(type);
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 9a2fe18..8a4c021 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -28,6 +28,9 @@
#include "src/tint/ir/continue.h"
#include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
#include "src/tint/ir/function.h"
#include "src/tint/ir/function_param.h"
#include "src/tint/ir/if.h"
@@ -35,7 +38,6 @@
#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"
#include "src/tint/ir/unary.h"
@@ -64,9 +66,6 @@
/// @returns a new block flow node
Block* CreateBlock();
- /// @returns a new root terminator flow node
- RootTerminator* CreateRootTerminator();
-
/// Creates a function flow node
/// @param name the function name
/// @param return_type the function return type
@@ -349,11 +348,21 @@
/// @returns the instruction
ir::Continue* Continue(Loop* loop);
- /// Creates a branch declaration
- /// @param to the node being branched too
+ /// Creates an exit switch instruction
+ /// @param sw the switch being exited
+ /// @returns the instruction
+ ir::ExitSwitch* ExitSwitch(Switch* sw);
+
+ /// Creates an exit loop instruction
+ /// @param loop the loop being exited
+ /// @returns the instruction
+ ir::ExitLoop* ExitLoop(Loop* loop);
+
+ /// Creates an exit if instruction
+ /// @param i the if being exited
/// @param args the branch arguments
/// @returns the instruction
- ir::Branch* Branch(Block* to, utils::VectorRef<Value*> args = {});
+ ir::ExitIf* ExitIf(If* i, utils::VectorRef<Value*> args = {});
/// Creates a new `BlockParam`
/// @param type the parameter type
diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc
deleted file mode 100644
index c719698..0000000
--- a/src/tint/ir/debug.cc
+++ /dev/null
@@ -1,95 +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.
-
-#include "src/tint/ir/debug.h"
-
-#include <unordered_map>
-#include <unordered_set>
-
-#include "src/tint/ir/block.h"
-#include "src/tint/ir/continue.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"
-
-namespace tint::ir {
-
-// static
-std::string Debug::AsDotGraph(const Module* mod) {
- size_t block_count = 0;
-
- std::unordered_set<const Block*> visited;
- std::unordered_set<const Block*> merge_blocks;
- std::unordered_map<const Block*, std::string> block_to_name;
- utils::StringStream out;
-
- auto name_for = [&](const Block* blk) -> std::string {
- if (block_to_name.count(blk) > 0) {
- return block_to_name[blk];
- }
-
- std::string name = "blk_" + std::to_string(block_count);
- block_count += 1;
-
- block_to_name[blk] = name;
- return name;
- };
-
- std::function<void(const Block*)> Graph = [&](const Block* blk) {
- if (visited.count(blk) > 0) {
- return;
- }
- visited.insert(blk);
-
- 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]";
- }
-
- out << std::endl;
-
- if (b->Branch()->Is<ir::Return>()) {
- return;
- } else if (auto* cont = b->Branch()->As<ir::Continue>()) {
- Graph(cont->Loop()->Continuing());
- } else {
- Graph(b->Branch()->To());
- }
- });
- };
-
- out << "digraph G {" << std::endl;
- for (const auto* func : mod->functions) {
- // Cluster each function to label and draw a box around it.
- 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;
- Graph(func->StartTarget());
- out << "}" << std::endl;
- }
- out << "}";
- return out.str();
-}
-
-} // namespace tint::ir
diff --git a/src/tint/ir/debug.h b/src/tint/ir/debug.h
deleted file mode 100644
index 2363776..0000000
--- a/src/tint/ir/debug.h
+++ /dev/null
@@ -1,40 +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_DEBUG_H_
-#define SRC_TINT_IR_DEBUG_H_
-
-#include <string>
-
-#include "src/tint/ir/module.h"
-
-namespace tint::ir {
-
-/// Helper class to debug IR.
-class Debug {
- public:
- /// Returns the module as a dot graph
- /// @param mod the module to emit
- /// @returns the dot graph for the given module
- static std::string AsDotGraph(const Module* mod);
-
- /// Returns the module as a string
- /// @param mod the module to emit
- /// @returns the string representation of the module
- static std::string AsString(const Module* mod);
-};
-
-} // namespace tint::ir
-
-#endif // SRC_TINT_IR_DEBUG_H_
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 6097726..d39438a 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -27,11 +27,13 @@
#include "src/tint/ir/continue.h"
#include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.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"
#include "src/tint/ir/user_call.h"
@@ -70,7 +72,6 @@
for (const auto* inst : b->Instructions()) {
Indent();
EmitInstruction(inst);
- out_ << std::endl;
}
}
@@ -92,8 +93,8 @@
std::string Disassembler::Disassemble() {
if (mod_.root_block) {
Indent() << "# Root block" << std::endl;
- Walk(mod_.root_block);
- Walk(mod_.root_block->Branch()->To());
+ WalkInternal(mod_.root_block);
+ out_ << std::endl;
}
for (auto* func : mod_.functions) {
@@ -108,41 +109,33 @@
}
visited_.Add(blk);
- tint::Switch(
- blk,
- [&](const ir::RootTerminator* t) {
- TINT_ASSERT(IR, !in_function_);
- Indent() << "%b" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
- },
- [&](const ir::Block* b) {
- // If this block is dead, nothing to do
- if (!b->HasBranchTarget()) {
- return;
- }
+ // If this block is dead, nothing to do
+ if (!blk->HasBranchTarget()) {
+ return;
+ }
- Indent() << "%b" << IdOf(b) << " = block";
- if (!b->Params().IsEmpty()) {
- out_ << " (";
- for (auto* p : b->Params()) {
- if (p != b->Params().Front()) {
- out_ << ", ";
- }
- EmitValue(p);
- }
- out_ << ")";
- }
+ WalkInternal(blk);
+}
- out_ << " {" << std::endl;
- {
- ScopedIndent si(indent_size_);
- EmitBlockInstructions(b);
+void Disassembler::WalkInternal(const Block* blk) {
+ Indent() << "%b" << IdOf(blk) << " = block";
+ if (!blk->Params().IsEmpty()) {
+ out_ << " (";
+ for (auto* p : blk->Params()) {
+ if (p != blk->Params().Front()) {
+ out_ << ", ";
}
- Indent() << "}" << std::endl;
+ EmitValue(p);
+ }
+ out_ << ")";
+ }
- if (!b->Branch()->Is<ir::Return>()) {
- out_ << std::endl;
- }
- });
+ out_ << " {" << std::endl;
+ {
+ ScopedIndent si(indent_size_);
+ EmitBlockInstructions(blk);
+ }
+ Indent() << "}" << std::endl;
}
void Disassembler::EmitFunction(const Function* func) {
@@ -258,33 +251,39 @@
EmitValueWithType(b);
out_ << " = bitcast ";
EmitArgs(b);
+ out_ << std::endl;
},
- [&](const ir::Discard*) { out_ << "discard"; },
+ [&](const ir::Discard*) { out_ << "discard" << std::endl; },
[&](const ir::Builtin* b) {
EmitValueWithType(b);
out_ << " = " << builtin::str(b->Func()) << " ";
EmitArgs(b);
+ out_ << std::endl;
},
[&](const ir::Construct* c) {
EmitValueWithType(c);
out_ << " = construct ";
EmitArgs(c);
+ out_ << std::endl;
},
[&](const ir::Convert* c) {
EmitValueWithType(c);
out_ << " = convert " << c->FromType()->FriendlyName() << ", ";
EmitArgs(c);
+ out_ << std::endl;
},
[&](const ir::Load* l) {
EmitValueWithType(l);
out_ << " = load ";
EmitValue(l->From());
+ out_ << std::endl;
},
[&](const ir::Store* s) {
out_ << "store ";
EmitValue(s->To());
out_ << ", ";
EmitValue(s->From());
+ out_ << std::endl;
},
[&](const ir::UserCall* uc) {
EmitValueWithType(uc);
@@ -293,6 +292,7 @@
out_ << ", ";
}
EmitArgs(uc);
+ out_ << std::endl;
},
[&](const ir::Var* v) {
EmitValueWithType(v);
@@ -301,6 +301,7 @@
out_ << ", ";
EmitValue(v->Initializer());
}
+ out_ << std::endl;
},
[&](const ir::Branch* b) { EmitBranch(b); },
[&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; });
@@ -332,15 +333,18 @@
ScopedIndent si(indent_size_);
Indent() << "# True block" << std::endl;
Walk(i->True());
+ out_ << std::endl;
}
if (has_false) {
ScopedIndent si(indent_size_);
Indent() << "# False block" << std::endl;
Walk(i->False());
+ out_ << std::endl;
}
if (i->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl;
Walk(i->Merge());
+ out_ << std::endl;
}
}
@@ -358,16 +362,19 @@
{
ScopedIndent si(indent_size_);
Walk(l->Start());
+ out_ << std::endl;
}
if (l->Continuing()->HasBranchTarget()) {
ScopedIndent si(indent_size_);
Indent() << "# Continuing block" << std::endl;
Walk(l->Continuing());
+ out_ << std::endl;
}
if (l->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl;
Walk(l->Merge());
+ out_ << std::endl;
}
}
@@ -402,29 +409,31 @@
ScopedIndent si(indent_size_);
Indent() << "# Case block" << std::endl;
Walk(c.Start());
+ out_ << std::endl;
}
if (s->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl;
Walk(s->Merge());
+ out_ << std::endl;
}
}
void Disassembler::EmitBranch(const Branch* b) {
- std::string suffix = "";
- if (b->Is<ir::Return>()) {
- out_ << "ret";
- } else if (auto* cont = b->As<ir::Continue>()) {
- out_ << "continue %b" << IdOf(cont->Loop()->Continuing());
- } else if (auto* bi = b->As<ir::BreakIf>()) {
- out_ << "break_if ";
- EmitValue(bi->Condition());
- out_ << " %b" << IdOf(bi->Loop()->Start());
- } else {
- out_ << "br %b" << IdOf(b->To());
- if (b->To()->Is<RootTerminator>()) {
- suffix = "root_end";
- }
- }
+ tint::Switch(
+ b, //
+ [&](const ir::Return*) { out_ << "ret"; },
+ [&](const ir::Continue* cont) {
+ out_ << "continue %b" << IdOf(cont->Loop()->Continuing());
+ },
+ [&](const ir::ExitIf* ei) { out_ << "exit_if %b" << IdOf(ei->If()->Merge()); },
+ [&](const ir::ExitSwitch* es) { out_ << "exit_switch %b" << IdOf(es->Switch()->Merge()); },
+ [&](const ir::ExitLoop* el) { out_ << "exit_loop %b" << IdOf(el->Loop()->Merge()); },
+ [&](const ir::BreakIf* bi) {
+ out_ << "break_if ";
+ EmitValue(bi->Condition());
+ out_ << " %b" << IdOf(bi->Loop()->Start());
+ },
+ [&](Default) { out_ << "Unknown branch " << b->TypeInfo().name; });
if (!b->Args().IsEmpty()) {
out_ << " ";
@@ -435,9 +444,7 @@
EmitValue(v);
}
}
- if (!suffix.empty()) {
- out_ << " # " << suffix;
- }
+ out_ << std::endl;
}
void Disassembler::EmitArgs(const Call* call) {
@@ -508,6 +515,7 @@
EmitValue(b->LHS());
out_ << ", ";
EmitValue(b->RHS());
+ out_ << std::endl;
}
void Disassembler::EmitUnary(const Unary* u) {
@@ -523,6 +531,7 @@
}
out_ << " ";
EmitValue(u->Val());
+ out_ << std::endl;
}
} // namespace tint::ir
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index f171031..06f5b08 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -57,6 +57,7 @@
std::string_view IdOf(const Value* node);
void Walk(const Block* blk);
+ void WalkInternal(const Block* blk);
void EmitFunction(const Function* func);
void EmitInstruction(const Instruction* inst);
void EmitValueWithType(const Value* val);
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/exit_if.cc
similarity index 63%
rename from src/tint/ir/root_terminator.cc
rename to src/tint/ir/exit_if.cc
index bfccf46..8b7de7f 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/exit_if.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,20 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/exit_if.h"
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/if.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitIf);
namespace tint::ir {
-RootTerminator::RootTerminator() : Base() {}
+ExitIf::ExitIf(ir::If* i, utils::VectorRef<Value*> args) : Base(args), if_(i) {
+ TINT_ASSERT(IR, if_);
+ if_->AddUsage(this);
+ if_->Merge()->AddInboundBranch(this);
+}
-RootTerminator::~RootTerminator() = default;
+ExitIf::~ExitIf() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/exit_if.h b/src/tint/ir/exit_if.h
new file mode 100644
index 0000000..9ba1421
--- /dev/null
+++ b/src/tint/ir/exit_if.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_EXIT_IF_H_
+#define SRC_TINT_IR_EXIT_IF_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class If;
+} // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit if instruction.
+class ExitIf : public utils::Castable<ExitIf, Branch> {
+ public:
+ /// Constructor
+ /// @param i the if being exited
+ /// @param args the branch arguments
+ explicit ExitIf(ir::If* i, utils::VectorRef<Value*> args = {});
+ ~ExitIf() override;
+
+ /// @returns the if being exited
+ const ir::If* If() const { return if_; }
+
+ private:
+ ir::If* if_ = nullptr;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_EXIT_IF_H_
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/exit_loop.cc
similarity index 63%
copy from src/tint/ir/root_terminator.cc
copy to src/tint/ir/exit_loop.cc
index bfccf46..5fe3910 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/exit_loop.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,20 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/exit_loop.h"
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/loop.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitLoop);
namespace tint::ir {
-RootTerminator::RootTerminator() : Base() {}
+ExitLoop::ExitLoop(ir::Loop* loop) : Base(utils::Empty), loop_(loop) {
+ TINT_ASSERT(IR, loop_);
+ loop_->AddUsage(this);
+ loop_->Merge()->AddInboundBranch(this);
+}
-RootTerminator::~RootTerminator() = default;
+ExitLoop::~ExitLoop() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/exit_loop.h b/src/tint/ir/exit_loop.h
new file mode 100644
index 0000000..1df1119
--- /dev/null
+++ b/src/tint/ir/exit_loop.h
@@ -0,0 +1,45 @@
+// 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_EXIT_LOOP_H_
+#define SRC_TINT_IR_EXIT_LOOP_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Loop;
+} // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit loop instruction.
+class ExitLoop : public utils::Castable<ExitLoop, Branch> {
+ public:
+ /// Constructor
+ /// @param loop the loop being exited
+ explicit ExitLoop(ir::Loop* loop);
+ ~ExitLoop() override;
+
+ /// @returns the loop being exited
+ const ir::Loop* Loop() const { return loop_; }
+
+ private:
+ ir::Loop* loop_ = nullptr;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_EXIT_LOOP_H_
diff --git a/src/tint/ir/root_terminator.cc b/src/tint/ir/exit_switch.cc
similarity index 61%
copy from src/tint/ir/root_terminator.cc
copy to src/tint/ir/exit_switch.cc
index bfccf46..ba6a178 100644
--- a/src/tint/ir/root_terminator.cc
+++ b/src/tint/ir/exit_switch.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,20 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/root_terminator.h"
+#include "src/tint/ir/exit_switch.h"
-TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator);
+#include "src/tint/ir/switch.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitSwitch);
namespace tint::ir {
-RootTerminator::RootTerminator() : Base() {}
+ExitSwitch::ExitSwitch(ir::Switch* sw) : Base(utils::Empty), switch_(sw) {
+ TINT_ASSERT(IR, switch_);
+ switch_->AddUsage(this);
+ switch_->Merge()->AddInboundBranch(this);
+}
-RootTerminator::~RootTerminator() = default;
+ExitSwitch::~ExitSwitch() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/exit_switch.h b/src/tint/ir/exit_switch.h
new file mode 100644
index 0000000..6b406fe
--- /dev/null
+++ b/src/tint/ir/exit_switch.h
@@ -0,0 +1,45 @@
+// 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_EXIT_SWITCH_H_
+#define SRC_TINT_IR_EXIT_SWITCH_H_
+
+#include "src/tint/ir/branch.h"
+#include "src/tint/utils/castable.h"
+
+// Forward declarations
+namespace tint::ir {
+class Switch;
+} // namespace tint::ir
+
+namespace tint::ir {
+
+/// A exit switch instruction.
+class ExitSwitch : public utils::Castable<ExitSwitch, Branch> {
+ public:
+ /// Constructor
+ /// @param sw the switch being exited
+ explicit ExitSwitch(ir::Switch* sw);
+ ~ExitSwitch() override;
+
+ /// @returns the switch being exited
+ const ir::Switch* Switch() const { return switch_; }
+
+ private:
+ ir::Switch* switch_ = nullptr;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_EXIT_SWITCH_H_
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index a4d95e0..67c9f3f 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -62,6 +62,9 @@
#include "src/tint/ast/while_statement.h"
#include "src/tint/ir/block_param.h"
#include "src/tint/ir/builder.h"
+#include "src/tint/ir/exit_if.h"
+#include "src/tint/ir/exit_loop.h"
+#include "src/tint/ir/exit_switch.h"
#include "src/tint/ir/function.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/loop.h"
@@ -98,11 +101,8 @@
using ResultType = utils::Result<Module, diag::List>;
-// For an `if` and `switch` block, the merge has a registered incoming branch instruction of the
-// `if` and `switch. So, to determine if the merge is connected to any of the branches that happend
-// in the `if` or `switch` we need a `count` value that is larger then 1.
-bool IsConnected(const Block* b, uint32_t count) {
- return b->InboundBranches().Length() > count;
+bool IsConnected(const Block* b) {
+ return b->InboundBranches().Length() > 0;
}
/// Impl is the private-implementation of FromProgram().
@@ -176,21 +176,6 @@
current_flow_block_ = nullptr;
}
- void BranchTo(Block* node, utils::VectorRef<Value*> args = {}) {
- TINT_ASSERT(IR, current_flow_block_);
- TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
-
- current_flow_block_->Instructions().Push(builder_.Branch(node, args));
- current_flow_block_ = nullptr;
- }
-
- void BranchToIfNeeded(Block* node) {
- if (!NeedBranch()) {
- return;
- }
- BranchTo(node);
- }
-
Branch* FindEnclosingControl(ControlFlags flags) {
for (auto it = control_stack_.rbegin(); it != control_stack_.rend(); ++it) {
if ((*it)->Is<Loop>()) {
@@ -238,11 +223,6 @@
});
}
- // Add the root terminator if needed
- if (mod.root_block) {
- mod.root_block->Instructions().Push(builder_.Branch(builder_.CreateRootTerminator()));
- }
-
if (diagnostics_.contains_errors()) {
return ResultType(std::move(diagnostics_));
}
@@ -541,7 +521,9 @@
EmitBlock(stmt->body);
// If the true branch did not execute control flow, then go to the Merge().target
- BranchToIfNeeded(if_inst->Merge());
+ if (NeedBranch()) {
+ SetBranch(builder_.ExitIf(if_inst));
+ }
current_flow_block_ = if_inst->False();
if (stmt->else_statement) {
@@ -549,14 +531,16 @@
}
// If the false branch did not execute control flow, then go to the Merge().target
- BranchToIfNeeded(if_inst->Merge());
+ if (NeedBranch()) {
+ SetBranch(builder_.ExitIf(if_inst));
+ }
}
current_flow_block_ = nullptr;
// If both branches went somewhere, then they both returned, continued or broke. So,
// there is no need for the if merge-block and there is nothing to branch to the merge
// block anyway.
- if (IsConnected(if_inst->Merge(), 1)) {
+ if (IsConnected(if_inst->Merge())) {
current_flow_block_ = if_inst->Merge();
}
}
@@ -580,7 +564,7 @@
SetBranch(builder_.Continue(loop_inst));
}
- if (IsConnected(loop_inst->Continuing(), 0)) {
+ if (IsConnected(loop_inst->Continuing())) {
// Note, even if there is no continuing block, we may have branched into the
// continue so we have to set the current block and then emit the branch if needed
// below otherwise empty continuing blocks will fail to branch back to the start
@@ -600,7 +584,7 @@
// target branches, eventually, to the merge, but nothing branched to the
// Continuing() block.
current_flow_block_ = loop_inst->Merge();
- if (!IsConnected(loop_inst->Merge(), 0)) {
+ if (!IsConnected(loop_inst->Merge())) {
current_flow_block_ = nullptr;
}
}
@@ -626,10 +610,14 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get());
- if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
- if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
current_flow_block_->Instructions().Push(if_inst);
+ current_flow_block_ = if_inst->True();
+ SetBranch(builder_.ExitIf(if_inst));
+
+ current_flow_block_ = if_inst->False();
+ SetBranch(builder_.ExitLoop(loop_inst));
+
current_flow_block_ = if_inst->Merge();
EmitBlock(stmt->body);
@@ -669,10 +657,14 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get());
- if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
- if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
current_flow_block_->Instructions().Push(if_inst);
+ current_flow_block_ = if_inst->True();
+ SetBranch(builder_.ExitIf(if_inst));
+
+ current_flow_block_ = if_inst->False();
+ SetBranch(builder_.ExitLoop(loop_inst));
+
current_flow_block_ = if_inst->Merge();
}
@@ -719,12 +711,14 @@
current_flow_block_ = builder_.CreateCase(switch_inst, selectors);
EmitBlock(c->Body()->Declaration());
- BranchToIfNeeded(switch_inst->Merge());
+ if (NeedBranch()) {
+ SetBranch(builder_.ExitSwitch(switch_inst));
+ }
}
}
current_flow_block_ = nullptr;
- if (IsConnected(switch_inst->Merge(), 1)) {
+ if (IsConnected(switch_inst->Merge())) {
current_flow_block_ = switch_inst->Merge();
}
}
@@ -746,9 +740,9 @@
TINT_ASSERT(IR, current_control);
if (auto* c = current_control->As<Loop>()) {
- BranchTo(c->Merge());
+ SetBranch(builder_.ExitLoop(c));
} else if (auto* s = current_control->As<Switch>()) {
- BranchTo(s->Merge());
+ SetBranch(builder_.ExitSwitch(s));
} else {
TINT_UNREACHABLE(IR, diagnostics_);
}
@@ -964,14 +958,14 @@
// If the lhs is false, then that is the result we want to pass to the merge
// block as our argument
current_flow_block_ = if_inst->False();
- BranchTo(if_inst->Merge(), std::move(alt_args));
+ SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
current_flow_block_ = if_inst->True();
} else {
// If the lhs is true, then that is the result we want to pass to the merge
// block as our argument
current_flow_block_ = if_inst->True();
- BranchTo(if_inst->Merge(), std::move(alt_args));
+ SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
current_flow_block_ = if_inst->False();
}
@@ -983,7 +977,7 @@
utils::Vector<Value*, 1> args;
args.Push(rhs.Get());
- BranchTo(if_inst->Merge(), std::move(args));
+ SetBranch(builder_.ExitIf(if_inst, std::move(args)));
}
current_flow_block_ = if_inst->Merge();
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 1a81f56..8cc5604 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -60,13 +60,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = add %3, 1u
store %v1, %4
@@ -87,13 +84,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = add %3, 1u
store %v1, %4
@@ -137,13 +131,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, i32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:i32 = sub %3, 1i
store %v1, %4
@@ -164,13 +155,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = sub %3, 1u
store %v1, %4
@@ -214,13 +202,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = mul %3, 1u
store %v1, %4
@@ -264,13 +249,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = div %3, 1u
store %v1, %4
@@ -314,13 +296,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = mod %3, 1u
store %v1, %4
@@ -364,13 +343,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, bool, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:bool = and %3, false
store %v1, %4
@@ -414,13 +390,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, bool, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:bool = or %3, false
store %v1, %4
@@ -464,13 +437,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = xor %3, 1u
store %v1, %4
@@ -499,12 +469,12 @@
if %3 [t: %b3, f: %b4, m: %b5]
# True block
%b3 = block {
- br %b5 false
+ exit_if %b5 false
}
# False block
%b4 = block {
- br %b5 %3
+ exit_if %b5 %3
}
# Merge block
@@ -512,12 +482,12 @@
if %4:bool [t: %b6, f: %b7, m: %b8]
# True block
%b6 = block {
- br %b8
+ exit_if %b8
}
# False block
%b7 = block {
- br %b8
+ exit_if %b8
}
# Merge block
@@ -527,9 +497,7 @@
}
-
}
-
}
)");
}
@@ -553,12 +521,12 @@
if %3 [t: %b3, f: %b4, m: %b5]
# True block
%b3 = block {
- br %b5 %3
+ exit_if %b5 %3
}
# False block
%b4 = block {
- br %b5 true
+ exit_if %b5 true
}
# Merge block
@@ -566,12 +534,12 @@
if %4:bool [t: %b6, f: %b7, m: %b8]
# True block
%b6 = block {
- br %b8
+ exit_if %b8
}
# False block
%b7 = block {
- br %b8
+ exit_if %b8
}
# Merge block
@@ -581,9 +549,7 @@
}
-
}
-
}
)");
}
@@ -760,13 +726,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = shiftl %3, 1u
store %v1, %4
@@ -810,13 +773,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v1:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 = load %v1
%4:u32 = shiftr %3, 1u
store %v1, %4
@@ -853,12 +813,12 @@
%7:f32 = mul 2.29999995231628417969f, %6
%8:f32 = div %5, %7
%9:bool = gt 2.5f, %8
- br %b5 %9
+ exit_if %b5 %9
}
# False block
%b4 = block {
- br %b5 %4
+ exit_if %b5 %4
}
# Merge block
@@ -867,7 +827,6 @@
}
}
-
}
)");
}
diff --git a/src/tint/ir/from_program_builtin_test.cc b/src/tint/ir/from_program_builtin_test.cc
index 3993fbb..acdbc8a 100644
--- a/src/tint/ir/from_program_builtin_test.cc
+++ b/src/tint/ir/from_program_builtin_test.cc
@@ -37,13 +37,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, f32, read_write> = var, 1.0f
- br %b2 # root_end
}
-%b2 = root_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 = load %i
%tint_symbol:f32 = asin %3
ret
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 43a9709..155e42c 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -102,13 +102,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, i32, read_write> = var, 1i
- br %b2 # root_end
}
-%b2 = root_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 = load %i
%tint_symbol:f32 = convert i32, %3
ret
@@ -127,11 +124,8 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
- br %b2 # root_end
}
-%b2 = root_terminator
-
)");
}
@@ -146,13 +140,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%i:ptr<private, f32, read_write> = var, 1.0f
- br %b2 # root_end
}
-%b2 = root_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 = load %i
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
ret
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index 2d3f3d7..38efea0 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -38,13 +38,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%a:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 {
store %a, 4u
ret
}
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index e0f161d..f2b9c19 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -141,7 +141,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -149,12 +149,12 @@
if true [t: %b2, f: %b3, m: %b4]
# True block
%b2 = block {
- br %b4
+ exit_if %b4
}
# False block
%b3 = block {
- br %b4
+ exit_if %b4
}
# Merge block
@@ -163,7 +163,6 @@
}
}
-
}
)");
}
@@ -182,7 +181,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -192,9 +191,10 @@
%b2 = block {
ret
}
+
# False block
%b3 = block {
- br %b4
+ exit_if %b4
}
# Merge block
@@ -203,7 +203,6 @@
}
}
-
}
)");
}
@@ -222,7 +221,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -230,20 +229,20 @@
if true [t: %b2, f: %b3, m: %b4]
# True block
%b2 = block {
- br %b4
+ exit_if %b4
}
# False block
%b3 = block {
ret
}
+
# Merge block
%b4 = block {
ret
}
}
-
}
)");
}
@@ -262,7 +261,7 @@
EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -272,13 +271,13 @@
%b2 = block {
ret
}
+
# False block
%b3 = block {
ret
}
}
-
}
)");
}
@@ -306,20 +305,19 @@
%b2 = block {
loop [s: %b5, m: %b6]
%b5 = block {
- br %b6
+ exit_loop %b6
}
# Merge block
%b6 = block {
- br %b4
+ exit_if %b4
}
-
}
# False block
%b3 = block {
- br %b4
+ exit_if %b4
}
# Merge block
@@ -328,7 +326,6 @@
}
}
-
}
)");
}
@@ -345,7 +342,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@@ -354,7 +351,7 @@
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
- br %b3
+ exit_loop %b3
}
# Merge block
@@ -363,7 +360,6 @@
}
}
-
}
)");
}
@@ -383,12 +379,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length());
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, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -398,12 +394,12 @@
if true [t: %b5, f: %b6, m: %b7]
# True block
%b5 = block {
- br %b4
+ exit_loop %b4
}
# False block
%b6 = block {
- br %b7
+ exit_if %b7
}
# Merge block
@@ -411,7 +407,6 @@
continue %b3
}
-
}
# Continuing block
@@ -425,7 +420,6 @@
}
}
-
}
)");
}
@@ -443,7 +437,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
@@ -466,7 +460,6 @@
}
}
-
}
)");
}
@@ -500,7 +493,6 @@
}
}
-
}
)");
}
@@ -519,12 +511,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
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, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -536,9 +528,10 @@
%b5 = block {
ret
}
+
# False block
%b6 = block {
- br %b7
+ exit_if %b7
}
# Merge block
@@ -546,7 +539,6 @@
continue %b3
}
-
}
# Continuing block
@@ -560,7 +552,6 @@
}
}
-
}
)");
}
@@ -577,7 +568,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
@@ -590,7 +581,6 @@
}
}
-
}
)");
}
@@ -616,7 +606,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
@@ -629,7 +619,6 @@
}
}
-
}
)");
}
@@ -648,12 +637,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length());
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(0u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -663,15 +652,14 @@
if true [t: %b4, f: %b5]
# True block
%b4 = block {
- br %b3
+ exit_loop %b3
}
# False block
%b5 = block {
- br %b3
+ exit_loop %b3
}
-
}
# Merge block
@@ -680,7 +668,6 @@
}
}
-
}
)");
}
@@ -712,12 +699,12 @@
if true [t: %b8, f: %b9, m: %b10]
# True block
%b8 = block {
- br %b7
+ exit_loop %b7
}
# False block
%b9 = block {
- br %b10
+ exit_if %b10
}
# Merge block
@@ -730,7 +717,7 @@
# False block
%b12 = block {
- br %b13
+ exit_if %b13
}
# Merge block
@@ -738,17 +725,15 @@
continue %b6
}
-
}
-
}
# Continuing block
%b6 = block {
loop [s: %b14, m: %b15]
%b14 = block {
- br %b15
+ exit_loop %b15
}
# Merge block
@@ -768,10 +753,8 @@
break_if false %b5
}
-
}
-
}
# Merge block
@@ -779,12 +762,12 @@
if true [t: %b19, f: %b20, m: %b21]
# True block
%b19 = block {
- br %b4
+ exit_loop %b4
}
# False block
%b20 = block {
- br %b21
+ exit_if %b21
}
# Merge block
@@ -792,10 +775,8 @@
continue %b3
}
-
}
-
}
# Continuing block
@@ -809,7 +790,6 @@
}
}
-
}
)");
}
@@ -830,12 +810,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
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, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -845,12 +825,12 @@
if false [t: %b5, f: %b6, m: %b7]
# True block
%b5 = block {
- br %b7
+ exit_if %b7
}
# False block
%b6 = block {
- br %b4
+ exit_loop %b4
}
# Merge block
@@ -858,7 +838,6 @@
continue %b3
}
-
}
# Continuing block
@@ -872,7 +851,6 @@
}
}
-
}
)");
}
@@ -893,12 +871,12 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
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, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -908,12 +886,12 @@
if true [t: %b5, f: %b6, m: %b7]
# True block
%b5 = block {
- br %b7
+ exit_if %b7
}
# False block
%b6 = block {
- br %b4
+ exit_loop %b4
}
# Merge block
@@ -934,7 +912,6 @@
}
}
-
}
)");
}
@@ -973,7 +950,7 @@
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
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, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), R"()");
}
@@ -990,7 +967,7 @@
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@@ -999,7 +976,7 @@
%b1 = block {
loop [s: %b2, m: %b3]
%b2 = block {
- br %b3
+ exit_loop %b3
}
# Merge block
@@ -1008,7 +985,6 @@
}
}
-
}
)");
}
@@ -1047,7 +1023,7 @@
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
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(3u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1055,17 +1031,17 @@
switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4), m: %b5]
# Case block
%b2 = block {
- br %b5
+ exit_switch %b5
}
# Case block
%b3 = block {
- br %b5
+ exit_switch %b5
}
# Case block
%b4 = block {
- br %b5
+ exit_switch %b5
}
# Merge block
@@ -1074,7 +1050,6 @@
}
}
-
}
)");
}
@@ -1109,7 +1084,7 @@
EXPECT_TRUE(cases[0].selectors[2].IsDefault());
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1117,7 +1092,7 @@
switch 1i [c: (0i 1i default, %b2), m: %b3]
# Case block
%b2 = block {
- br %b3
+ exit_switch %b3
}
# Merge block
@@ -1126,7 +1101,6 @@
}
}
-
}
)");
}
@@ -1149,7 +1123,7 @@
EXPECT_TRUE(cases[0].selectors[0].IsDefault());
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1157,7 +1131,7 @@
switch 1i [c: (default, %b2), m: %b3]
# Case block
%b2 = block {
- br %b3
+ exit_switch %b3
}
# Merge block
@@ -1166,7 +1140,6 @@
}
}
-
}
)");
}
@@ -1197,7 +1170,7 @@
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
- EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
// This is 1 because the if is dead-code eliminated and the return doesn't happen.
EXPECT_EQ(Disassemble(m),
@@ -1206,12 +1179,12 @@
switch 1i [c: (0i, %b2), c: (default, %b3), m: %b4]
# Case block
%b2 = block {
- br %b4
+ exit_switch %b4
}
# Case block
%b3 = block {
- br %b4
+ exit_switch %b4
}
# Merge block
@@ -1220,7 +1193,6 @@
}
}
-
}
)");
}
@@ -1254,7 +1226,7 @@
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(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@@ -1264,13 +1236,13 @@
%b2 = block {
ret
}
+
# Case block
%b3 = block {
ret
}
}
-
}
)");
}
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index 774be8c..bb58c02 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -107,13 +107,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v2:ptr<private, i32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 {
ret
}
}
@@ -134,13 +131,10 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%v3:ptr<private, i32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_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 {
store %v3, 42i
ret
}
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index 29377e4..6cecf47 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -35,11 +35,8 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%a:ptr<private, u32, read_write> = var
- br %b2 # root_end
}
-%b2 = root_terminator
-
)");
}
@@ -53,11 +50,8 @@
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block {
%a:ptr<private, u32, read_write> = var, 2u
- br %b2 # root_end
}
-%b2 = root_terminator
-
)");
}
diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc
index a89a51f..2069c6d 100644
--- a/src/tint/ir/if.cc
+++ b/src/tint/ir/if.cc
@@ -19,7 +19,7 @@
namespace tint::ir {
If::If(Value* cond, Block* t, Block* f, Block* m)
- : Base(m), condition_(cond), true_(t), false_(f), merge_(m) {
+ : Base(utils::Empty), condition_(cond), true_(t), false_(f), merge_(m) {
TINT_ASSERT(IR, true_);
TINT_ASSERT(IR, false_);
TINT_ASSERT(IR, merge_);
diff --git a/src/tint/ir/loop.cc b/src/tint/ir/loop.cc
index 0bbb710..fe34283 100644
--- a/src/tint/ir/loop.cc
+++ b/src/tint/ir/loop.cc
@@ -18,7 +18,8 @@
namespace tint::ir {
-Loop::Loop(Block* s, Block* c, Block* m) : Base(s), start_(s), continuing_(c), merge_(m) {
+Loop::Loop(Block* s, Block* c, Block* m)
+ : Base(utils::Empty), start_(s), continuing_(c), merge_(m) {
TINT_ASSERT(IR, start_);
TINT_ASSERT(IR, continuing_);
TINT_ASSERT(IR, merge_);
diff --git a/src/tint/ir/root_terminator.h b/src/tint/ir/root_terminator.h
deleted file mode 100644
index 4a52b32..0000000
--- a/src/tint/ir/root_terminator.h
+++ /dev/null
@@ -1,32 +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_ROOT_TERMINATOR_H_
-#define SRC_TINT_IR_ROOT_TERMINATOR_H_
-
-#include "src/tint/ir/block.h"
-
-namespace tint::ir {
-
-/// Block used as the end of a root block. There are no instructions in this block.
-class RootTerminator : public utils::Castable<RootTerminator, Block> {
- public:
- /// Constructor
- RootTerminator();
- ~RootTerminator() override;
-};
-
-} // namespace tint::ir
-
-#endif // SRC_TINT_IR_ROOT_TERMINATOR_H_
diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc
index a28666a..003feb2 100644
--- a/src/tint/ir/switch.cc
+++ b/src/tint/ir/switch.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-Switch::Switch(Value* cond, Block* m) : Base(m), condition_(cond), merge_(m) {
+Switch::Switch(Value* cond, Block* m) : Base(utils::Empty), condition_(cond), merge_(m) {
TINT_ASSERT(IR, condition_);
TINT_ASSERT(IR, merge_);
condition_->AddUsage(this);
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index af4a957..08436f3 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -20,6 +20,7 @@
#include "src/tint/ir/block.h"
#include "src/tint/ir/call.h"
#include "src/tint/ir/constant.h"
+#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/load.h"
@@ -121,46 +122,27 @@
while (block) {
TINT_ASSERT(IR, block->HasBranchTarget());
- enum Status { kContinue, kStop, kError };
-
- Status status = tint::Switch(
- block,
-
- [&](const ir::Block* blk) {
- for (auto* inst : blk->Instructions()) {
- auto stmt = Stmt(inst);
- if (TINT_UNLIKELY(!stmt)) {
- return kError;
- }
- if (auto* s = stmt.Get()) {
- stmts.Push(s);
- }
- }
- if (auto* if_ = blk->Branch()->As<ir::If>()) {
- if (if_->Merge()->HasBranchTarget()) {
- block = if_->Merge();
- return kContinue;
- }
- } else if (auto* switch_ = blk->Branch()->As<ir::Switch>()) {
- if (switch_->Merge()->HasBranchTarget()) {
- block = switch_->Merge();
- return kContinue;
- }
- }
- return kStop;
- },
-
- [&](Default) {
- UNHANDLED_CASE(block);
- return kError;
- });
-
- if (TINT_UNLIKELY(status == kError)) {
- return nullptr;
+ for (auto* inst : block->Instructions()) {
+ auto stmt = Stmt(inst);
+ if (TINT_UNLIKELY(!stmt)) {
+ return nullptr;
+ }
+ if (auto* s = stmt.Get()) {
+ stmts.Push(s);
+ }
}
- if (status == kStop) {
- break;
+ if (auto* if_ = block->Branch()->As<ir::If>()) {
+ if (if_->Merge()->HasBranchTarget()) {
+ block = if_->Merge();
+ continue;
+ }
+ } else if (auto* switch_ = block->Branch()->As<ir::Switch>()) {
+ if (switch_->Merge()->HasBranchTarget()) {
+ block = switch_->Merge();
+ continue;
+ }
}
+ break;
}
return b.Block(std::move(stmts));
@@ -174,16 +156,20 @@
return nullptr;
}
- if (!IsEmpty(i->False(), i->Merge())) {
+ auto* false_blk = i->False();
+ if (false_blk->Instructions().Length() > 1 ||
+ (false_blk->Instructions().Length() == 1 && false_blk->HasBranchTarget() &&
+ !false_blk->Branch()->Is<ir::ExitIf>())) {
// If the else target is an `if` which has a merge target that just bounces to the outer
// if merge target then emit an 'else if' instead of a block statement for the else.
- if (auto* inst = i->False()->Instructions().Front()->As<ir::If>();
- inst && inst->Merge()->IsTrampoline(i->Merge())) {
- auto* f = If(inst);
- if (!f) {
- return nullptr;
+ if (auto* inst = i->False()->Instructions().Front()->As<ir::If>()) {
+ if (auto* br = inst->Merge()->Branch()->As<ir::ExitIf>(); br && br->If() == i) {
+ auto* f = If(inst);
+ if (!f) {
+ return nullptr;
+ }
+ return b.If(cond, t, b.Else(f));
}
- return b.If(cond, t, b.Else(f));
} else {
auto* f = BlockGraph(i->False());
if (!f) {
@@ -192,7 +178,6 @@
return b.If(cond, t, b.Else(f));
}
}
-
return b.If(cond, t);
}
@@ -265,17 +250,6 @@
return b.Return(val);
}
- /// @return true if there are no instructions between @p node and and @p stop_at
- bool IsEmpty(const ir::Block* node, const ir::Block* stop_at) {
- if (node->Instructions().IsEmpty()) {
- return true;
- }
- if (auto* br = node->Instructions().Front()->As<Branch>()) {
- return !br->Is<ir::Return>() && br->To() == stop_at;
- }
- return false;
- }
-
utils::Result<const ast::Statement*> Stmt(const ir::Instruction* inst) {
return tint::Switch<utils::Result<const ast::Statement*>>(
inst, //
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
index 29c89e6..7da87c1 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -19,6 +19,7 @@
#include "spirv/unified1/spirv.h"
#include "src/tint/ir/binary.h"
#include "src/tint/ir/block.h"
+#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/load.h"
#include "src/tint/ir/module.h"
@@ -354,25 +355,24 @@
}
void GeneratorImplIr::EmitBranch(const ir::Branch* b) {
- if (b->Is<ir::Return>()) {
- 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::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); },
+ tint::Switch( //
+ b, //
+ [&](const ir::Return*) {
+ 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;
+ },
+ [&](const ir::ExitIf* if_) {
+ current_function_.push_inst(spv::Op::OpBranch, {Label(if_->If()->Merge())});
+ },
[&](Default) {
- // A block may not have an outward branch (e.g. an unreachable merge
- // block).
- current_function_.push_inst(spv::Op::OpUnreachable, {});
+ TINT_ICE(Writer, diagnostics_) << "unimplemented branch: " << b->TypeInfo().name;
});
}
@@ -388,10 +388,12 @@
uint32_t merge_label = Label(merge_block);
uint32_t true_label = merge_label;
uint32_t false_label = merge_label;
- if (true_block->Instructions().Length() > 1 || true_block->Branch()->To() != merge_block) {
+ if (true_block->Instructions().Length() > 1 ||
+ (true_block->HasBranchTarget() && !true_block->Branch()->Is<ir::ExitIf>())) {
true_label = Label(true_block);
}
- if (false_block->Instructions().Length() > 1 || false_block->Branch()->To() != merge_block) {
+ if (false_block->Instructions().Length() > 1 ||
+ (false_block->HasBranchTarget() && !false_block->Branch()->Is<ir::ExitIf>())) {
false_label = Label(false_block);
}
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 b526c91..f43a09a 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
@@ -23,8 +23,8 @@
auto* func = b.CreateFunction("foo", mod.Types().void_());
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->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
+ i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
func->StartTarget()->SetInstructions(utils::Vector{i});
@@ -49,12 +49,12 @@
auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true));
- i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+ i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
auto* true_block = i->True();
- true_block->SetInstructions(utils::Vector{
- b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
+ true_block->SetInstructions(
+ utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
func->StartTarget()->SetInstructions(utils::Vector{i});
@@ -83,12 +83,12 @@
auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true));
- i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())});
+ i->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
auto* false_block = i->False();
- false_block->SetInstructions(utils::Vector{
- b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())});
+ false_block->SetInstructions(
+ utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
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 d66df98..574a0b0 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
@@ -100,7 +100,7 @@
v->SetInitializer(b.Constant(42_i));
auto* i = b.CreateIf(b.Constant(true));
- i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())});
+ i->True()->SetInstructions(utils::Vector{v, b.ExitIf(i)});
i->False()->SetInstructions(utils::Vector{b.Return(func)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
