From 68b4e6460ffba5e8f19d4886b31293838f4b75f8 Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Tue, 23 May 2023 22:26:24 +0000 Subject: [PATCH] [ir] Dissolve the flow graph The ir::Value objects each have a list of instructions in which their used. These lists allow us to determine all the places the value is used. Currently this is unable to track the usage of a value in an `if` or `switch` condition. It is also unable to track the usage of a value as a branch argument. In order to facilitate this tracking, the flow graph has been resolved. Branches are moved to branch instructions (and jump instructions). A jump is walk continue branch. A branch is a walk terminating branch. The `if`, `switch` and `loop` flow nodes are moved to instructions as well. Bug: tint:1718 Change-Id: I8e4cc4688bb1bdd5c7eecc72d366e6531ec685b3 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/133840 Kokoro: Kokoro Commit-Queue: Dan Sinclair Reviewed-by: James Price --- src/tint/BUILD.gn | 4 + src/tint/CMakeLists.txt | 4 + src/tint/ir/block.cc | 9 - src/tint/ir/block.h | 36 +- src/tint/ir/branch.cc | 35 + src/tint/ir/branch.h | 33 +- src/tint/ir/builder.cc | 46 +- src/tint/ir/builder.h | 13 + src/tint/ir/debug.cc | 72 +- src/tint/ir/disassembler.cc | 499 +++-- src/tint/ir/disassembler.h | 14 +- src/tint/ir/flow_node.h | 13 +- src/tint/ir/from_program.cc | 362 ++-- src/tint/ir/from_program_binary_test.cc | 743 ++++---- src/tint/ir/from_program_call_test.cc | 85 +- src/tint/ir/from_program_materialize_test.cc | 9 +- src/tint/ir/from_program_store_test.cc | 13 +- src/tint/ir/from_program_test.cc | 1610 ++++++++--------- src/tint/ir/from_program_unary_test.cc | 90 +- src/tint/ir/from_program_var_test.cc | 26 +- src/tint/ir/if.cc | 11 +- src/tint/ir/if.h | 33 +- src/tint/ir/jump.cc | 25 + src/tint/ir/jump.h | 37 + src/tint/ir/loop.cc | 6 +- src/tint/ir/loop.h | 26 +- src/tint/ir/switch.cc | 6 +- src/tint/ir/switch.h | 22 +- src/tint/ir/to_program.cc | 159 +- src/tint/ir/to_program_roundtrip_test.cc | 3 +- .../ir/transform/add_empty_entry_point.cc | 2 +- .../transform/add_empty_entry_point_test.cc | 20 +- src/tint/ir/unary.h | 2 + src/tint/ir/unary_test.cc | 6 +- src/tint/transform/manager_test.cc | 4 +- src/tint/writer/spirv/ir/generator_impl_ir.cc | 51 +- src/tint/writer/spirv/ir/generator_impl_ir.h | 5 + .../spirv/ir/generator_impl_ir_binary_test.cc | 45 +- .../ir/generator_impl_ir_function_test.cc | 16 +- .../spirv/ir/generator_impl_ir_if_test.cc | 41 +- .../spirv/ir/generator_impl_ir_var_test.cc | 28 +- 41 files changed, 2088 insertions(+), 2176 deletions(-) create mode 100644 src/tint/ir/branch.cc create mode 100644 src/tint/ir/jump.cc create mode 100644 src/tint/ir/jump.h diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index a704487592..b14851f8c6 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -1210,6 +1210,8 @@ if (tint_build_ir) { "ir/block.h", "ir/block_param.cc", "ir/block_param.h", + "ir/branch.cc", + "ir/branch.h", "ir/builder.cc", "ir/builder.h", "ir/builtin.cc", @@ -1240,6 +1242,8 @@ if (tint_build_ir) { "ir/if.h", "ir/instruction.cc", "ir/instruction.h", + "ir/jump.cc", + "ir/jump.h", "ir/load.cc", "ir/load.h", "ir/loop.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index 8ec573a376..979c4342cb 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -718,6 +718,8 @@ if(${TINT_BUILD_IR}) ir/block.h ir/block_param.cc ir/block_param.h + ir/branch.cc + ir/branch.h ir/builder.cc ir/builder.h ir/builtin.cc @@ -750,6 +752,8 @@ if(${TINT_BUILD_IR}) ir/if.h ir/instruction.cc ir/instruction.h + ir/jump.cc + ir/jump.h ir/load.cc ir/load.h ir/loop.cc diff --git a/src/tint/ir/block.cc b/src/tint/ir/block.cc index f5b5bfb077..2030c19c66 100644 --- a/src/tint/ir/block.cc +++ b/src/tint/ir/block.cc @@ -22,13 +22,4 @@ Block::Block() : Base() {} Block::~Block() = default; -void Block::BranchTo(FlowNode* to, utils::VectorRef args) { - TINT_ASSERT(IR, to); - branch_.target = to; - branch_.args = args; - if (to) { - to->AddInboundBranch(this); - } -} - } // namespace tint::ir diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h index ce85162b5e..abdf2a3277 100644 --- a/src/tint/ir/block.h +++ b/src/tint/ir/block.h @@ -34,16 +34,30 @@ class Block : public utils::Castable { Block(); ~Block() override; - /// Sets the blocks branch target to the given node. - /// @param to the node to branch too - /// @param args the branch arguments - void BranchTo(FlowNode* to, utils::VectorRef args = {}); - /// @returns true if this is block has a branch target set - bool HasBranchTarget() const override { return branch_.target != nullptr; } + bool HasBranchTarget() const override { + return !instructions_.IsEmpty() && instructions_.Back()->Is(); + } - /// @return the node this block branches too. - const ir::Branch& Branch() const { return branch_; } + /// @return the node this block branches to or nullptr if the block doesn't branch + const ir::Branch* Branch() const { + if (!HasBranchTarget()) { + return nullptr; + } + return instructions_.Back()->As(); + } + + /// @param target the block to see if we trampoline too + /// @returns if this block just branches to the provided target. + bool IsTrampoline(const FlowNode* target) const { + if (instructions_.Length() != 1) { + return false; + } + if (auto* inst = instructions_.Front()->As()) { + return inst->To() == target; + } + return false; + } /// Sets the instructions in the block /// @param instructions the instructions to set @@ -59,14 +73,12 @@ class Block : public utils::Castable { /// Sets the params to the block /// @param params the params for the block void SetParams(utils::VectorRef params) { params_ = std::move(params); } + /// @return the parameters passed into the block + utils::VectorRef Params() const { return params_; } /// @returns the params to the block utils::Vector& Params() { return params_; } - /// @return the parameters passed into the block - utils::VectorRef Params() const { return params_; } - private: - ir::Branch branch_ = {}; utils::Vector instructions_; utils::Vector params_; }; diff --git a/src/tint/ir/branch.cc b/src/tint/ir/branch.cc new file mode 100644 index 0000000000..a16b7ae95f --- /dev/null +++ b/src/tint/ir/branch.cc @@ -0,0 +1,35 @@ +// 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. + +#include "src/tint/ir/branch.h" + +#include + +#include "src/tint/ir/flow_node.h" + +TINT_INSTANTIATE_TYPEINFO(tint::ir::Branch); + +namespace tint::ir { + +Branch::Branch(FlowNode* to, utils::VectorRef args) : to_(to), args_(std::move(args)) { + TINT_ASSERT(IR, to_); + to_->AddInboundBranch(this); + for (auto* arg : args) { + arg->AddUsage(this); + } +} + +Branch::~Branch() = default; + +} // namespace tint::ir diff --git a/src/tint/ir/branch.h b/src/tint/ir/branch.h index 667f131f87..fcb12563ef 100644 --- a/src/tint/ir/branch.h +++ b/src/tint/ir/branch.h @@ -15,20 +15,35 @@ #ifndef SRC_TINT_IR_BRANCH_H_ #define SRC_TINT_IR_BRANCH_H_ -#include "src/tint/ir/flow_node.h" +#include "src/tint/ir/instruction.h" #include "src/tint/ir/value.h" +#include "src/tint/utils/castable.h" + +// Forward declarations +namespace tint::ir { +class FlowNode; +} // namespace tint::ir namespace tint::ir { -/// A information on a branch to another block -struct Branch { - /// The block being branched too. - FlowNode* target = nullptr; +/// A branch instruction. A branch is a walk terminating jump. +class Branch : public utils::Castable { + public: + /// Constructor + /// @param to the block to branch too + /// @param args the branch arguments + explicit Branch(FlowNode* to, utils::VectorRef args = {}); + ~Branch() override; - /// The arguments provided for that branch. These arguments could be the - /// 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; + /// @returns the block being branched too. + const FlowNode* To() const { return to_; } + + /// @returns the branch arguments + utils::VectorRef Args() const { return args_; } + + private: + FlowNode* to_; + utils::Vector args_; }; } // namespace tint::ir diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc index b0b04d676d..b502de191c 100644 --- a/src/tint/ir/builder.cc +++ b/src/tint/ir/builder.cc @@ -29,10 +29,6 @@ Builder::~Builder() = default; ir::Block* Builder::CreateRootBlockIfNeeded() { if (!ir.root_block) { 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 root terminator for the root-block now. - ir.root_block->BranchTo(CreateRootTerminator()); } return ir.root_block; } @@ -59,50 +55,26 @@ Function* Builder::CreateFunction(Symbol name, ir_func->SetStartTarget(CreateBlock()); ir_func->SetEndTarget(CreateFunctionTerminator()); - // Function is always branching into the Start().target - ir_func->StartTarget()->AddInboundBranch(ir_func); - return ir_func; } If* Builder::CreateIf(Value* condition) { TINT_ASSERT(IR, condition); - - auto* ir_if = ir.flow_nodes.Create(condition); - ir_if->True().target = CreateBlock(); - ir_if->False().target = CreateBlock(); - ir_if->Merge().target = CreateBlock(); - - // An if always branches to both the true and false block. - ir_if->True().target->AddInboundBranch(ir_if); - ir_if->False().target->AddInboundBranch(ir_if); - - return ir_if; + return ir.values.Create(condition, CreateBlock(), CreateBlock(), CreateBlock()); } Loop* Builder::CreateLoop() { - auto* ir_loop = ir.flow_nodes.Create(); - ir_loop->Start().target = CreateBlock(); - ir_loop->Continuing().target = CreateBlock(); - ir_loop->Merge().target = CreateBlock(); - - // A loop always branches to the start block. - ir_loop->Start().target->AddInboundBranch(ir_loop); - - return ir_loop; + return ir.values.Create(CreateBlock(), CreateBlock(), CreateBlock()); } Switch* Builder::CreateSwitch(Value* condition) { - auto* ir_switch = ir.flow_nodes.Create(condition); - ir_switch->Merge().target = CreateBlock(); - return ir_switch; + return ir.values.Create(condition, CreateBlock()); } Block* Builder::CreateCase(Switch* s, utils::VectorRef selectors) { - s->Cases().Push(Switch::Case{selectors, {CreateBlock(), utils::Empty}}); + s->Cases().Push(Switch::Case{std::move(selectors), CreateBlock()}); - Block* b = s->Cases().Back().Start().target->As(); - // Switch branches into the case block + Block* b = s->Cases().Back().Start(); b->AddInboundBranch(s); return b; } @@ -238,6 +210,14 @@ ir::Var* Builder::Declare(const type::Type* type) { return ir.values.Create(type); } +ir::Branch* Builder::Branch(FlowNode* to, utils::VectorRef args) { + return ir.values.Create(to, args); +} + +ir::Jump* Builder::Jump(FlowNode* to, utils::VectorRef args) { + return ir.values.Create(to, args); +} + ir::BlockParam* Builder::BlockParam(const type::Type* type) { return ir.values.Create(type); } diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h index ebd7a87280..ebbab7ba18 100644 --- a/src/tint/ir/builder.h +++ b/src/tint/ir/builder.h @@ -30,6 +30,7 @@ #include "src/tint/ir/function_param.h" #include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" +#include "src/tint/ir/jump.h" #include "src/tint/ir/load.h" #include "src/tint/ir/loop.h" #include "src/tint/ir/module.h" @@ -351,6 +352,18 @@ class Builder { /// @returns the instruction ir::Var* Declare(const type::Type* type); + /// Creates a branch declaration + /// @param to the node being branched too + /// @param args the branch arguments + /// @returns the instruction + ir::Branch* Branch(FlowNode* to, utils::VectorRef args = {}); + + /// Creates a jump declaration + /// @param to the node being branched too + /// @param args the branch arguments + /// @returns the instruction + ir::Jump* Jump(FlowNode* to, utils::VectorRef args = {}); + /// Creates a new `BlockParam` /// @param type the parameter type /// @returns the value diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc index 85666ff628..7e8155df64 100644 --- a/src/tint/ir/debug.cc +++ b/src/tint/ir/debug.cc @@ -60,81 +60,15 @@ std::string Debug::AsDotGraph(const Module* mod) { if (node_to_name.count(b) == 0) { out << name_for(b) << R"( [label="block"])" << std::endl; } - out << name_for(b) << " -> " << name_for(b->Branch().target); + out << name_for(b) << " -> " << name_for(b->Branch()->To()); // Dashed lines to merge blocks - if (merge_nodes.count(b->Branch().target) != 0) { + if (merge_nodes.count(b->Branch()->To()) != 0) { out << " [style=dashed]"; } out << std::endl; - Graph(b->Branch().target); - }, - [&](const ir::Switch* s) { - out << name_for(s) << R"( [label="switch"])" << std::endl; - out << name_for(s->Merge().target) << R"( [label="switch merge"])" << std::endl; - merge_nodes.insert(s->Merge().target); - - size_t i = 0; - for (const auto& c : s->Cases()) { - out << name_for(c.Start().target) - << R"( [label="case )" + std::to_string(i++) + R"("])" << std::endl; - } - out << name_for(s) << " -> {"; - for (const auto& c : s->Cases()) { - if (&c != &(s->Cases().Front())) { - out << ", "; - } - out << name_for(c.Start().target); - } - out << "}" << std::endl; - - for (const auto& c : s->Cases()) { - Graph(c.Start().target); - } - Graph(s->Merge().target); - }, - [&](const ir::If* i) { - out << name_for(i) << R"( [label="if"])" << std::endl; - out << name_for(i->True().target) << R"( [label="true"])" << std::endl; - out << name_for(i->False().target) << R"( [label="false"])" << std::endl; - out << name_for(i->Merge().target) << R"( [label="if merge"])" << std::endl; - merge_nodes.insert(i->Merge().target); - - out << name_for(i) << " -> {"; - out << name_for(i->True().target) << ", " << name_for(i->False().target); - out << "}" << std::endl; - - // Subgraph if true/false branches so they draw on the same line - out << "subgraph sub_" << name_for(i) << " {" << std::endl; - out << R"(rank="same")" << std::endl; - out << name_for(i->True().target) << std::endl; - out << name_for(i->False().target) << std::endl; - out << "}" << std::endl; - - Graph(i->True().target); - Graph(i->False().target); - Graph(i->Merge().target); - }, - [&](const ir::Loop* l) { - out << name_for(l) << R"( [label="loop"])" << std::endl; - out << name_for(l->Start().target) << R"( [label="start"])" << std::endl; - out << name_for(l->Continuing().target) << R"( [label="continuing"])" << std::endl; - out << name_for(l->Merge().target) << R"( [label="loop merge"])" << std::endl; - merge_nodes.insert(l->Merge().target); - - // Subgraph the continuing and merge so they get drawn on the same line - out << "subgraph sub_" << name_for(l) << " {" << std::endl; - out << R"(rank="same")" << std::endl; - out << name_for(l->Continuing().target) << std::endl; - out << name_for(l->Merge().target) << std::endl; - out << "}" << std::endl; - - out << name_for(l) << " -> " << name_for(l->Start().target) << std::endl; - - Graph(l->Start().target); - Graph(l->Continuing().target); - Graph(l->Merge().target); + Graph(b->Branch()->To()); }, [&](const ir::FunctionTerminator*) { // Already done diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc index 76af74091b..8eb0974f38 100644 --- a/src/tint/ir/disassembler.cc +++ b/src/tint/ir/disassembler.cc @@ -27,6 +27,7 @@ #include "src/tint/ir/discard.h" #include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" +#include "src/tint/ir/jump.h" #include "src/tint/ir/load.h" #include "src/tint/ir/loop.h" #include "src/tint/ir/root_terminator.h" @@ -41,22 +42,6 @@ namespace tint::ir { namespace { -class ScopedStopNode { - static constexpr size_t N = 32; - - public: - ScopedStopNode(utils::Hashset& stop_nodes, const FlowNode* node) - : stop_nodes_(stop_nodes), node_(node) { - stop_nodes_.Add(node_); - } - - ~ScopedStopNode() { stop_nodes_.Remove(node_); } - - private: - utils::Hashset& stop_nodes_; - const FlowNode* node_; -}; - class ScopedIndent { public: explicit ScopedIndent(uint32_t& indent) : indent_(indent) { indent_ += 2; } @@ -103,261 +88,115 @@ std::string_view Disassembler::IdOf(const Value* value) { }); } -void Disassembler::Walk(const FlowNode* node) { - if (visited_.Contains(node) || stop_nodes_.Contains(node)) { - return; - } - visited_.Add(node); - - tint::Switch( - node, - [&](const ir::Function* f) { - TINT_SCOPED_ASSIGNMENT(in_function_, true); - - Indent() << "%fn" << IdOf(f) << " = func " << f->Name().Name() << "("; - for (auto* p : f->Params()) { - if (p != f->Params().Front()) { - out_ << ", "; - } - out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName(); - } - out_ << "):" << f->ReturnType()->FriendlyName(); - - if (f->Stage() != Function::PipelineStage::kUndefined) { - out_ << " [@" << f->Stage(); - - if (f->WorkgroupSize()) { - auto arr = f->WorkgroupSize().value(); - out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] - << ")"; - } - - if (!f->ReturnAttributes().IsEmpty()) { - out_ << " ra:"; - - for (auto attr : f->ReturnAttributes()) { - out_ << " @" << attr; - if (attr == Function::ReturnAttribute::kLocation) { - out_ << "(" << f->ReturnLocation().value() << ")"; - } - } - } - - out_ << "]"; - } - out_ << " {" << std::endl; - - { - ScopedIndent func_indent(indent_size_); - ScopedStopNode scope(stop_nodes_, f->EndTarget()); - Walk(f->StartTarget()); - } - out_ << "} "; - Walk(f->EndTarget()); - }, - [&](const ir::Block* b) { - // If this block is dead, nothing to do - if (!b->HasBranchTarget()) { - return; - } - - Indent() << "%fn" << IdOf(b) << " = block"; - if (!b->Params().IsEmpty()) { - out_ << " ("; - for (const auto* p : b->Params()) { - if (p != b->Params().Front()) { - out_ << ", "; - } - EmitValue(p); - } - out_ << ")"; - } - - out_ << " {" << std::endl; - { - ScopedIndent si(indent_size_); - EmitBlockInstructions(b); - } - Indent() << "}"; - - std::string suffix = ""; - if (b->Branch().target->Is()) { - out_ << " -> %func_end"; - suffix = "return"; - } else if (b->Branch().target->Is()) { - // Nothing to do - } else { - out_ << " -> " - << "%fn" << IdOf(b->Branch().target); - suffix = "branch"; - } - if (!b->Branch().args.IsEmpty()) { - out_ << " "; - for (const auto* v : b->Branch().args) { - if (v != b->Branch().args.Front()) { - out_ << ", "; - } - EmitValue(v); - } - } - if (!suffix.empty()) { - out_ << " # " << suffix; - } - out_ << std::endl; - - if (!b->Branch().target->Is()) { - out_ << std::endl; - } - - Walk(b->Branch().target); - }, - [&](const ir::Switch* s) { - Indent() << "%fn" << IdOf(s) << " = switch "; - EmitValue(s->Condition()); - out_ << " ["; - for (const auto& c : s->Cases()) { - if (&c != &s->Cases().Front()) { - out_ << ", "; - } - out_ << "c: ("; - for (const auto& selector : c.selectors) { - if (&selector != &c.selectors.Front()) { - out_ << " "; - } - - if (selector.IsDefault()) { - out_ << "default"; - } else { - EmitValue(selector.val); - } - } - out_ << ", %fn" << IdOf(c.Start().target) << ")"; - } - if (s->Merge().target->IsConnected()) { - out_ << ", m: %fn" << IdOf(s->Merge().target); - } - out_ << "]" << std::endl; - - { - ScopedIndent switch_indent(indent_size_); - ScopedStopNode scope(stop_nodes_, s->Merge().target); - for (const auto& c : s->Cases()) { - Indent() << "# case "; - for (const auto& selector : c.selectors) { - if (&selector != &c.selectors.Front()) { - out_ << " "; - } - - if (selector.IsDefault()) { - out_ << "default"; - } else { - EmitValue(selector.val); - } - } - out_ << std::endl; - Walk(c.Start().target); - } - } - - if (s->Merge().target->IsConnected()) { - Indent() << "# switch merge" << std::endl; - Walk(s->Merge().target); - } - }, - [&](const ir::If* i) { - Indent() << "%fn" << IdOf(i) << " = if "; - EmitValue(i->Condition()); - - bool has_true = i->True().target->HasBranchTarget(); - bool has_false = i->False().target->HasBranchTarget(); - - out_ << " ["; - if (has_true) { - out_ << "t: %fn" << IdOf(i->True().target); - } - if (has_false) { - if (has_true) { - out_ << ", "; - } - out_ << "f: %fn" << IdOf(i->False().target); - } - if (i->Merge().target->IsConnected()) { - out_ << ", m: %fn" << IdOf(i->Merge().target); - } - out_ << "]" << std::endl; - - { - ScopedIndent if_indent(indent_size_); - ScopedStopNode scope(stop_nodes_, i->Merge().target); - - if (has_true) { - Indent() << "# true branch" << std::endl; - Walk(i->True().target); - } - - if (has_false) { - Indent() << "# false branch" << std::endl; - Walk(i->False().target); - } - } - - if (i->Merge().target->IsConnected()) { - Indent() << "# if merge" << std::endl; - Walk(i->Merge().target); - } - }, - [&](const ir::Loop* l) { - Indent() << "%fn" << IdOf(l) << " = loop [s: %fn" << IdOf(l->Start().target); - - if (l->Continuing().target->IsConnected()) { - out_ << ", c: %fn" << IdOf(l->Continuing().target); - } - if (l->Merge().target->IsConnected()) { - out_ << ", m: %fn" << IdOf(l->Merge().target); - } - out_ << "]" << std::endl; - - { - ScopedStopNode loop_scope(stop_nodes_, l->Merge().target); - ScopedIndent loop_indent(indent_size_); - { - ScopedStopNode inner_scope(stop_nodes_, l->Continuing().target); - Indent() << "# loop start" << std::endl; - Walk(l->Start().target); - } - - if (l->Continuing().target->IsConnected()) { - Indent() << "# loop continuing" << std::endl; - Walk(l->Continuing().target); - } - } - - if (l->Merge().target->IsConnected()) { - Indent() << "# loop merge" << std::endl; - Walk(l->Merge().target); - } - }, - [&](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; - }); -} - std::string Disassembler::Disassemble() { if (mod_.root_block) { - Walk(mod_.root_block); + walk_list_.push_back(mod_.root_block); + Walk(); + TINT_ASSERT(IR, walk_list_.empty()); } - for (const auto* func : mod_.functions) { - Walk(func); + for (auto* func : mod_.functions) { + walk_list_.push_back(func); + Walk(); + TINT_ASSERT(IR, walk_list_.empty()); } return out_.str(); } +void Disassembler::Walk() { + utils::Hashset visited_; + + while (!walk_list_.empty()) { + const FlowNode* node = walk_list_.front(); + walk_list_.pop_front(); + + if (visited_.Contains(node)) { + continue; + } + visited_.Add(node); + + tint::Switch( + node, + [&](const ir::Function* f) { + in_function_ = true; + + Indent() << "%fn" << IdOf(f) << " = func " << f->Name().Name() << "("; + for (auto* p : f->Params()) { + if (p != f->Params().Front()) { + out_ << ", "; + } + out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName(); + } + out_ << "):" << f->ReturnType()->FriendlyName(); + + if (f->Stage() != Function::PipelineStage::kUndefined) { + out_ << " [@" << f->Stage(); + + if (f->WorkgroupSize()) { + auto arr = f->WorkgroupSize().value(); + out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] + << ")"; + } + + if (!f->ReturnAttributes().IsEmpty()) { + out_ << " ra:"; + + for (auto attr : f->ReturnAttributes()) { + out_ << " @" << attr; + if (attr == Function::ReturnAttribute::kLocation) { + out_ << "(" << f->ReturnLocation().value() << ")"; + } + } + } + + out_ << "]"; + } + out_ << " -> %fn" << IdOf(f->StartTarget()) << std::endl; + walk_list_.push_back(f->StartTarget()); + }, + [&](const ir::Block* b) { + // If this block is dead, nothing to do + if (!b->HasBranchTarget()) { + return; + } + + Indent() << "%fn" << IdOf(b) << " = block"; + if (!b->Params().IsEmpty()) { + out_ << " ("; + for (auto* p : b->Params()) { + if (p != b->Params().Front()) { + out_ << ", "; + } + EmitValue(p); + } + out_ << ")"; + } + + out_ << " {" << std::endl; + { + ScopedIndent si(indent_size_); + EmitBlockInstructions(b); + } + Indent() << "}" << std::endl; + + if (!b->Branch()->To()->Is()) { + out_ << std::endl; + } + + walk_list_.push_back(b->Branch()->To()); + }, + [&](const ir::FunctionTerminator* t) { + TINT_ASSERT(IR, in_function_); + Indent() << "%fn" << IdOf(t) << " = func_terminator" << std::endl << std::endl; + in_function_ = false; + }, + [&](const ir::RootTerminator* t) { + TINT_ASSERT(IR, !in_function_); + Indent() << "%fn" << IdOf(t) << " = root_terminator" << std::endl << std::endl; + }); + } +} + void Disassembler::EmitValueWithType(const Value* val) { EmitValue(val); if (auto* i = val->As(); i->Type() != nullptr) { @@ -419,8 +258,12 @@ void Disassembler::EmitValue(const Value* val) { void Disassembler::EmitInstruction(const Instruction* inst) { tint::Switch( - inst, // - [&](const ir::Binary* b) { EmitBinary(b); }, [&](const ir::Unary* u) { EmitUnary(u); }, + inst, // + [&](const ir::Switch* s) { EmitSwitch(s); }, // + [&](const ir::If* i) { EmitIf(i); }, // + [&](const ir::Loop* l) { EmitLoop(l); }, // + [&](const ir::Binary* b) { EmitBinary(b); }, // + [&](const ir::Unary* u) { EmitUnary(u); }, [&](const ir::Bitcast* b) { EmitValueWithType(b); out_ << " = bitcast "; @@ -468,7 +311,131 @@ void Disassembler::EmitInstruction(const Instruction* inst) { out_ << ", "; EmitValue(v->Initializer()); } - }); + }, + [&](const ir::Branch* b) { EmitBranch(b); }, + [&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; }); +} + +void Disassembler::EmitIf(const If* i) { + out_ << "if "; + EmitValue(i->Condition()); + + bool has_true = i->True()->HasBranchTarget(); + bool has_false = i->False()->HasBranchTarget(); + + out_ << " ["; + if (has_true) { + out_ << "t: %fn" << IdOf(i->True()); + } + if (has_false) { + if (has_true) { + out_ << ", "; + } + out_ << "f: %fn" << IdOf(i->False()); + } + if (i->Merge()->IsConnected()) { + out_ << ", m: %fn" << IdOf(i->Merge()); + } + out_ << "]"; + + if (has_true) { + walk_list_.push_back(i->True()); + } + if (has_false) { + walk_list_.push_back(i->False()); + } + if (i->Merge()->IsConnected()) { + walk_list_.push_back(i->Merge()); + } +} + +void Disassembler::EmitLoop(const Loop* l) { + out_ << "loop [s: %fn" << IdOf(l->Start()); + + if (l->Continuing()->IsConnected()) { + out_ << ", c: %fn" << IdOf(l->Continuing()); + } + if (l->Merge()->IsConnected()) { + out_ << ", m: %fn" << IdOf(l->Merge()); + } + out_ << "]"; + + { walk_list_.push_back(l->Start()); } + + if (l->Continuing()->IsConnected()) { + walk_list_.push_back(l->Continuing()); + } + if (l->Merge()->IsConnected()) { + walk_list_.push_back(l->Merge()); + } +} + +void Disassembler::EmitSwitch(const Switch* s) { + out_ << "switch "; + EmitValue(s->Condition()); + out_ << " ["; + for (const auto& c : s->Cases()) { + if (&c != &s->Cases().Front()) { + out_ << ", "; + } + out_ << "c: ("; + for (const auto& selector : c.selectors) { + if (&selector != &c.selectors.Front()) { + out_ << " "; + } + + if (selector.IsDefault()) { + out_ << "default"; + } else { + EmitValue(selector.val); + } + } + out_ << ", %fn" << IdOf(c.Start()) << ")"; + } + if (s->Merge()->IsConnected()) { + out_ << ", m: %fn" << IdOf(s->Merge()); + } + out_ << "]"; + + for (auto& c : s->Cases()) { + walk_list_.push_back(c.Start()); + } + if (s->Merge()->IsConnected()) { + walk_list_.push_back(s->Merge()); + } +} + +void Disassembler::EmitBranch(const Branch* b) { + if (b->Is()) { + out_ << "jmp "; + + // Stuff the thing we're jumping too into the front of the walk list so it will be emitted + // next. + walk_list_.push_front(b->To()); + } else { + out_ << "br "; + } + + std::string suffix = ""; + out_ << "%fn" << IdOf(b->To()); + if (b->To()->Is()) { + suffix = "return"; + } else if (b->To()->Is()) { + suffix = "root_end"; + } + + if (!b->Args().IsEmpty()) { + out_ << " "; + for (auto* v : b->Args()) { + if (v != b->Args().Front()) { + out_ << ", "; + } + EmitValue(v); + } + } + if (!suffix.empty()) { + out_ << " # " << suffix; + } } void Disassembler::EmitArgs(const Call* call) { diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h index c8953dba82..7b9e4c575e 100644 --- a/src/tint/ir/disassembler.h +++ b/src/tint/ir/disassembler.h @@ -15,15 +15,18 @@ #ifndef SRC_TINT_IR_DISASSEMBLER_H_ #define SRC_TINT_IR_DISASSEMBLER_H_ +#include #include #include "src/tint/ir/binary.h" #include "src/tint/ir/call.h" #include "src/tint/ir/flow_node.h" +#include "src/tint/ir/if.h" +#include "src/tint/ir/loop.h" #include "src/tint/ir/module.h" +#include "src/tint/ir/switch.h" #include "src/tint/ir/unary.h" #include "src/tint/utils/hashmap.h" -#include "src/tint/utils/hashset.h" #include "src/tint/utils/string_stream.h" namespace tint::ir { @@ -53,18 +56,21 @@ class Disassembler { size_t IdOf(const FlowNode* node); std::string_view IdOf(const Value* node); - void Walk(const FlowNode* node); + void Walk(); void EmitInstruction(const Instruction* inst); void EmitValueWithType(const Value* val); void EmitValue(const Value* val); void EmitArgs(const Call* call); void EmitBinary(const Binary* b); void EmitUnary(const Unary* b); + void EmitBranch(const Branch* b); + void EmitSwitch(const Switch* s); + void EmitLoop(const Loop* l); + void EmitIf(const If* i); const Module& mod_; utils::StringStream out_; - utils::Hashset visited_; - utils::Hashset stop_nodes_; + std::deque walk_list_; utils::Hashmap flow_node_ids_; utils::Hashmap value_ids_; uint32_t indent_size_ = 0; diff --git a/src/tint/ir/flow_node.h b/src/tint/ir/flow_node.h index 289873b1c5..b0729640d1 100644 --- a/src/tint/ir/flow_node.h +++ b/src/tint/ir/flow_node.h @@ -18,6 +18,11 @@ #include "src/tint/utils/castable.h" #include "src/tint/utils/vector.h" +// Forward Declarations +namespace tint::ir { +class Branch; +} // namespace tint::ir + namespace tint::ir { /// Base class for flow nodes @@ -26,17 +31,17 @@ class FlowNode : public utils::Castable { ~FlowNode() override; /// @returns true if this node has inbound branches and branches out - bool IsConnected() const { return HasBranchTarget() && !inbound_branches_.IsEmpty(); } + bool IsConnected() const { return HasBranchTarget(); } /// @returns true if the node has a branch target virtual bool HasBranchTarget() const { return false; } /// @returns the inbound branch list for the flow node - utils::VectorRef InboundBranches() const { return inbound_branches_; } + utils::VectorRef InboundBranches() const { return inbound_branches_; } /// Adds the given node to the inbound branches /// @param node the node to add - void AddInboundBranch(FlowNode* node) { inbound_branches_.Push(node); } + void AddInboundBranch(Branch* node) { inbound_branches_.Push(node); } protected: /// Constructor @@ -48,7 +53,7 @@ class FlowNode : public utils::Castable { /// - Node is a start node /// - Node is a merge target outside control flow (e.g. an if that returns in both branches) /// - Node is a continue target outside control flow (e.g. a loop that returns) - utils::Vector inbound_branches_; + utils::Vector inbound_branches_; }; } // namespace tint::ir diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc index 595db584d7..212efea613 100644 --- a/src/tint/ir/from_program.cc +++ b/src/tint/ir/from_program.cc @@ -98,19 +98,15 @@ namespace { using ResultType = utils::Result; -bool IsConnected(const FlowNode* b) { +// 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 FlowNode* b, uint32_t count) { // Function is always connected as it's the start. if (b->Is()) { return true; } - - for (auto* parent : b->InboundBranches()) { - if (IsConnected(parent)) { - return true; - } - } - // Getting here means all the incoming branches are disconnected. - return false; + return b->InboundBranches().Length() > count; } /// Impl is the private-implementation of FromProgram(). @@ -145,8 +141,8 @@ class Impl { /* dst */ {&builder_.ir.constants_arena}, }; - /// The stack of flow control blocks. - utils::Vector flow_stack_; + /// The stack of control blocks. + utils::Vector control_stack_; /// The current flow block for expressions. Block* current_flow_block_ = nullptr; @@ -160,15 +156,11 @@ class Impl { /// The diagnostic that have been raised. diag::List diagnostics_; - /// Map from ast nodes to flow nodes, used to retrieve the flow node for a given AST node. - /// Used for testing purposes. - std::unordered_map ast_to_flow_; - - class FlowStackScope { + class ControlStackScope { public: - FlowStackScope(Impl* impl, FlowNode* node) : impl_(impl) { impl_->flow_stack_.Push(node); } + ControlStackScope(Impl* impl, Branch* b) : impl_(impl) { impl_->control_stack_.Push(b); } - ~FlowStackScope() { impl_->flow_stack_.Pop(); } + ~ControlStackScope() { impl_->control_stack_.Pop(); } private: Impl* impl_; @@ -178,11 +170,25 @@ class Impl { diagnostics_.add_error(tint::diag::System::IR, err, s); } + void JumpTo(FlowNode* node, utils::VectorRef args = {}) { + TINT_ASSERT(IR, current_flow_block_); + TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget()); + + current_flow_block_->Instructions().Push(builder_.Jump(node, args)); + current_flow_block_ = nullptr; + } + void JumpToIfNeeded(FlowNode* node) { + if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) { + return; + } + JumpTo(node); + } + void BranchTo(FlowNode* node, utils::VectorRef args = {}) { TINT_ASSERT(IR, current_flow_block_); TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget()); - current_flow_block_->BranchTo(node, args); + current_flow_block_->Instructions().Push(builder_.Branch(node, args)); current_flow_block_ = nullptr; } @@ -193,8 +199,8 @@ class Impl { BranchTo(node); } - FlowNode* FindEnclosingControl(ControlFlags flags) { - for (auto it = flow_stack_.rbegin(); it != flow_stack_.rend(); ++it) { + Branch* FindEnclosingControl(ControlFlags flags) { + for (auto it = control_stack_.rbegin(); it != control_stack_.rend(); ++it) { if ((*it)->Is()) { return *it; } @@ -244,6 +250,11 @@ class Impl { }); } + // 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_)); } @@ -253,7 +264,7 @@ class Impl { void EmitFunction(const ast::Function* ast_func) { // The flow stack should have been emptied when the previous function finished building. - TINT_ASSERT(IR, flow_stack_.IsEmpty()); + TINT_ASSERT(IR, control_stack_.IsEmpty()); const auto* sem = program_->Sem().Get(ast_func); @@ -262,8 +273,6 @@ class Impl { current_function_ = ir_func; builder_.ir.functions.Push(ir_func); - ast_to_flow_[ast_func] = ir_func; - if (ast_func->IsEntryPoint()) { switch (ast_func->PipelineStage()) { case ast::PipelineStage::kVertex: @@ -343,17 +352,15 @@ class Impl { ir_func->SetParams(params); { - FlowStackScope scope(this, ir_func); - current_flow_block_ = ir_func->StartTarget(); EmitBlock(ast_func->body); - // If the branch target has already been set then a `return` was called. Only set in the - // case where `return` wasn't called. - BranchToIfNeeded(current_function_->EndTarget()); + // If the branch target has already been set then a `return` was called. Only set in + // the case where `return` wasn't called. + JumpToIfNeeded(current_function_->EndTarget()); } - TINT_ASSERT(IR, flow_stack_.IsEmpty()); + TINT_ASSERT(IR, control_stack_.IsEmpty()); current_flow_block_ = nullptr; current_function_ = nullptr; } @@ -362,8 +369,8 @@ class Impl { for (auto* s : stmts) { EmitStatement(s); - // If the current flow block has a branch target then the rest of the statements in this - // block are dead code. Skip them. + // If the current flow block has a branch target then the rest of the statements in + // this block are dead code. Skip them. if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) { break; } @@ -399,11 +406,11 @@ class Impl { } void EmitAssignment(const ast::AssignmentStatement* stmt) { - // If assigning to a phony, just generate the RHS and we're done. Note that, because this - // isn't used, a subsequent transform could remove it due to it being dead code. This could - // then change the interface for the program (i.e. a global var no longer used). If that - // happens we have to either fix this to store to a phony value, or make sure we pull the - // interface before doing the dead code elimination. + // If assigning to a phony, just generate the RHS and we're done. Note that, because + // this isn't used, a subsequent transform could remove it due to it being dead code. + // This could then change the interface for the program (i.e. a global var no longer + // used). If that happens we have to either fix this to store to a phony value, or make + // sure we pull the interface before doing the dead code elimination. if (stmt->lhs->Is()) { (void)EmitExpression(stmt->rhs); return; @@ -523,8 +530,8 @@ class Impl { TINT_DEFER(scopes_.Pop()); // Note, this doesn't need to emit a Block as the current block flow node should be - // sufficient as the blocks all get flattened. Each flow control node will inject the basic - // blocks it requires. + // sufficient as the blocks all get flattened. Each flow control node will inject the + // basic blocks it requires. EmitStatements(block->statements); } @@ -534,50 +541,43 @@ class Impl { if (!reg) { return; } - auto* if_node = builder_.CreateIf(reg.Get()); - - BranchTo(if_node); - - ast_to_flow_[stmt] = if_node; + auto* if_inst = builder_.CreateIf(reg.Get()); + current_flow_block_->Instructions().Push(if_inst); { - FlowStackScope scope(this, if_node); + ControlStackScope scope(this, if_inst); - current_flow_block_ = if_node->True().target->As(); + current_flow_block_ = if_inst->True(); EmitBlock(stmt->body); // If the true branch did not execute control flow, then go to the Merge().target - BranchToIfNeeded(if_node->Merge().target); + BranchToIfNeeded(if_inst->Merge()); - current_flow_block_ = if_node->False().target->As(); + current_flow_block_ = if_inst->False(); if (stmt->else_statement) { EmitStatement(stmt->else_statement); } // If the false branch did not execute control flow, then go to the Merge().target - BranchToIfNeeded(if_node->Merge().target); + BranchToIfNeeded(if_inst->Merge()); } 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_node->Merge().target)) { - current_flow_block_ = if_node->Merge().target->As(); + // 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)) { + current_flow_block_ = if_inst->Merge(); } } void EmitLoop(const ast::LoopStatement* stmt) { - auto* loop_node = builder_.CreateLoop(); - - BranchTo(loop_node); - - ast_to_flow_[stmt] = loop_node; + auto* loop_inst = builder_.CreateLoop(); + current_flow_block_->Instructions().Push(loop_inst); { - FlowStackScope scope(this, loop_node); - - current_flow_block_ = loop_node->Start().target->As(); + ControlStackScope scope(this, loop_inst); + current_flow_block_ = loop_inst->Start(); // The loop doesn't use EmitBlock because it needs the scope stack to not get popped // until after the continuing block. @@ -585,41 +585,39 @@ class Impl { TINT_DEFER(scopes_.Pop()); EmitStatements(stmt->body->statements); - // The current block didn't `break`, `return` or `continue`, go to the continuing block. - BranchToIfNeeded(loop_node->Continuing().target); + // The current block didn't `break`, `return` or `continue`, go to the continuing + // block. + JumpToIfNeeded(loop_inst->Continuing()); - current_flow_block_ = loop_node->Continuing().target->As(); + current_flow_block_ = loop_inst->Continuing(); if (stmt->continuing) { EmitBlock(stmt->continuing); } // Branch back to the start node if the continue target didn't branch out already - BranchToIfNeeded(loop_node->Start().target); + BranchToIfNeeded(loop_inst->Start()); } // The loop merge can get disconnected if the loop returns directly, or the continuing // target branches, eventually, to the merge, but nothing branched to the // Continuing().target. - current_flow_block_ = loop_node->Merge().target->As(); - if (!IsConnected(loop_node->Merge().target)) { + current_flow_block_ = loop_inst->Merge(); + if (!IsConnected(loop_inst->Merge(), 0)) { current_flow_block_ = nullptr; } } void EmitWhile(const ast::WhileStatement* stmt) { - auto* loop_node = builder_.CreateLoop(); + auto* loop_inst = builder_.CreateLoop(); + current_flow_block_->Instructions().Push(loop_inst); + // Continue is always empty, just go back to the start - TINT_ASSERT(IR, loop_node->Continuing().target->Is()); - loop_node->Continuing().target->As()->BranchTo(loop_node->Start().target); - - BranchTo(loop_node); - - ast_to_flow_[stmt] = loop_node; + loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start())); { - FlowStackScope scope(this, loop_node); + ControlStackScope scope(this, loop_inst); - current_flow_block_ = loop_node->Start().target->As(); + current_flow_block_ = loop_inst->Start(); // Emit the while condition into the Start().target of the loop auto reg = EmitExpression(stmt->condition); @@ -628,25 +626,26 @@ class Impl { } // Create an `if (cond) {} else {break;}` control flow - auto* if_node = builder_.CreateIf(reg.Get()); - if_node->True().target->As()->BranchTo(if_node->Merge().target); - if_node->False().target->As()->BranchTo(loop_node->Merge().target); + 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); - BranchTo(if_node); - - current_flow_block_ = if_node->Merge().target->As(); + current_flow_block_ = if_inst->Merge(); EmitBlock(stmt->body); - BranchToIfNeeded(loop_node->Continuing().target); + JumpToIfNeeded(loop_inst->Continuing()); } // The while loop always has a path to the Merge().target as the break statement comes // before anything inside the loop. - current_flow_block_ = loop_node->Merge().target->As(); + current_flow_block_ = loop_inst->Merge(); } void EmitForLoop(const ast::ForLoopStatement* stmt) { - auto* loop_node = builder_.CreateLoop(); - loop_node->Continuing().target->As()->BranchTo(loop_node->Start().target); + auto* loop_inst = builder_.CreateLoop(); + current_flow_block_->Instructions().Push(loop_inst); + + loop_inst->Continuing()->Instructions().Push(builder_.Branch(loop_inst->Start())); // Make sure the initializer ends up in a contained scope scopes_.Push(); @@ -657,14 +656,10 @@ class Impl { EmitStatement(stmt->initializer); } - BranchTo(loop_node); - - ast_to_flow_[stmt] = loop_node; - { - FlowStackScope scope(this, loop_node); + ControlStackScope scope(this, loop_inst); - current_flow_block_ = loop_node->Start().target->As(); + current_flow_block_ = loop_inst->Start(); if (stmt->condition) { // Emit the condition into the target target of the loop @@ -674,26 +669,26 @@ class Impl { } // Create an `if (cond) {} else {break;}` control flow - auto* if_node = builder_.CreateIf(reg.Get()); - if_node->True().target->As()->BranchTo(if_node->Merge().target); - if_node->False().target->As()->BranchTo(loop_node->Merge().target); + 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); - BranchTo(if_node); - current_flow_block_ = if_node->Merge().target->As(); + current_flow_block_ = if_inst->Merge(); } EmitBlock(stmt->body); - BranchToIfNeeded(loop_node->Continuing().target); + JumpToIfNeeded(loop_inst->Continuing()); if (stmt->continuing) { - current_flow_block_ = loop_node->Continuing().target->As(); + current_flow_block_ = loop_inst->Continuing(); EmitStatement(stmt->continuing); } } // The while loop always has a path to the Merge().target as the break statement comes // before anything inside the loop. - current_flow_block_ = loop_node->Merge().target->As(); + current_flow_block_ = loop_inst->Merge(); } void EmitSwitch(const ast::SwitchStatement* stmt) { @@ -702,14 +697,11 @@ class Impl { if (!reg) { return; } - auto* switch_node = builder_.CreateSwitch(reg.Get()); - - BranchTo(switch_node); - - ast_to_flow_[stmt] = switch_node; + auto* switch_inst = builder_.CreateSwitch(reg.Get()); + current_flow_block_->Instructions().Push(switch_inst); { - FlowStackScope scope(this, switch_node); + ControlStackScope scope(this, switch_inst); const auto* sem = program_->Sem().Get(stmt); for (const auto* c : sem->Cases()) { @@ -722,16 +714,16 @@ class Impl { } } - current_flow_block_ = builder_.CreateCase(switch_node, selectors); + current_flow_block_ = builder_.CreateCase(switch_inst, selectors); EmitBlock(c->Body()->Declaration()); - BranchToIfNeeded(switch_node->Merge().target); + BranchToIfNeeded(switch_inst->Merge()); } } current_flow_block_ = nullptr; - if (IsConnected(switch_node->Merge().target)) { - current_flow_block_ = switch_node->Merge().target->As(); + if (IsConnected(switch_inst->Merge(), 1)) { + current_flow_block_ = switch_inst->Merge(); } } @@ -753,9 +745,9 @@ class Impl { TINT_ASSERT(IR, current_control); if (auto* c = current_control->As()) { - BranchTo(c->Merge().target); + BranchTo(c->Merge()); } else if (auto* s = current_control->As()) { - BranchTo(s->Merge().target); + BranchTo(s->Merge()); } else { TINT_UNREACHABLE(IR, diagnostics_); } @@ -766,14 +758,14 @@ class Impl { TINT_ASSERT(IR, current_control); if (auto* c = current_control->As()) { - BranchTo(c->Continuing().target); + BranchTo(c->Continuing()); } else { TINT_UNREACHABLE(IR, diagnostics_); } } - // Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so the - // code has to continue as before it just predicates writes. If WGSL grows some kind of + // Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so + // the code has to continue as before it just predicates writes. If WGSL grows some kind of // terminating discard that would probably make sense as a FlowNode but would then require // figuring out the multi-level exit that is triggered. void EmitDiscard(const ast::DiscardStatement*) { @@ -787,11 +779,8 @@ class Impl { if (!reg) { return; } - auto* if_node = builder_.CreateIf(reg.Get()); - - BranchTo(if_node); - - ast_to_flow_[stmt] = if_node; + auto* if_inst = builder_.CreateIf(reg.Get()); + current_flow_block_->Instructions().Push(if_inst); auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch); TINT_ASSERT(IR, current_control); @@ -799,17 +788,17 @@ class Impl { auto* loop = current_control->As(); - current_flow_block_ = if_node->True().target->As(); - BranchTo(loop->Merge().target); + current_flow_block_ = if_inst->True(); + BranchTo(loop->Merge()); - current_flow_block_ = if_node->False().target->As(); - BranchTo(if_node->Merge().target); + current_flow_block_ = if_inst->False(); + BranchTo(if_inst->Merge()); - current_flow_block_ = if_node->Merge().target->As(); + current_flow_block_ = if_inst->Merge(); - // The `break-if` has to be the last item in the continuing block. The false branch of the - // `break-if` will always take us back to the start of the loop. - BranchTo(loop->Start().target); + // The `break-if` has to be the last item in the continuing block. The false branch of + // the `break-if` will always take us back to the start of the loop. + BranchTo(loop->Start()); } utils::Result EmitExpression(const ast::Expression* expr) { @@ -845,8 +834,8 @@ class Impl { // TODO(dsinclair): Implement // }, [&](const ast::UnaryOpExpression* u) { return EmitUnary(u); }, - // Note, ast::PhonyExpression is explicitly not handled here as it should never get into - // this method. The assignment statement should have filtered it out already. + // Note, ast::PhonyExpression is explicitly not handled here as it should never get + // into this method. The assignment statement should have filtered it out already. [&](Default) { add_error(expr->source, "unknown expression type: " + std::string(expr->TypeInfo().name)); @@ -891,8 +880,8 @@ class Impl { builder_.ir.SetName(val, v->name->symbol.Name()); }, [&](const ast::Let* l) { - // A `let` doesn't exist as a standalone item in the IR, it's just the result of the - // initializer. + // A `let` doesn't exist as a standalone item in the IR, it's just the result of + // the initializer. auto init = EmitExpression(l->initializer); if (!init) { return; @@ -911,12 +900,12 @@ class Impl { }, [&](const ast::Const*) { // Skip. This should be handled by const-eval already, so the const will be a - // `constant::` value at the usage sites. Can just ignore the `const` variable as it - // should never be used. + // `constant::` value at the usage sites. Can just ignore the `const` variable + // as it should never be used. // - // TODO(dsinclair): Probably want to store the const variable somewhere and then in - // identifier expression log an error if we ever see a const identifier. Add this - // when identifiers and variables are supported. + // TODO(dsinclair): Probably want to store the const variable somewhere and then + // in identifier expression log an error if we ever see a const identifier. Add + // this when identifiers and variables are supported. }, [&](Default) { add_error(var->source, "unknown variable: " + std::string(var->TypeInfo().name)); @@ -953,8 +942,8 @@ class Impl { return inst; } - // A short-circut needs special treatment. The short-circuit is decomposed into the relevant if - // statements and declarations. + // A short-circut needs special treatment. The short-circuit is decomposed into the relevant + // if statements and declarations. utils::Result EmitShortCircuit(const ast::BinaryExpression* expr) { switch (expr->op) { case ast::BinaryOp::kLogicalAnd: @@ -972,15 +961,15 @@ class Impl { return utils::Failure; } - auto* if_node = builder_.CreateIf(lhs.Get()); - BranchTo(if_node); + auto* if_inst = builder_.CreateIf(lhs.Get()); + current_flow_block_->Instructions().Push(if_inst); auto* result = builder_.BlockParam(builder_.ir.types.Get()); - if_node->Merge().target->As()->SetParams(utils::Vector{result}); + if_inst->Merge()->SetParams(utils::Vector{result}); utils::Result rhs; { - FlowStackScope scope(this, if_node); + ControlStackScope scope(this, if_inst); utils::Vector alt_args; alt_args.Push(lhs.Get()); @@ -988,19 +977,19 @@ class Impl { // If this is an `&&` then we only evaluate the RHS expression in the true block. // If this is an `||` then we only evaluate the RHS expression in the false block. if (expr->op == ast::BinaryOp::kLogicalAnd) { - // 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_node->False().target->As(); - BranchTo(if_node->Merge().target, std::move(alt_args)); + // 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)); - current_flow_block_ = if_node->True().target->As(); + 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_node->True().target->As(); - BranchTo(if_node->Merge().target, std::move(alt_args)); + // 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)); - current_flow_block_ = if_node->False().target->As(); + current_flow_block_ = if_inst->False(); } rhs = EmitExpression(expr->rhs); @@ -1010,9 +999,9 @@ class Impl { utils::Vector args; args.Push(rhs.Get()); - BranchTo(if_node->Merge().target, std::move(args)); + BranchTo(if_inst->Merge(), std::move(args)); } - current_flow_block_ = if_node->Merge().target->As(); + current_flow_block_ = if_inst->Merge(); return result; } @@ -1191,67 +1180,6 @@ class Impl { } return builder_.Constant(cv); } - - // void EmitAttributes(utils::VectorRef attrs) { - // for (auto* attr : attrs) { - // EmitAttribute(attr); - // } - // } - // - // void EmitAttribute(const ast::Attribute* attr) { - // tint::Switch( // - // attr, - // [&](const ast::WorkgroupAttribute* wg) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::StageAttribute* s) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::BindingAttribute* b) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::GroupAttribute* g) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::LocationAttribute* l) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::BuiltinAttribute* b) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::InterpolateAttribute* i) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::InvariantAttribute* i) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::MustUseAttribute* i) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::IdAttribute*) { - // add_error(attr->source, - // "found an `Id` attribute. The SubstituteOverrides transform " - // "must be run before converting to IR"); - // }, - // [&](const ast::StructMemberSizeAttribute*) { - // TINT_ICE(IR, diagnostics_) - // << "StructMemberSizeAttribute encountered during IR conversion"; - // }, - // [&](const ast::StructMemberAlignAttribute*) { - // TINT_ICE(IR, diagnostics_) - // << "StructMemberAlignAttribute encountered during IR conversion"; - // }, - // [&](const ast::StrideAttribute* s) { - // // TODO(dsinclair): Implement - // }, - // [&](const ast::InternalAttribute *i) { - // // TODO(dsinclair): Implement - // }, - // [&](Default) { - // add_error(attr->source, "unknown attribute: " + - // std::string(attr->TypeInfo().name)); - // }); - // } }; } // namespace diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc index b74a8c88a0..6e63e400a8 100644 --- a/src/tint/ir/from_program_binary_test.cc +++ b/src/tint/ir/from_program_binary_test.cc @@ -34,17 +34,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Add) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = add %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = add %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -59,16 +61,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = add %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = add %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -83,16 +88,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = add %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = add %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -105,17 +113,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Subtract) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = sub %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = sub %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -130,16 +140,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:i32 = load %v1 - %3:i32 = sub %2, 1i - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:i32 = load %v1 + %3:i32 = sub %2, 1i + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -154,16 +167,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = sub %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = sub %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -176,17 +192,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Multiply) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = mul %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = mul %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -201,16 +219,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = mul %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = mul %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -223,17 +244,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Div) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = div %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = div %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -248,16 +271,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = div %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = div %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -270,17 +296,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Modulo) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = mod %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = mod %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -295,16 +323,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = mod %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = mod %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -317,17 +348,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_And) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = and %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = and %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -342,16 +375,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:bool = load %v1 - %3:bool = and %2, false - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:bool = load %v1 + %3:bool = and %2, false + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -364,17 +400,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Or) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = or %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = or %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -389,16 +427,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:bool = load %v1 - %3:bool = or %2, false - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:bool = load %v1 + %3:bool = or %2, false + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -411,17 +452,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Xor) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = xor %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = xor %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -436,16 +479,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = xor %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = xor %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -458,42 +504,42 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool { - %fn2 = block { - } -> %func_end true # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool -> %fn2 +%fn2 = block { + br %fn3 true # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:bool = call my_func - } -> %fn5 # branch +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:bool = call my_func + if %1 [t: %fn6, f: %fn7, m: %fn8] +} - %fn5 = if %1 [t: %fn6, f: %fn7, m: %fn8] - # true branch - %fn6 = block { - } -> %fn8 false # branch +%fn6 = block { + br %fn8 false +} - # false branch - %fn7 = block { - } -> %fn8 %1 # branch +%fn7 = block { + br %fn8 %1 +} - # if merge - %fn8 = block (%2:bool) { - } -> %fn9 # branch +%fn8 = block (%2:bool) { + if %2:bool [t: %fn9, f: %fn10, m: %fn11] +} - %fn9 = if %2:bool [t: %fn10, f: %fn11, m: %fn12] - # true branch - %fn10 = block { - } -> %fn12 # branch +%fn9 = block { + br %fn11 +} - # false branch - %fn11 = block { - } -> %fn12 # branch +%fn10 = block { + br %fn11 +} - # if merge - %fn12 = block { - } -> %func_end # return -} %func_end +%fn11 = block { + jmp %fn12 # return +} +%fn12 = func_terminator )"); } @@ -506,42 +552,42 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool { - %fn2 = block { - } -> %func_end true # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool -> %fn2 +%fn2 = block { + br %fn3 true # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:bool = call my_func - } -> %fn5 # branch +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:bool = call my_func + if %1 [t: %fn6, f: %fn7, m: %fn8] +} - %fn5 = if %1 [t: %fn6, f: %fn7, m: %fn8] - # true branch - %fn6 = block { - } -> %fn8 %1 # branch +%fn6 = block { + br %fn8 %1 +} - # false branch - %fn7 = block { - } -> %fn8 true # branch +%fn7 = block { + br %fn8 true +} - # if merge - %fn8 = block (%2:bool) { - } -> %fn9 # branch +%fn8 = block (%2:bool) { + if %2:bool [t: %fn9, f: %fn10, m: %fn11] +} - %fn9 = if %2:bool [t: %fn10, f: %fn11, m: %fn12] - # true branch - %fn10 = block { - } -> %fn12 # branch +%fn9 = block { + br %fn11 +} - # false branch - %fn11 = block { - } -> %fn12 # branch +%fn10 = block { + br %fn11 +} - # if merge - %fn12 = block { - } -> %func_end # return -} %func_end +%fn11 = block { + jmp %fn12 # return +} +%fn12 = func_terminator )"); } @@ -554,17 +600,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Equal) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:bool = eq %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:bool = eq %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -577,17 +625,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_NotEqual) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:bool = neq %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:bool = neq %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -600,17 +650,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThan) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:bool = lt %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:bool = lt %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -623,17 +675,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThan) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:bool = gt %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:bool = gt %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -646,17 +700,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThanEqual) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:bool = lte %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:bool = lte %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -669,17 +725,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThanEqual) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:bool = gte %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:bool = gte %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -692,17 +750,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftLeft) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = shiftl %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = shiftl %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -717,16 +777,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = shiftl %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = shiftl %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -739,17 +802,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftRight) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 0u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 0u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = shiftr %1, 4u - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = shiftr %1, 4u + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -764,16 +829,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v1:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:u32 = load %v1 - %3:u32 = shiftr %2, 1u - store %v1, %3 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:u32 = load %v1 + %3:u32 = shiftr %2, 1u + store %v1, %3 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -788,35 +856,36 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 { - %fn2 = block { - } -> %func_end 0.0f # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 -> %fn2 +%fn2 = block { + br %fn3 0.0f # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:f32 = call my_func - %2:bool = lt %1, 2.0f - } -> %fn5 # branch +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:f32 = call my_func + %2:bool = lt %1, 2.0f + if %2 [t: %fn6, f: %fn7, m: %fn8] +} - %fn5 = if %2 [t: %fn6, f: %fn7, m: %fn8] - # true branch - %fn6 = block { - %3:f32 = call my_func - %4:f32 = call my_func - %5:f32 = mul 2.29999995231628417969f, %4 - %6:f32 = div %3, %5 - %7:bool = gt 2.5f, %6 - } -> %fn8 %7 # branch +%fn6 = block { + %3:f32 = call my_func + %4:f32 = call my_func + %5:f32 = mul 2.29999995231628417969f, %4 + %6:f32 = div %3, %5 + %7:bool = gt 2.5f, %6 + br %fn8 %7 +} - # false branch - %fn7 = block { - } -> %fn8 %2 # branch +%fn7 = block { + br %fn8 %2 +} - # if merge - %fn8 = block (%tint_symbol:bool) { - } -> %func_end # return -} %func_end +%fn8 = block (%tint_symbol:bool) { + jmp %fn9 # return +} +%fn9 = func_terminator )"); } @@ -830,16 +899,18 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:bool):bool { - %fn2 = block { - } -> %func_end true # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:bool):bool -> %fn2 +%fn2 = block { + br %fn3 true # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %tint_symbol:bool = call my_func, false - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %tint_symbol:bool = call my_func, false + jmp %fn6 # return +} +%fn6 = func_terminator )"); } diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc index 6339c223fd..2f74508987 100644 --- a/src/tint/ir/from_program_call_test.cc +++ b/src/tint/ir/from_program_call_test.cc @@ -35,17 +35,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 { - %fn2 = block { - } -> %func_end 0.0f # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 -> %fn2 +%fn2 = block { + br %fn3 0.0f # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:f32 = call my_func - %tint_symbol:f32 = bitcast %1 - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:f32 = call my_func + %tint_symbol:f32 = bitcast %1 + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -60,11 +62,12 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Discard) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():void [@fragment] { - %fn2 = block { - discard - } -> %func_end # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():void [@fragment] -> %fn2 +%fn2 = block { + discard + jmp %fn3 # return +} +%fn3 = func_terminator )"); } @@ -77,16 +80,18 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void { - %fn2 = block { - } -> %func_end # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void -> %fn2 +%fn2 = block { + jmp %fn3 # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %2:void = call my_func, 6.0f - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %2:void = call my_func, 6.0f + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -101,15 +106,18 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %i:ptr = var, 1i + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:i32 = load %i - %tint_symbol:f32 = convert i32, %2 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:i32 = load %i + %tint_symbol:f32 = convert i32, %2 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -123,8 +131,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %i:ptr, read_write> = var, vec3 0.0f + br %fn2 # root_end } +%fn2 = root_terminator )"); } @@ -139,15 +149,18 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %i:ptr = var, 1.0f + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - %2:f32 = load %i - %tint_symbol:vec3 = construct 2.0f, 3.0f, %2 - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + %2:f32 = load %i + %tint_symbol:vec3 = construct 2.0f, 3.0f, %2 + jmp %fn5 # return +} +%fn5 = func_terminator )"); } diff --git a/src/tint/ir/from_program_materialize_test.cc b/src/tint/ir/from_program_materialize_test.cc index ba293a7257..009d417377 100644 --- a/src/tint/ir/from_program_materialize_test.cc +++ b/src/tint/ir/from_program_materialize_test.cc @@ -34,10 +34,11 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():f32 { - %fn2 = block { - } -> %func_end 2.0f # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():f32 -> %fn2 +%fn2 = block { + br %fn3 2.0f # return +} +%fn3 = func_terminator )"); } diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc index 5bb63988f0..ab0945b747 100644 --- a/src/tint/ir/from_program_store_test.cc +++ b/src/tint/ir/from_program_store_test.cc @@ -37,14 +37,17 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %a:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - store %a, 4u - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + store %a, 4u + jmp %fn5 # return +} +%fn5 = func_terminator )"); } diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc index 1049386c3e..197e7e38f5 100644 --- a/src/tint/ir/from_program_test.cc +++ b/src/tint/ir/from_program_test.cc @@ -32,10 +32,10 @@ namespace { /// If multiple flow nodes are found with the type T, then an error is raised and the first is /// returned. template -const T* FindSingleFlowNode(const Module& mod) { - const T* found = nullptr; +T* FindSingleValue(Module& mod) { + T* found = nullptr; size_t count = 0; - for (auto* node : mod.flow_nodes.Objects()) { + for (auto* node : mod.values.Objects()) { if (auto* as = node->As()) { count++; if (!found) { @@ -44,7 +44,7 @@ const T* FindSingleFlowNode(const Module& mod) { } } if (count > 1) { - ADD_FAILURE() << "FindSingleFlowNode() found " << count << " nodes of type " + ADD_FAILURE() << "FindSingleValue() found " << count << " nodes of type " << utils::TypeInfo::Of().name; } return found; @@ -66,15 +66,15 @@ TEST_F(IR_BuilderImplTest, Func) { ASSERT_NE(f->StartTarget(), nullptr); ASSERT_NE(f->EndTarget(), nullptr); - EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length()); EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length()); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f():void { - %fn2 = block { - } -> %func_end # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f():void -> %fn2 +%fn2 = block { + jmp %fn3 # return +} +%fn3 = func_terminator )"); } @@ -91,15 +91,15 @@ TEST_F(IR_BuilderImplTest, Func_WithParam) { ASSERT_NE(f->StartTarget(), nullptr); ASSERT_NE(f->EndTarget(), nullptr); - EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length()); EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length()); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32):u32 { - %fn2 = block { - } -> %func_end %a # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32):u32 -> %fn2 +%fn2 = block { + br %fn3 %a # return +} +%fn3 = func_terminator )"); } @@ -117,15 +117,15 @@ TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) { ASSERT_NE(f->StartTarget(), nullptr); ASSERT_NE(f->EndTarget(), nullptr); - EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length()); EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length()); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32, %b:i32, %c:bool):void { - %fn2 = block { - } -> %func_end # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32, %b:i32, %c:bool):void -> %fn2 +%fn2 = block { + jmp %fn3 # return +} +%fn3 = func_terminator )"); } @@ -144,42 +144,38 @@ TEST_F(IR_BuilderImplTest, IfStatement) { auto* ast_if = If(true, Block(), Else(Block())); WrapInFunction(ast_if); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->True().target, nullptr); - ASSERT_NE(flow->False().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->False().target->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); + EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + if true [t: %fn3, f: %fn4, m: %fn5] +} - %fn3 = if true [t: %fn4, f: %fn5, m: %fn6] - # true branch - %fn4 = block { - } -> %fn6 # branch +%fn3 = block { + br %fn5 +} - # false branch - %fn5 = block { - } -> %fn6 # branch +%fn4 = block { + br %fn5 +} - # if merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn5 = block { + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -188,41 +184,37 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) { auto* ast_if = If(true, Block(Return())); WrapInFunction(ast_if); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->True().target, nullptr); - ASSERT_NE(flow->False().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + if true [t: %fn3, f: %fn4, m: %fn5] +} - %fn3 = if true [t: %fn4, f: %fn5, m: %fn6] - # true branch - %fn4 = block { - } -> %func_end # return - # false branch - %fn5 = block { - } -> %fn6 # branch +%fn3 = block { + br %fn6 # return +} +%fn4 = block { + br %fn5 +} - # if merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn5 = block { + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -231,41 +223,37 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) { auto* ast_if = If(true, Block(), Else(Block(Return()))); WrapInFunction(ast_if); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->True().target, nullptr); - ASSERT_NE(flow->False().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + if true [t: %fn3, f: %fn4, m: %fn5] +} - %fn3 = if true [t: %fn4, f: %fn5, m: %fn6] - # true branch - %fn4 = block { - } -> %fn6 # branch +%fn3 = block { + br %fn5 +} - # false branch - %fn5 = block { - } -> %func_end # return - # if merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn4 = block { + br %fn6 # return +} +%fn5 = block { + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -274,37 +262,33 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) { auto* ast_if = If(true, Block(Return()), Else(Block(Return()))); WrapInFunction(ast_if); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->True().target, nullptr); - ASSERT_NE(flow->False().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->False().target->InboundBranches().Length()); - EXPECT_EQ(0u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + if true [t: %fn3, f: %fn4] +} - %fn3 = if true [t: %fn4, f: %fn5] - # true branch - %fn4 = block { - } -> %func_end # return - # false branch - %fn5 = block { - } -> %func_end # return -} %func_end +%fn3 = block { + br %fn5 # return +} +%fn4 = block { + br %fn5 # return +} +%fn5 = func_terminator )"); } @@ -314,47 +298,46 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) { auto* ast_if = If(true, Block(ast_loop)); WrapInFunction(ast_if); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* if_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(if_flow->True().target, nullptr); - ASSERT_NE(if_flow->False().target, nullptr); - ASSERT_NE(if_flow->Merge().target, nullptr); + auto m = res.Move(); + auto* if_flow = FindSingleValue(m); + ASSERT_NE(if_flow, nullptr); - auto* loop_flow = FindSingleFlowNode(m.Get()); + auto* loop_flow = FindSingleValue(m); ASSERT_NE(loop_flow, nullptr); - ASSERT_NE(loop_flow->Start().target, nullptr); - ASSERT_NE(loop_flow->Continuing().target, nullptr); - ASSERT_NE(loop_flow->Merge().target, nullptr); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + if true [t: %fn3, f: %fn4, m: %fn5] +} - %fn3 = if true [t: %fn4, f: %fn5, m: %fn6] - # true branch - %fn4 = block { - } -> %fn7 # branch +%fn3 = block { + loop [s: %fn6, c: %fn7, m: %fn8] +} - %fn7 = loop [s: %fn8, m: %fn9] - # loop start - %fn8 = block { - } -> %fn9 # branch +%fn4 = block { + br %fn5 +} - # loop merge - %fn9 = block { - } -> %fn6 # branch +%fn5 = block { + jmp %fn9 # return +} +%fn9 = func_terminator - # false branch - %fn5 = block { - } -> %fn6 # branch +%fn6 = block { + br %fn8 +} - # if merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn7 = block { + br %fn6 +} + +%fn8 = block { + br %fn5 +} )"); } @@ -363,38 +346,38 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) { auto* ast_loop = Loop(Block(Break())); WrapInFunction(ast_loop); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Start().target, nullptr); - ASSERT_NE(flow->Continuing().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); + EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, m: %fn5] - # loop start - %fn4 = block { - } -> %fn5 # branch +%fn3 = block { + br %fn5 +} - # loop merge - %fn5 = block { - } -> %func_end # return -} %func_end +%fn4 = block { + br %fn3 +} + +%fn5 = block { + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -404,64 +387,55 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) { auto* ast_loop = Loop(Block(ast_if, Continue())); WrapInFunction(ast_loop); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* loop_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(loop_flow->Start().target, nullptr); - ASSERT_NE(loop_flow->Continuing().target, nullptr); - ASSERT_NE(loop_flow->Merge().target, nullptr); + auto m = res.Move(); + auto* loop_flow = FindSingleValue(m); - auto* if_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(if_flow->True().target, nullptr); - ASSERT_NE(if_flow->False().target, nullptr); - ASSERT_NE(if_flow->Merge().target, nullptr); + auto* if_flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, loop_flow->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(2u, 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, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, c: %fn5, m: %fn6] - # loop start - %fn4 = block { - } -> %fn7 # branch +%fn3 = block { + if true [t: %fn6, f: %fn7, m: %fn8] +} - %fn7 = if true [t: %fn8, f: %fn9, m: %fn10] - # true branch - %fn8 = block { - } -> %fn6 # branch +%fn4 = block { + br %fn3 +} - # false branch - %fn9 = block { - } -> %fn10 # branch +%fn5 = block { + jmp %fn9 # return +} +%fn9 = func_terminator - # if merge - %fn10 = block { - } -> %fn5 # branch +%fn6 = block { + br %fn5 +} - # loop continuing - %fn5 = block { - } -> %fn4 # branch +%fn7 = block { + br %fn8 +} - # loop merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn8 = block { + br %fn4 +} )"); } @@ -471,64 +445,54 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) { auto* ast_loop = Loop(Block(), Block(ast_break_if)); WrapInFunction(ast_loop); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* loop_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(loop_flow->Start().target, nullptr); - ASSERT_NE(loop_flow->Continuing().target, nullptr); - ASSERT_NE(loop_flow->Merge().target, nullptr); + auto m = res.Move(); + auto* loop_flow = FindSingleValue(m); + auto* break_if_flow = FindSingleValue(m); - auto* break_if_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(break_if_flow->True().target, nullptr); - ASSERT_NE(break_if_flow->False().target, nullptr); - ASSERT_NE(break_if_flow->Merge().target, nullptr); + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; - - EXPECT_EQ(1u, loop_flow->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, break_if_flow->InboundBranches().Length()); - EXPECT_EQ(1u, break_if_flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, break_if_flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, break_if_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); + EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length()); + EXPECT_EQ(1u, break_if_flow->True()->InboundBranches().Length()); + EXPECT_EQ(1u, break_if_flow->False()->InboundBranches().Length()); + EXPECT_EQ(2u, break_if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, c: %fn5, m: %fn6] - # loop start - %fn4 = block { - } -> %fn5 # branch +%fn3 = block { + jmp %fn4 +} - # loop continuing - %fn5 = block { - } -> %fn7 # branch +%fn4 = block { + if true [t: %fn6, f: %fn7, m: %fn8] +} - %fn7 = if true [t: %fn8, f: %fn9, m: %fn10] - # true branch - %fn8 = block { - } -> %fn6 # branch +%fn5 = block { + jmp %fn9 # return +} +%fn9 = func_terminator - # false branch - %fn9 = block { - } -> %fn10 # branch +%fn6 = block { + br %fn5 +} - # if merge - %fn10 = block { - } -> %fn4 # branch +%fn7 = block { + br %fn8 +} - # loop merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn8 = block { + br %fn3 +} )"); } @@ -539,40 +503,40 @@ TEST_F(IR_BuilderImplTest, Loop_Continuing_Body_Scope) { auto* ast_loop = Loop(Block(a), Block(ast_break_if)); WrapInFunction(ast_loop); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + auto m = res.Move(); + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, c: %fn5, m: %fn6] - # loop start - %fn4 = block { - } -> %fn5 # branch +%fn3 = block { + jmp %fn4 +} - # loop continuing - %fn5 = block { - } -> %fn7 # branch +%fn4 = block { + if true [t: %fn6, f: %fn7, m: %fn8] +} - %fn7 = if true [t: %fn8, f: %fn9, m: %fn10] - # true branch - %fn8 = block { - } -> %fn6 # branch +%fn5 = block { + jmp %fn9 # return +} +%fn9 = func_terminator - # false branch - %fn9 = block { - } -> %fn10 # branch +%fn6 = block { + br %fn5 +} - # if merge - %fn10 = block { - } -> %fn4 # branch +%fn7 = block { + br %fn8 +} - # loop merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn8 = block { + br %fn3 +} )"); } @@ -582,60 +546,50 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) { auto* ast_loop = Loop(Block(ast_if, Continue())); WrapInFunction(ast_loop); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* loop_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(loop_flow->Start().target, nullptr); - ASSERT_NE(loop_flow->Continuing().target, nullptr); - ASSERT_NE(loop_flow->Merge().target, nullptr); + auto m = res.Move(); + auto* loop_flow = FindSingleValue(m); + auto* if_flow = FindSingleValue(m); - auto* if_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(if_flow->True().target, nullptr); - ASSERT_NE(if_flow->False().target, nullptr); - ASSERT_NE(if_flow->Merge().target, nullptr); + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; - - EXPECT_EQ(1u, loop_flow->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(0u, loop_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); + EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(0u, 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, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4] +} - %fn3 = loop [s: %fn4, c: %fn5] - # loop start - %fn4 = block { - } -> %fn6 # branch +%fn3 = block { + if true [t: %fn5, f: %fn6, m: %fn7] +} - %fn6 = if true [t: %fn7, f: %fn8, m: %fn9] - # true branch - %fn7 = block { - } -> %func_end # return - # false branch - %fn8 = block { - } -> %fn9 # branch +%fn4 = block { + br %fn3 +} - # if merge - %fn9 = block { - } -> %fn5 # branch +%fn5 = block { + br %fn8 # return +} +%fn6 = block { + br %fn7 +} - # loop continuing - %fn5 = block { - } -> %fn4 # branch +%fn7 = block { + br %fn4 +} -} %func_end +%fn8 = func_terminator )"); } @@ -644,34 +598,34 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) { auto* ast_loop = Loop(Block(Return(), Continue())); WrapInFunction(ast_loop, If(true, Block(Return()))); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* loop_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(loop_flow->Start().target, nullptr); - ASSERT_NE(loop_flow->Continuing().target, nullptr); - ASSERT_NE(loop_flow->Merge().target, nullptr); + auto m = res.Move(); + auto* loop_flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, loop_flow->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(0u, loop_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); + EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4] +} - %fn3 = loop [s: %fn4] - # loop start - %fn4 = block { - } -> %func_end # return -} %func_end +%fn3 = block { + br %fn5 # return +} +%fn4 = block { + br %fn3 +} + +%fn5 = func_terminator )"); } @@ -689,41 +643,61 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) { auto* ast_if = If(true, Block(Return())); WrapInFunction(Block(ast_loop, ast_if)); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* loop_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(loop_flow->Start().target, nullptr); - ASSERT_NE(loop_flow->Continuing().target, nullptr); - ASSERT_NE(loop_flow->Merge().target, nullptr); + auto m = res.Move(); + auto* loop_flow = FindSingleValue(m); - auto* break_if_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(break_if_flow->True().target, nullptr); - ASSERT_NE(break_if_flow->False().target, nullptr); - ASSERT_NE(break_if_flow->Merge().target, nullptr); + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); + EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length()); + EXPECT_EQ(3u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); - // This is 1 because only the loop branch happens. The subsequent if return is dead code. - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch +%fn3 = block { + br %fn6 # return +} +%fn4 = block { + if true [t: %fn7, f: %fn8, m: %fn9] +} - %fn3 = loop [s: %fn4] - # loop start - %fn4 = block { - } -> %func_end # return -} %func_end +%fn5 = block { + if true [t: %fn10, f: %fn11, m: %fn12] +} +%fn6 = func_terminator + +%fn7 = block { + br %fn5 +} + +%fn8 = block { + br %fn9 +} + +%fn9 = block { + br %fn3 +} + +%fn10 = block { + br %fn6 # return +} +%fn11 = block { + br %fn12 +} + +%fn12 = block { + jmp %fn6 # return +} )"); } @@ -732,56 +706,50 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) { auto* ast_loop = Loop(Block(ast_if, Continue())); WrapInFunction(ast_loop); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* loop_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(loop_flow->Start().target, nullptr); - ASSERT_NE(loop_flow->Continuing().target, nullptr); - ASSERT_NE(loop_flow->Merge().target, nullptr); + auto m = res.Move(); + auto* loop_flow = FindSingleValue(m); + auto* if_flow = FindSingleValue(m); - auto* if_flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(if_flow->True().target, nullptr); - ASSERT_NE(if_flow->False().target, nullptr); - ASSERT_NE(if_flow->Merge().target, nullptr); + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; - - EXPECT_EQ(1u, loop_flow->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length()); - EXPECT_EQ(0u, if_flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); + EXPECT_EQ(2u, 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(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, m: %fn5] - # loop start - %fn4 = block { - } -> %fn6 # branch +%fn3 = block { + if true [t: %fn6, f: %fn7] +} - %fn6 = if true [t: %fn7, f: %fn8] - # true branch - %fn7 = block { - } -> %fn5 # branch +%fn4 = block { + br %fn3 +} - # false branch - %fn8 = block { - } -> %fn5 # branch +%fn5 = block { + jmp %fn8 # return +} +%fn8 = func_terminator - # loop merge - %fn5 = block { - } -> %func_end # return -} %func_end +%fn6 = block { + br %fn5 +} + +%fn7 = block { + br %fn5 +} )"); } @@ -803,206 +771,108 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - ASSERT_EQ(1u, m->functions.Length()); - - auto block_exit = [&](const ir::FlowNode* node) -> const ir::FlowNode* { - if (auto* block = As(node)) { - return block->Branch().target; - } - return nullptr; - }; - - auto* loop_flow_a = As(m->functions[0]->StartTarget()->Branch().target); - ASSERT_NE(loop_flow_a, nullptr); - ASSERT_NE(loop_flow_a->Start().target, nullptr); - ASSERT_NE(loop_flow_a->Continuing().target, nullptr); - ASSERT_NE(loop_flow_a->Merge().target, nullptr); - - auto* loop_flow_b = As(block_exit(loop_flow_a->Start().target)); - ASSERT_NE(loop_flow_b, nullptr); - ASSERT_NE(loop_flow_b->Start().target, nullptr); - ASSERT_NE(loop_flow_b->Continuing().target, nullptr); - ASSERT_NE(loop_flow_b->Merge().target, nullptr); - - auto* if_flow_a = As(block_exit(loop_flow_b->Start().target)); - ASSERT_NE(if_flow_a, nullptr); - ASSERT_NE(if_flow_a->True().target, nullptr); - ASSERT_NE(if_flow_a->False().target, nullptr); - ASSERT_NE(if_flow_a->Merge().target, nullptr); - - auto* if_flow_b = As(block_exit(if_flow_a->Merge().target)); - ASSERT_NE(if_flow_b, nullptr); - ASSERT_NE(if_flow_b->True().target, nullptr); - ASSERT_NE(if_flow_b->False().target, nullptr); - ASSERT_NE(if_flow_b->Merge().target, nullptr); - - auto* loop_flow_c = As(block_exit(loop_flow_b->Continuing().target)); - ASSERT_NE(loop_flow_c, nullptr); - ASSERT_NE(loop_flow_c->Start().target, nullptr); - ASSERT_NE(loop_flow_c->Continuing().target, nullptr); - ASSERT_NE(loop_flow_c->Merge().target, nullptr); - - auto* loop_flow_d = As(block_exit(loop_flow_c->Merge().target)); - ASSERT_NE(loop_flow_d, nullptr); - ASSERT_NE(loop_flow_d->Start().target, nullptr); - ASSERT_NE(loop_flow_d->Continuing().target, nullptr); - ASSERT_NE(loop_flow_d->Merge().target, nullptr); - - auto* if_flow_c = As(block_exit(loop_flow_d->Continuing().target)); - ASSERT_NE(if_flow_c, nullptr); - ASSERT_NE(if_flow_c->True().target, nullptr); - ASSERT_NE(if_flow_c->False().target, nullptr); - ASSERT_NE(if_flow_c->Merge().target, nullptr); - - auto* if_flow_d = As(block_exit(loop_flow_b->Merge().target)); - ASSERT_NE(if_flow_d, nullptr); - ASSERT_NE(if_flow_d->True().target, nullptr); - ASSERT_NE(if_flow_d->False().target, nullptr); - ASSERT_NE(if_flow_d->Merge().target, nullptr); - - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; - - EXPECT_EQ(1u, loop_flow_a->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow_a->Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_a->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_a->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_b->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow_b->Start().target->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow_b->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_b->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_c->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow_c->Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, loop_flow_c->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_c->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_d->InboundBranches().Length()); - EXPECT_EQ(2u, loop_flow_d->Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_d->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, loop_flow_d->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_a->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_a->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_a->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_a->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_b->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_b->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_b->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_b->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_c->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_c->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_c->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_c->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_d->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_d->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_d->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow_d->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length()); - EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, c: %fn5, m: %fn6] - # loop start - %fn4 = block { - } -> %fn7 # branch +%fn3 = block { + loop [s: %fn6, c: %fn7, m: %fn8] +} - %fn7 = loop [s: %fn8, c: %fn9, m: %fn10] - # loop start - %fn8 = block { - } -> %fn11 # branch +%fn4 = block { + br %fn3 +} - %fn11 = if true [t: %fn12, f: %fn13, m: %fn14] - # true branch - %fn12 = block { - } -> %fn10 # branch +%fn5 = block { + jmp %fn9 # return +} +%fn9 = func_terminator - # false branch - %fn13 = block { - } -> %fn14 # branch +%fn6 = block { + if true [t: %fn10, f: %fn11, m: %fn12] +} - # if merge - %fn14 = block { - } -> %fn15 # branch +%fn7 = block { + loop [s: %fn13, c: %fn14, m: %fn15] +} - %fn15 = if true [t: %fn16, f: %fn17, m: %fn18] - # true branch - %fn16 = block { - } -> %fn9 # branch +%fn8 = block { + if true [t: %fn16, f: %fn17, m: %fn18] +} - # false branch - %fn17 = block { - } -> %fn18 # branch +%fn10 = block { + br %fn8 +} - # if merge - %fn18 = block { - } -> %fn9 # branch +%fn11 = block { + br %fn12 +} - # loop continuing - %fn9 = block { - } -> %fn19 # branch +%fn12 = block { + if true [t: %fn19, f: %fn20, m: %fn21] +} - %fn19 = loop [s: %fn20, m: %fn21] - # loop start - %fn20 = block { - } -> %fn21 # branch +%fn13 = block { + br %fn15 +} - # loop merge - %fn21 = block { - } -> %fn22 # branch +%fn14 = block { + br %fn13 +} - %fn22 = loop [s: %fn23, c: %fn24, m: %fn25] - # loop start - %fn23 = block { - } -> %fn24 # branch +%fn15 = block { + loop [s: %fn22, c: %fn23, m: %fn24] +} - # loop continuing - %fn24 = block { - } -> %fn26 # branch +%fn16 = block { + br %fn5 +} - %fn26 = if true [t: %fn27, f: %fn28, m: %fn29] - # true branch - %fn27 = block { - } -> %fn25 # branch +%fn17 = block { + br %fn18 +} - # false branch - %fn28 = block { - } -> %fn29 # branch +%fn18 = block { + jmp %fn4 +} - # if merge - %fn29 = block { - } -> %fn23 # branch +%fn19 = block { + br %fn7 +} - # loop merge - %fn25 = block { - } -> %fn8 # branch +%fn20 = block { + br %fn21 +} - # loop merge - %fn10 = block { - } -> %fn30 # branch +%fn21 = block { + jmp %fn7 +} - %fn30 = if true [t: %fn31, f: %fn32, m: %fn33] - # true branch - %fn31 = block { - } -> %fn6 # branch +%fn22 = block { + jmp %fn23 +} - # false branch - %fn32 = block { - } -> %fn33 # branch +%fn23 = block { + if true [t: %fn25, f: %fn26, m: %fn27] +} - # if merge - %fn33 = block { - } -> %fn5 # branch +%fn24 = block { + br %fn6 +} - # loop continuing - %fn5 = block { - } -> %fn4 # branch +%fn25 = block { + br %fn24 +} - # loop merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn26 = block { + br %fn27 +} + +%fn27 = block { + br %fn22 +} )"); } @@ -1011,64 +881,57 @@ TEST_F(IR_BuilderImplTest, While) { auto* ast_while = While(false, Block()); WrapInFunction(ast_while); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Start().target, nullptr); - ASSERT_NE(flow->Continuing().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_NE(flow->Start().target->As()->Branch().target, nullptr); - ASSERT_TRUE(flow->Start().target->As()->Branch().target->Is()); - auto* if_flow = flow->Start().target->As()->Branch().target->As(); - ASSERT_NE(if_flow->True().target, nullptr); - ASSERT_NE(if_flow->False().target, nullptr); - ASSERT_NE(if_flow->Merge().target, nullptr); + ASSERT_NE(flow->Start()->Branch(), nullptr); + ASSERT_TRUE(flow->Start()->Branch()->Is()); + auto* if_flow = flow->Start()->Branch()->As(); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(1u, 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(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, c: %fn5, m: %fn6] - # loop start - %fn4 = block { - } -> %fn7 # branch +%fn3 = block { + if false [t: %fn6, f: %fn7, m: %fn8] +} - %fn7 = if false [t: %fn8, f: %fn9, m: %fn10] - # true branch - %fn8 = block { - } -> %fn10 # branch +%fn4 = block { + br %fn3 +} - # false branch - %fn9 = block { - } -> %fn6 # branch +%fn5 = block { + jmp %fn9 # return +} +%fn9 = func_terminator - # if merge - %fn10 = block { - } -> %fn5 # branch +%fn6 = block { + br %fn8 +} - # loop continuing - %fn5 = block { - } -> %fn4 # branch +%fn7 = block { + br %fn5 +} - # loop merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn8 = block { + jmp %fn4 +} )"); } @@ -1077,60 +940,57 @@ TEST_F(IR_BuilderImplTest, While_Return) { auto* ast_while = While(true, Block(Return())); WrapInFunction(ast_while); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Start().target, nullptr); - ASSERT_NE(flow->Continuing().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_NE(flow->Start().target->As()->Branch().target, nullptr); - ASSERT_TRUE(flow->Start().target->As()->Branch().target->Is()); - auto* if_flow = flow->Start().target->As()->Branch().target->As(); - ASSERT_NE(if_flow->True().target, nullptr); - ASSERT_NE(if_flow->False().target, nullptr); - ASSERT_NE(if_flow->Merge().target, nullptr); + ASSERT_NE(flow->Start()->Branch(), nullptr); + ASSERT_TRUE(flow->Start()->Branch()->Is()); + auto* if_flow = flow->Start()->Branch()->As(); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); + EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(1u, 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(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, m: %fn5] - # loop start - %fn4 = block { - } -> %fn6 # branch +%fn3 = block { + if true [t: %fn6, f: %fn7, m: %fn8] +} - %fn6 = if true [t: %fn7, f: %fn8, m: %fn9] - # true branch - %fn7 = block { - } -> %fn9 # branch +%fn4 = block { + br %fn3 +} - # false branch - %fn8 = block { - } -> %fn5 # branch +%fn5 = block { + jmp %fn9 # return +} +%fn9 = func_terminator - # if merge - %fn9 = block { - } -> %func_end # return - # loop merge - %fn5 = block { - } -> %func_end # return -} %func_end +%fn6 = block { + br %fn8 +} +%fn7 = block { + br %fn5 +} + +%fn8 = block { + br %fn9 # return +} )"); } @@ -1151,71 +1011,66 @@ TEST_F(IR_BuilderImplTest, DISABLED_For) { auto* ast_for = For(Decl(Var("i", ty.i32())), LessThan("i", 10_a), Increment("i"), Block()); WrapInFunction(ast_for); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Start().target, nullptr); - ASSERT_NE(flow->Continuing().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_NE(flow->Start().target->As()->Branch().target, nullptr); - ASSERT_TRUE(flow->Start().target->As()->Branch().target->Is()); - auto* if_flow = flow->Start().target->As()->Branch().target->As(); - ASSERT_NE(if_flow->True().target, nullptr); - ASSERT_NE(if_flow->False().target, nullptr); - ASSERT_NE(if_flow->Merge().target, nullptr); + ASSERT_NE(flow->Start()->Branch(), nullptr); + ASSERT_TRUE(flow->Start()->Branch()->Is()); + auto* if_flow = flow->Start()->Branch()->As(); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length()); - EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(2u, 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(Disassemble(m.Get()), R"()"); + EXPECT_EQ(Disassemble(m), R"()"); } TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) { auto* ast_for = For(nullptr, nullptr, nullptr, Block(Break())); WrapInFunction(ast_for); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Start().target, nullptr); - ASSERT_NE(flow->Continuing().target, nullptr); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); + EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + loop [s: %fn3, c: %fn4, m: %fn5] +} - %fn3 = loop [s: %fn4, m: %fn5] - # loop start - %fn4 = block { - } -> %fn5 # branch +%fn3 = block { + br %fn5 +} - # loop merge - %fn5 = block { - } -> %func_end # return -} %func_end +%fn4 = block { + br %fn3 +} + +%fn5 = block { + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -1227,14 +1082,14 @@ TEST_F(IR_BuilderImplTest, Switch) { WrapInFunction(ast_switch); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(3u, cases.Length()); @@ -1252,35 +1107,34 @@ TEST_F(IR_BuilderImplTest, Switch) { ASSERT_EQ(1u, cases[2].selectors.Length()); EXPECT_TRUE(cases[2].selectors[0].IsDefault()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, cases[2].Start().target->InboundBranches().Length()); - EXPECT_EQ(3u, flow->Merge().target->InboundBranches().Length()); + 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(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + switch 1i [c: (0i, %fn3), c: (1i, %fn4), c: (default, %fn5), m: %fn6] +} - %fn3 = switch 1i [c: (0i, %fn4), c: (1i, %fn5), c: (default, %fn6), m: %fn7] - # case 0i - %fn4 = block { - } -> %fn7 # branch +%fn3 = block { + br %fn6 +} - # case 1i - %fn5 = block { - } -> %fn7 # branch +%fn4 = block { + br %fn6 +} - # case default - %fn6 = block { - } -> %fn7 # branch +%fn5 = block { + br %fn6 +} - # switch merge - %fn7 = block { - } -> %func_end # return -} %func_end +%fn6 = block { + jmp %fn7 # return +} +%fn7 = func_terminator )"); } @@ -1293,14 +1147,14 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) { WrapInFunction(ast_switch); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(1u, cases.Length()); @@ -1315,25 +1169,24 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) { EXPECT_TRUE(cases[0].selectors[2].IsDefault()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + switch 1i [c: (0i 1i default, %fn3), m: %fn4] +} - %fn3 = switch 1i [c: (0i 1i default, %fn4), m: %fn5] - # case 0i 1i default - %fn4 = block { - } -> %fn5 # branch +%fn3 = block { + br %fn4 +} - # switch merge - %fn5 = block { - } -> %func_end # return -} %func_end +%fn4 = block { + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -1342,39 +1195,38 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) { auto* ast_switch = Switch(1_i, utils::Vector{DefaultCase(Block())}); WrapInFunction(ast_switch); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(1u, cases.Length()); ASSERT_EQ(1u, cases[0].selectors.Length()); EXPECT_TRUE(cases[0].selectors[0].IsDefault()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); + EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + switch 1i [c: (default, %fn3), m: %fn4] +} - %fn3 = switch 1i [c: (default, %fn4), m: %fn5] - # case default - %fn4 = block { - } -> %fn5 # branch +%fn3 = block { + br %fn4 +} - # switch merge - %fn5 = block { - } -> %func_end # return -} %func_end +%fn4 = block { + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -1385,14 +1237,14 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) { DefaultCase(Block())}); WrapInFunction(ast_switch); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Merge().target, nullptr); + auto m = res.Move(); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(2u, cases.Length()); @@ -1404,31 +1256,30 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) { ASSERT_EQ(1u, cases[1].selectors.Length()); EXPECT_TRUE(cases[1].selectors[0].IsDefault()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length()); - EXPECT_EQ(2u, flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); + EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); + EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length()); // This is 1 because the if is dead-code eliminated and the return doesn't happen. EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + switch 1i [c: (0i, %fn3), c: (default, %fn4), m: %fn5] +} - %fn3 = switch 1i [c: (0i, %fn4), c: (default, %fn5), m: %fn6] - # case 0i - %fn4 = block { - } -> %fn6 # branch +%fn3 = block { + br %fn5 +} - # case default - %fn5 = block { - } -> %fn6 # branch +%fn4 = block { + br %fn5 +} - # switch merge - %fn6 = block { - } -> %func_end # return -} %func_end +%fn5 = block { + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -1440,16 +1291,16 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) { auto* ast_if = If(true, Block(Return())); WrapInFunction(ast_switch, ast_if); - auto m = Build(); - ASSERT_TRUE(m) << (!m ? m.Failure() : ""); + auto res = Build(); + ASSERT_TRUE(res) << (!res ? res.Failure() : ""); - ASSERT_EQ(FindSingleFlowNode(m.Get()), nullptr); + auto m = res.Move(); + ASSERT_EQ(FindSingleValue(m), nullptr); - auto* flow = FindSingleFlowNode(m.Get()); - ASSERT_NE(flow->Merge().target, nullptr); + auto* flow = FindSingleValue(m); - ASSERT_EQ(1u, m->functions.Length()); - auto* func = m->functions[0]; + ASSERT_EQ(1u, m.functions.Length()); + auto* func = m.functions[0]; auto cases = flow->Cases(); ASSERT_EQ(2u, cases.Length()); @@ -1461,25 +1312,24 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) { ASSERT_EQ(1u, cases[1].selectors.Length()); EXPECT_TRUE(cases[1].selectors[0].IsDefault()); - EXPECT_EQ(1u, flow->InboundBranches().Length()); - EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length()); - EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length()); - EXPECT_EQ(0u, flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); + EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); + EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); - EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %fn3 # branch + EXPECT_EQ(Disassemble(m), + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + switch 1i [c: (0i, %fn3), c: (default, %fn4)] +} - %fn3 = switch 1i [c: (0i, %fn4), c: (default, %fn5)] - # case 0i - %fn4 = block { - } -> %func_end # return - # case default - %fn5 = block { - } -> %func_end # return -} %func_end +%fn3 = block { + br %fn5 # return +} +%fn4 = block { + br %fn5 # return +} +%fn5 = func_terminator )"); } @@ -1492,16 +1342,18 @@ TEST_F(IR_BuilderImplTest, Emit_Phony) { ASSERT_TRUE(m) << (!m ? m.Failure() : ""); EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func b():i32 { - %fn2 = block { - } -> %func_end 1i # return -} %func_end + R"(%fn1 = func b():i32 -> %fn2 +%fn2 = block { + br %fn3 1i # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:i32 = call b - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:i32 = call b + jmp %fn6 # return +} +%fn6 = func_terminator )"); } diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc index 2afd2a2b03..a47acf66b8 100644 --- a/src/tint/ir/from_program_unary_test.cc +++ b/src/tint/ir/from_program_unary_test.cc @@ -34,17 +34,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Not) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool { - %fn2 = block { - } -> %func_end false # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool -> %fn2 +%fn2 = block { + br %fn3 false # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:bool = call my_func - %tint_symbol:bool = eq %1, false - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:bool = call my_func + %tint_symbol:bool = eq %1, false + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -57,17 +59,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Complement) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 { - %fn2 = block { - } -> %func_end 1u # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 +%fn2 = block { + br %fn3 1u # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:u32 = call my_func - %tint_symbol:u32 = complement %1 - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:u32 = call my_func + %tint_symbol:u32 = complement %1 + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -80,17 +84,19 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Negation) { auto m = Build(); ASSERT_TRUE(m) << (!m ? m.Failure() : ""); - EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():i32 { - %fn2 = block { - } -> %func_end 1i # return -} %func_end + EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():i32 -> %fn2 +%fn2 = block { + br %fn3 1i # return +} +%fn3 = func_terminator -%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn4 = block { - %1:i32 = call my_func - %tint_symbol:i32 = negation %1 - } -> %func_end # return -} %func_end +%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 +%fn5 = block { + %1:i32 = call my_func + %tint_symbol:i32 = negation %1 + jmp %fn6 # return +} +%fn6 = func_terminator )"); } @@ -106,13 +112,16 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v2:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + jmp %fn5 # return +} +%fn5 = func_terminator )"); } @@ -130,14 +139,17 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %v3:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator -%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn3 = block { - store %v3, 42i - } -> %func_end # return -} %func_end +%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 +%fn4 = block { + store %v3, 42i + jmp %fn5 # return +} +%fn5 = func_terminator )"); } diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc index e235f8880e..2e4926580c 100644 --- a/src/tint/ir/from_program_var_test.cc +++ b/src/tint/ir/from_program_var_test.cc @@ -34,8 +34,10 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %a:ptr = var + br %fn2 # root_end } +%fn2 = root_terminator )"); } @@ -49,8 +51,10 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) { EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block { %a:ptr = var, 2u + br %fn2 # root_end } +%fn2 = root_terminator )"); } @@ -63,11 +67,12 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) { ASSERT_TRUE(m) << (!m ? m.Failure() : ""); EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - %a:ptr = var - } -> %func_end # return -} %func_end + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + %a:ptr = var + jmp %fn3 # return +} +%fn3 = func_terminator )"); } @@ -81,11 +86,12 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) { ASSERT_TRUE(m) << (!m ? m.Failure() : ""); EXPECT_EQ(Disassemble(m.Get()), - R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - %a:ptr = var, 2u - } -> %func_end # return -} %func_end + R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + %a:ptr = var, 2u + jmp %fn3 # return +} +%fn3 = func_terminator )"); } diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc index d235b5efd3..a89a51f765 100644 --- a/src/tint/ir/if.cc +++ b/src/tint/ir/if.cc @@ -18,7 +18,16 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::If); namespace tint::ir { -If::If(Value* cond) : Base(), condition_(cond) {} +If::If(Value* cond, Block* t, Block* f, Block* m) + : Base(m), condition_(cond), true_(t), false_(f), merge_(m) { + TINT_ASSERT(IR, true_); + TINT_ASSERT(IR, false_); + TINT_ASSERT(IR, merge_); + + condition_->AddUsage(this); + true_->AddInboundBranch(this); + false_->AddInboundBranch(this); +} If::~If() = default; diff --git a/src/tint/ir/if.h b/src/tint/ir/if.h index d404c94a6d..ad4db1de87 100644 --- a/src/tint/ir/if.h +++ b/src/tint/ir/if.h @@ -15,8 +15,8 @@ #ifndef SRC_TINT_IR_IF_H_ #define SRC_TINT_IR_IF_H_ +#include "src/tint/ir/block.h" #include "src/tint/ir/branch.h" -#include "src/tint/ir/flow_node.h" #include "src/tint/ir/value.h" // Forward declarations @@ -26,37 +26,42 @@ class Block; namespace tint::ir { -/// A flow node representing an if statement. -class If : public utils::Castable { +/// An if instruction +class If : public utils::Castable { public: /// Constructor /// @param cond the if condition - explicit If(Value* cond); + /// @param t the true block + /// @param f the false block + /// @param m the merge block + explicit If(Value* cond, Block* t, Block* f, Block* m); ~If() override; /// @returns the if condition const Value* Condition() const { return condition_; } + /// @returns the if condition + Value* Condition() { return condition_; } /// @returns the true branch block - const Branch& True() const { return true_; } + const Block* True() const { return true_; } /// @returns the true branch block - Branch& True() { return true_; } + Block* True() { return true_; } /// @returns the false branch block - const Branch& False() const { return false_; } + const Block* False() const { return false_; } /// @returns the false branch block - Branch& False() { return false_; } + Block* False() { return false_; } /// @returns the merge branch block - const Branch& Merge() const { return merge_; } + const Block* Merge() const { return merge_; } /// @returns the merge branch block - Branch& Merge() { return merge_; } + Block* Merge() { return merge_; } private: - Branch true_ = {}; - Branch false_ = {}; - Branch merge_ = {}; - Value* condition_; + Value* condition_ = nullptr; + Block* true_ = nullptr; + Block* false_ = nullptr; + Block* merge_ = nullptr; }; } // namespace tint::ir diff --git a/src/tint/ir/jump.cc b/src/tint/ir/jump.cc new file mode 100644 index 0000000000..cda2f06db4 --- /dev/null +++ b/src/tint/ir/jump.cc @@ -0,0 +1,25 @@ +// 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. + +#include "src/tint/ir/jump.h" + +TINT_INSTANTIATE_TYPEINFO(tint::ir::Jump); + +namespace tint::ir { + +Jump::Jump(FlowNode* to, utils::VectorRef args) : Base(to, args) {} + +Jump::~Jump() = default; + +} // namespace tint::ir diff --git a/src/tint/ir/jump.h b/src/tint/ir/jump.h new file mode 100644 index 0000000000..3159755c5a --- /dev/null +++ b/src/tint/ir/jump.h @@ -0,0 +1,37 @@ +// 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_JUMP_H_ +#define SRC_TINT_IR_JUMP_H_ + +#include "src/tint/ir/block.h" +#include "src/tint/ir/branch.h" +#include "src/tint/ir/value.h" +#include "src/tint/utils/castable.h" + +namespace tint::ir { + +/// A jump instruction. A jump is walk continuing. +class Jump : public utils::Castable { + public: + /// Constructor + /// @param to the block to branch too + /// @param args the branch arguments + explicit Jump(FlowNode* to, utils::VectorRef args = {}); + ~Jump() override; +}; + +} // namespace tint::ir + +#endif // SRC_TINT_IR_JUMP_H_ diff --git a/src/tint/ir/loop.cc b/src/tint/ir/loop.cc index 9a1af45a8d..0bbb710e54 100644 --- a/src/tint/ir/loop.cc +++ b/src/tint/ir/loop.cc @@ -18,7 +18,11 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Loop); namespace tint::ir { -Loop::Loop() : Base() {} +Loop::Loop(Block* s, Block* c, Block* m) : Base(s), start_(s), continuing_(c), merge_(m) { + TINT_ASSERT(IR, start_); + TINT_ASSERT(IR, continuing_); + TINT_ASSERT(IR, merge_); +} Loop::~Loop() = default; diff --git a/src/tint/ir/loop.h b/src/tint/ir/loop.h index 44f7881221..954a64aa57 100644 --- a/src/tint/ir/loop.h +++ b/src/tint/ir/loop.h @@ -17,36 +17,38 @@ #include "src/tint/ir/block.h" #include "src/tint/ir/branch.h" -#include "src/tint/ir/flow_node.h" namespace tint::ir { /// Flow node describing a loop. -class Loop : public utils::Castable { +class Loop : public utils::Castable { public: /// Constructor - Loop(); + /// @param s the start block + /// @param c the continuing block + /// @param m the merge block + Loop(Block* s, Block* c, Block* m); ~Loop() override; /// @returns the switch start branch - const Branch& Start() const { return start_; } + const Block* Start() const { return start_; } /// @returns the switch start branch - Branch& Start() { return start_; } + Block* Start() { return start_; } /// @returns the switch continuing branch - const Branch& Continuing() const { return continuing_; } + const Block* Continuing() const { return continuing_; } /// @returns the switch continuing branch - Branch& Continuing() { return continuing_; } + Block* Continuing() { return continuing_; } /// @returns the switch merge branch - const Branch& Merge() const { return merge_; } + const Block* Merge() const { return merge_; } /// @returns the switch merge branch - Branch& Merge() { return merge_; } + Block* Merge() { return merge_; } private: - Branch start_ = {}; - Branch continuing_ = {}; - Branch merge_ = {}; + Block* start_ = nullptr; + Block* continuing_ = nullptr; + Block* merge_ = nullptr; }; } // namespace tint::ir diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc index 3bccc83677..a28666a6b1 100644 --- a/src/tint/ir/switch.cc +++ b/src/tint/ir/switch.cc @@ -18,7 +18,11 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Switch); namespace tint::ir { -Switch::Switch(Value* cond) : Base(), condition_(cond) {} +Switch::Switch(Value* cond, Block* m) : Base(m), condition_(cond), merge_(m) { + TINT_ASSERT(IR, condition_); + TINT_ASSERT(IR, merge_); + condition_->AddUsage(this); +} Switch::~Switch() = default; diff --git a/src/tint/ir/switch.h b/src/tint/ir/switch.h index b4719993f4..588fee9f1d 100644 --- a/src/tint/ir/switch.h +++ b/src/tint/ir/switch.h @@ -18,13 +18,12 @@ #include "src/tint/ir/block.h" #include "src/tint/ir/branch.h" #include "src/tint/ir/constant.h" -#include "src/tint/ir/flow_node.h" #include "src/tint/ir/value.h" namespace tint::ir { /// Flow node representing a switch statement -class Switch : public utils::Castable { +class Switch : public utils::Castable { public: /// A case selector struct CaseSelector { @@ -40,23 +39,24 @@ class Switch : public utils::Castable { /// The case selector for this node utils::Vector selectors; /// The start block for the case block. - Branch start = {}; + Block* start = nullptr; /// @returns the case start target - const Branch& Start() const { return start; } + const Block* Start() const { return start; } /// @returns the case start target - Branch& Start() { return start; } + Block* Start() { return start; } }; /// Constructor /// @param cond the condition - explicit Switch(Value* cond); + /// @param m the merge block + explicit Switch(Value* cond, Block* m); ~Switch() override; /// @returns the switch merge branch - const Branch& Merge() const { return merge_; } + const Block* Merge() const { return merge_; } /// @returns the switch merge branch - Branch& Merge() { return merge_; } + Block* Merge() { return merge_; } /// @returns the switch cases utils::VectorRef Cases() const { return cases_; } @@ -65,11 +65,13 @@ class Switch : public utils::Castable { /// @returns the condition const Value* Condition() const { return condition_; } + /// @returns the condition + Value* Condition() { return condition_; } private: - Branch merge_ = {}; + Value* condition_ = nullptr; + Block* merge_ = nullptr; utils::Vector cases_; - Value* condition_; }; } // namespace tint::ir diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc index 9be189c216..c1cb9f9125 100644 --- a/src/tint/ir/to_program.cc +++ b/src/tint/ir/to_program.cc @@ -23,6 +23,7 @@ #include "src/tint/ir/function_terminator.h" #include "src/tint/ir/if.h" #include "src/tint/ir/instruction.h" +#include "src/tint/ir/jump.h" #include "src/tint/ir/load.h" #include "src/tint/ir/module.h" #include "src/tint/ir/store.h" @@ -108,25 +109,26 @@ class State { std::move(ret_attrs)); } - const ast::BlockStatement* FlowNodeGraph(ir::FlowNode* start_node, - ir::FlowNode* stop_at = nullptr) { + const ast::BlockStatement* FlowNodeGraph(const ir::Block* start_node) { // TODO(crbug.com/tint/1902): Check if the block is dead utils::Vector stmts; - ir::Branch root_branch{start_node, {}}; - const ir::Branch* branch = &root_branch; + const ir::FlowNode* block = start_node; // TODO(crbug.com/tint/1902): Handle block arguments. - while (branch->target != stop_at) { - enum Status { kContinue, kStop, kError }; - Status status = tint::Switch( - branch->target, + while (block) { + TINT_ASSERT(IR, block->HasBranchTarget()); - [&](const ir::Block* block) { - for (const auto* inst : block->Instructions()) { + 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; @@ -135,43 +137,27 @@ class State { stmts.Push(s); } } - branch = &block->Branch(); - return kContinue; - }, - - [&](const ir::If* if_) { - auto* stmt = If(if_); - if (TINT_UNLIKELY(!stmt)) { - return kError; - } - stmts.Push(stmt); - branch = &if_->Merge(); - return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue; - }, - - [&](const ir::Switch* switch_) { - auto* stmt = Switch(switch_); - if (TINT_UNLIKELY(!stmt)) { - return kError; - } - stmts.Push(stmt); - branch = &switch_->Merge(); - return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue; - }, - - [&](const ir::FunctionTerminator*) { - auto res = FunctionTerminator(branch); - if (TINT_UNLIKELY(!res)) { - return kError; - } - if (auto* stmt = res.Get()) { - stmts.Push(stmt); + if (blk->Branch()->Is() && blk->Branch()->To()->Is()) { + block = blk->Branch()->To()->As(); + return kContinue; + } else if (auto* if_ = blk->Branch()->As()) { + if (if_->Merge()->HasBranchTarget()) { + block = if_->Merge(); + return kContinue; + } + } else if (auto* switch_ = blk->Branch()->As()) { + if (switch_->Merge()->HasBranchTarget()) { + block = switch_->Merge(); + return kContinue; + } } return kStop; }, + [&](const ir::FunctionTerminator*) { return kStop; }, + [&](Default) { - UNHANDLED_CASE(branch->target); + UNHANDLED_CASE(block); return kError; }); @@ -188,26 +174,24 @@ class State { const ast::IfStatement* If(const ir::If* i) { SCOPED_NESTING(); - auto* cond = Expr(i->Condition()); - auto* t = FlowNodeGraph(i->True().target, i->Merge().target); + auto* t = FlowNodeGraph(i->True()); if (TINT_UNLIKELY(!t)) { return nullptr; } - if (!IsEmpty(i->False().target, i->Merge().target)) { - // If the else target is an if flow node with the same Merge().target as this if, then - // emit an 'else if' instead of a block statement for the else. - if (auto* else_if = As(NextNonEmptyNode(i->False().target)); - else_if && - NextNonEmptyNode(i->Merge().target) == NextNonEmptyNode(else_if->Merge().target)) { - auto* f = If(else_if); + if (!IsEmpty(i->False(), i->Merge())) { + // 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(); + inst && inst->Merge()->IsTrampoline(i->Merge())) { + auto* f = If(inst); if (!f) { return nullptr; } return b.If(cond, t, b.Else(f)); } else { - auto* f = FlowNodeGraph(i->False().target, i->Merge().target); + auto* f = FlowNodeGraph(i->False()); if (!f) { return nullptr; } @@ -226,11 +210,11 @@ class State { return nullptr; } - auto cases = utils::Transform<1>( + auto cases = utils::Transform<2>( s->Cases(), // - [&](const ir::Switch::Case& c) -> const tint::ast::CaseStatement* { + [&](const ir::Switch::Case c) -> const tint::ast::CaseStatement* { SCOPED_NESTING(); - auto* body = FlowNodeGraph(c.start.target, s->Merge().target); + auto* body = FlowNodeGraph(c.start); if (!body) { return nullptr; } @@ -261,26 +245,27 @@ class State { } utils::Result FunctionTerminator(const ir::Branch* branch) { - if (branch->args.IsEmpty()) { + if (branch->Args().IsEmpty()) { // Branch to function terminator has no arguments. - // If this block is nested withing some control flow, then we must emit a - // 'return' statement, otherwise we've just naturally reached the end of the - // function where the 'return' is redundant. + // If this block is nested withing some control flow, then we must + // emit a 'return' statement, otherwise we've just naturally reached + // the end of the function where the 'return' is redundant. if (nesting_depth_ > 1) { return b.Return(); } return nullptr; } - // Branch to function terminator has arguments - this is the return value. - if (branch->args.Length() != 1) { - TINT_ICE(IR, b.Diagnostics()) - << "expected 1 value for function terminator (return value), got " - << branch->args.Length(); + // Branch to function terminator has arguments - this is the return + // value. + if (branch->Args().Length() != 1) { + TINT_ICE(IR, b.Diagnostics()) << "expected 1 value for function " + "terminator (return value), got " + << branch->Args().Length(); return utils::Failure; } - auto* val = Expr(branch->args.Front()); + auto* val = Expr(branch->Args().Front()); if (TINT_UNLIKELY(!val)) { return utils::Failure; } @@ -289,36 +274,16 @@ class State { } /// @return true if there are no instructions between @p node and and @p stop_at - bool IsEmpty(const ir::FlowNode* node, const ir::FlowNode* stop_at) { - while (node != stop_at) { - if (auto* block = node->As()) { - if (!block->Instructions().IsEmpty()) { - return false; - } - node = block->Branch().target; - } else { - return false; - } + bool IsEmpty(const ir::Block* node, const ir::FlowNode* stop_at) { + if (node->Instructions().IsEmpty()) { + return true; } - return true; - } - - /// @return the next flow node that isn't an empty block - const ir::FlowNode* NextNonEmptyNode(const ir::FlowNode* node) { - while (node) { - if (auto* block = node->As()) { - for (const auto* inst : block->Instructions()) { - // Load instructions will be inlined, so ignore them. - if (!inst->Is()) { - return node; - } - } - node = block->Branch().target; - } else { - return node; - } + if (auto* br = node->Instructions().Front()->As()) { + return br->To() == stop_at; } - return nullptr; + // TODO(dsinclair): This should possibly walk over Jump instructions that + // just jump to empty blocks if we want to be comprehensive. + return false; } utils::Result Stmt(const ir::Instruction* inst) { @@ -328,6 +293,14 @@ class State { [&](const ir::Var* i) { return Var(i); }, // [&](const ir::Load*) { return nullptr; }, [&](const ir::Store* i) { return Store(i); }, // + [&](const ir::If* if_) { return If(if_); }, + [&](const ir::Switch* switch_) { return Switch(switch_); }, + [&](const ir::Branch* branch) { + if (branch->To()->Is()) { + return utils::Result{FunctionTerminator(branch)}; + } + return utils::Result{nullptr}; + }, [&](Default) { UNHANDLED_CASE(inst); return utils::Failure; diff --git a/src/tint/ir/to_program_roundtrip_test.cc b/src/tint/ir/to_program_roundtrip_test.cc index 8f267924d2..c8af503bcc 100644 --- a/src/tint/ir/to_program_roundtrip_test.cc +++ b/src/tint/ir/to_program_roundtrip_test.cc @@ -229,10 +229,9 @@ fn c() { fn f() { var cond_a : bool = true; - var cond_b : bool = true; if (cond_a) { a(); - } else if (cond_b) { + } else if (false) { b(); } c(); diff --git a/src/tint/ir/transform/add_empty_entry_point.cc b/src/tint/ir/transform/add_empty_entry_point.cc index 6788d7c375..f40c60e524 100644 --- a/src/tint/ir/transform/add_empty_entry_point.cc +++ b/src/tint/ir/transform/add_empty_entry_point.cc @@ -38,7 +38,7 @@ void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const { auto* ep = builder.CreateFunction(ir->symbols.New("unused_entry_point"), ir->types.Get(), Function::PipelineStage::kCompute, std::array{1u, 1u, 1u}); - ep->StartTarget()->BranchTo(ep->EndTarget()); + ep->StartTarget()->SetInstructions(utils::Vector{builder.Branch(ep->EndTarget())}); ir->functions.Push(ep); } diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc index a363d83640..3b951645c2 100644 --- a/src/tint/ir/transform/add_empty_entry_point_test.cc +++ b/src/tint/ir/transform/add_empty_entry_point_test.cc @@ -25,10 +25,11 @@ using IR_AddEmptyEntryPointTest = TransformTest; TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) { auto* expect = R"( -%fn1 = func unused_entry_point():void [@compute @workgroup_size(1, 1, 1)] { - %fn2 = block { - } -> %func_end # return -} %func_end +%fn1 = func unused_entry_point():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 +%fn2 = block { + br %fn3 # return +} +%fn3 = func_terminator )"; @@ -40,14 +41,15 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) { TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) { auto* ep = b.CreateFunction(mod.symbols.New("main"), mod.types.Get(), Function::PipelineStage::kFragment); - ep->StartTarget()->BranchTo(ep->EndTarget()); + ep->StartTarget()->SetInstructions(utils::Vector{b.Branch(ep->EndTarget())}); mod.functions.Push(ep); auto* expect = R"( -%fn1 = func main():void [@fragment] { - %fn2 = block { - } -> %func_end # return -} %func_end +%fn1 = func main():void [@fragment] -> %fn2 +%fn2 = block { + br %fn3 # return +} +%fn3 = func_terminator )"; diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h index 3b64b797a5..698413d309 100644 --- a/src/tint/ir/unary.h +++ b/src/tint/ir/unary.h @@ -41,6 +41,8 @@ class Unary : public utils::Castable { /// @returns the value for the instruction const Value* Val() const { return val_; } + /// @returns the value for the instruction + Value* Val() { return val_; } /// @returns the kind of unary instruction enum Kind Kind() const { return kind_; } diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc index 6cefed38a1..bb0b4f2921 100644 --- a/src/tint/ir/unary_test.cc +++ b/src/tint/ir/unary_test.cc @@ -26,7 +26,7 @@ using IR_InstructionTest = TestHelper; TEST_F(IR_InstructionTest, CreateComplement) { Module mod; Builder b{mod}; - const auto* inst = b.Complement(b.ir.types.Get(), b.Constant(4_i)); + auto* inst = b.Complement(b.ir.types.Get(), b.Constant(4_i)); ASSERT_TRUE(inst->Is()); EXPECT_EQ(inst->Kind(), Unary::Kind::kComplement); @@ -40,7 +40,7 @@ TEST_F(IR_InstructionTest, CreateComplement) { TEST_F(IR_InstructionTest, CreateNegation) { Module mod; Builder b{mod}; - const auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); + auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); ASSERT_TRUE(inst->Is()); EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation); @@ -54,7 +54,7 @@ TEST_F(IR_InstructionTest, CreateNegation) { TEST_F(IR_InstructionTest, Unary_Usage) { Module mod; Builder b{mod}; - const auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); + auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation); diff --git a/src/tint/transform/manager_test.cc b/src/tint/transform/manager_test.cc index a81f7bf07c..d1f63337aa 100644 --- a/src/tint/transform/manager_test.cc +++ b/src/tint/transform/manager_test.cc @@ -52,7 +52,7 @@ class IR_AddFunction final : public ir::transform::Transform { ir::Builder builder(*mod); auto* func = builder.CreateFunction(mod->symbols.New("ir_func"), mod->types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); + func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())}); mod->functions.Push(func); } }; @@ -70,7 +70,7 @@ ir::Module MakeIR() { ir::Builder builder(mod); auto* func = builder.CreateFunction(builder.ir.symbols.New("main"), builder.ir.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); + func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())}); builder.ir.functions.Push(func); return mod; } diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc index bd9a735bc5..44bead9118 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc @@ -292,8 +292,15 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) { current_function_.push_inst(spv::Op::OpLabel, {Label(block)}); } + // If there are no instructions in the block, it's a dead end, so we shouldn't be able to get + // here to begin with. + if (block->Instructions().IsEmpty()) { + current_function_.push_inst(spv::Op::OpUnreachable, {}); + return; + } + // Emit the instructions. - for (const auto* inst : block->Instructions()) { + for (auto* inst : block->Instructions()) { auto result = Switch( inst, // [&](const ir::Binary* b) { return EmitBinary(b); }, @@ -303,6 +310,14 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) { return 0u; }, [&](const ir::Var* v) { return EmitVar(v); }, + [&](const ir::If* i) { + EmitIf(i); + return 0u; + }, + [&](const ir::Branch* b) { + EmitBranch(b); + return 0u; + }, [&](Default) { TINT_ICE(Writer, diagnostics_) << "unimplemented instruction: " << inst->TypeInfo().name; @@ -310,46 +325,42 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) { }); instructions_.Add(inst, result); } +} - // Handle the branch at the end of the block. +void GeneratorImplIr::EmitBranch(const ir::Branch* b) { Switch( - block->Branch().target, - [&](const ir::Block* b) { current_function_.push_inst(spv::Op::OpBranch, {Label(b)}); }, - [&](const ir::If* i) { EmitIf(i); }, + b->To(), + [&](const ir::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); }, [&](const ir::FunctionTerminator*) { // TODO(jrprice): Handle the return value, which will be a branch argument. - if (!block->Branch().args.IsEmpty()) { + if (!b->Args().IsEmpty()) { TINT_ICE(Writer, diagnostics_) << "unimplemented return value"; } current_function_.push_inst(spv::Op::OpReturn, {}); }, [&](Default) { - if (!block->Branch().target) { - // A block may not have an outward branch (e.g. an unreachable merge block). - current_function_.push_inst(spv::Op::OpUnreachable, {}); - } else { - TINT_ICE(Writer, diagnostics_) - << "unimplemented branch target: " << block->Branch().target->TypeInfo().name; - } + // A block may not have an outward branch (e.g. an unreachable merge + // block). + current_function_.push_inst(spv::Op::OpUnreachable, {}); }); } void GeneratorImplIr::EmitIf(const ir::If* i) { - auto* merge_block = i->Merge().target->As(); - auto* true_block = i->True().target->As(); - auto* false_block = i->False().target->As(); + auto* merge_block = i->Merge(); + auto* true_block = i->True(); + auto* false_block = i->False(); // Generate labels for the blocks. We emit the true or false block if it: - // 1. contains instructions, or - // 2. branches somewhere other then the Merge().target. + // 1. contains instructions other then the branch, or + // 2. branches somewhere other then the Merge(). // Otherwise we skip them and branch straight to the merge block. uint32_t merge_label = Label(merge_block); uint32_t true_label = merge_label; uint32_t false_label = merge_label; - if (!true_block->Instructions().IsEmpty() || true_block->Branch().target != merge_block) { + if (true_block->Instructions().Length() > 1 || true_block->Branch()->To() != merge_block) { true_label = Label(true_block); } - if (!false_block->Instructions().IsEmpty() || false_block->Branch().target != merge_block) { + if (false_block->Instructions().Length() > 1 || false_block->Branch()->To() != merge_block) { false_label = Label(false_block); } diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.h b/src/tint/writer/spirv/ir/generator_impl_ir.h index b72201f5c6..ccd09ad75b 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir.h +++ b/src/tint/writer/spirv/ir/generator_impl_ir.h @@ -30,6 +30,7 @@ namespace tint::ir { class Binary; class Block; +class Branch; class If; class Function; class Load; @@ -121,6 +122,10 @@ class GeneratorImplIr { /// @returns the result ID of the instruction uint32_t EmitVar(const ir::Var* var); + /// Emit a branch instruction. + /// @param b the branch instruction to emit + void EmitBranch(const ir::Branch* b); + private: /// Get the result ID of the constant `constant`, emitting its instruction if necessary. /// @param constant the constant to get the ID for diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc index e9231f32b4..bc8757aa71 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc @@ -21,10 +21,9 @@ namespace { TEST_F(SpvGeneratorImplTest, Binary_Add_I32) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - func->StartTarget()->SetInstructions( - utils::Vector{b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(2_i))}); + utils::Vector{b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(2_i)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -43,10 +42,9 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Add_U32) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - func->StartTarget()->SetInstructions( - utils::Vector{b.Add(mod.types.Get(), b.Constant(1_u), b.Constant(2_u))}); + utils::Vector{b.Add(mod.types.Get(), b.Constant(1_u), b.Constant(2_u)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -65,10 +63,9 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Add_F32) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - func->StartTarget()->SetInstructions( - utils::Vector{b.Add(mod.types.Get(), b.Constant(1_f), b.Constant(2_f))}); + utils::Vector{b.Add(mod.types.Get(), b.Constant(1_f), b.Constant(2_f)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -87,10 +84,9 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Sub_I32) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - func->StartTarget()->SetInstructions( - utils::Vector{b.Subtract(mod.types.Get(), b.Constant(1_i), b.Constant(2_i))}); + utils::Vector{b.Subtract(mod.types.Get(), b.Constant(1_i), b.Constant(2_i)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -109,10 +105,9 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Sub_U32) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - func->StartTarget()->SetInstructions( - utils::Vector{b.Subtract(mod.types.Get(), b.Constant(1_u), b.Constant(2_u))}); + utils::Vector{b.Subtract(mod.types.Get(), b.Constant(1_u), b.Constant(2_u)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -131,10 +126,9 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Sub_F32) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - func->StartTarget()->SetInstructions( - utils::Vector{b.Subtract(mod.types.Get(), b.Constant(1_f), b.Constant(2_f))}); + utils::Vector{b.Subtract(mod.types.Get(), b.Constant(1_f), b.Constant(2_f)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -153,8 +147,6 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec2i) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - auto* lhs = mod.constants_arena.Create( mod.types.Get(mod.types.Get(), 2u), utils::Vector{b.Constant(42_i)->Value(), b.Constant(-1_i)->Value()}, false, false); @@ -163,7 +155,8 @@ TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec2i) { utils::Vector{b.Constant(0_i)->Value(), b.Constant(-43_i)->Value()}, false, false); func->StartTarget()->SetInstructions( utils::Vector{b.Subtract(mod.types.Get(mod.types.Get(), 2u), - b.Constant(lhs), b.Constant(rhs))}); + b.Constant(lhs), b.Constant(rhs)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -187,8 +180,6 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec4f) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - auto* lhs = mod.constants_arena.Create( mod.types.Get(mod.types.Get(), 4u), utils::Vector{b.Constant(42_f)->Value(), b.Constant(-1_f)->Value(), @@ -201,7 +192,8 @@ TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec4f) { false, false); func->StartTarget()->SetInstructions( utils::Vector{b.Subtract(mod.types.Get(mod.types.Get(), 4u), - b.Constant(lhs), b.Constant(rhs))}); + b.Constant(lhs), b.Constant(rhs)), + b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -227,10 +219,9 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Binary_Chain) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); - auto* a = b.Subtract(mod.types.Get(), b.Constant(1_i), b.Constant(2_i)); - func->StartTarget()->SetInstructions(utils::Vector{a, b.Add(mod.types.Get(), a, a)}); + func->StartTarget()->SetInstructions( + utils::Vector{a, b.Add(mod.types.Get(), a, a), b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc index d2af246a79..34722012d1 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc @@ -19,7 +19,7 @@ namespace { TEST_F(SpvGeneratorImplTest, Function_Empty) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); + func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -35,7 +35,7 @@ OpFunctionEnd // Test that we do not emit the same function type more than once. TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); + func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); generator_.EmitFunction(func); generator_.EmitFunction(func); @@ -48,7 +48,7 @@ TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) { TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Compute) { auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get(), ir::Function::PipelineStage::kCompute, {{32, 4, 1}}); - func->StartTarget()->BranchTo(func->EndTarget()); + func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main" @@ -66,7 +66,7 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Fragment) { auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get(), ir::Function::PipelineStage::kFragment); - func->StartTarget()->BranchTo(func->EndTarget()); + func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Fragment %1 "main" @@ -84,7 +84,7 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Vertex) { auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get(), ir::Function::PipelineStage::kVertex); - func->StartTarget()->BranchTo(func->EndTarget()); + func->StartTarget()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Vertex %1 "main" @@ -101,15 +101,15 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Multiple) { auto* f1 = b.CreateFunction(mod.symbols.Register("main1"), mod.types.Get(), ir::Function::PipelineStage::kCompute, {{32, 4, 1}}); - f1->StartTarget()->BranchTo(f1->EndTarget()); + f1->StartTarget()->SetInstructions(utils::Vector{b.Branch(f1->EndTarget())}); auto* f2 = b.CreateFunction(mod.symbols.Register("main2"), mod.types.Get(), ir::Function::PipelineStage::kCompute, {{8, 2, 16}}); - f2->StartTarget()->BranchTo(f2->EndTarget()); + f2->StartTarget()->SetInstructions(utils::Vector{b.Branch(f2->EndTarget())}); auto* f3 = b.CreateFunction(mod.symbols.Register("main3"), mod.types.Get(), ir::Function::PipelineStage::kFragment); - f3->StartTarget()->BranchTo(f3->EndTarget()); + f3->StartTarget()->SetInstructions(utils::Vector{b.Branch(f3->EndTarget())}); generator_.EmitFunction(f1); generator_.EmitFunction(f2); 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 fa8cc6bc11..2ed8661c61 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,11 +23,11 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); auto* i = b.CreateIf(b.Constant(true)); - i->True().target->As()->BranchTo(i->Merge().target); - i->False().target->As()->BranchTo(i->Merge().target); - i->Merge().target->As()->BranchTo(func->EndTarget()); + i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); + i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); + i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); - func->StartTarget()->BranchTo(i); + func->StartTarget()->SetInstructions(utils::Vector{i}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -49,15 +49,14 @@ TEST_F(SpvGeneratorImplTest, If_FalseEmpty) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); auto* i = b.CreateIf(b.Constant(true)); - i->False().target->As()->BranchTo(i->Merge().target); - i->Merge().target->As()->BranchTo(func->EndTarget()); + i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); + i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); - auto* true_block = i->True().target->As(); - true_block->SetInstructions( - utils::Vector{b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(1_i))}); - true_block->BranchTo(i->Merge().target); + auto* true_block = i->True(); + true_block->SetInstructions(utils::Vector{ + b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())}); - func->StartTarget()->BranchTo(i); + func->StartTarget()->SetInstructions(utils::Vector{i}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -84,15 +83,14 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); auto* i = b.CreateIf(b.Constant(true)); - i->True().target->As()->BranchTo(i->Merge().target); - i->Merge().target->As()->BranchTo(func->EndTarget()); + i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); + i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); - auto* false_block = i->False().target->As(); - false_block->SetInstructions( - utils::Vector{b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(1_i))}); - false_block->BranchTo(i->Merge().target); + auto* false_block = i->False(); + false_block->SetInstructions(utils::Vector{ + b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())}); - func->StartTarget()->BranchTo(i); + func->StartTarget()->SetInstructions(utils::Vector{i}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -119,11 +117,10 @@ TEST_F(SpvGeneratorImplTest, If_BothBranchesReturn) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); auto* i = b.CreateIf(b.Constant(true)); - i->True().target->As()->BranchTo(func->EndTarget()); - i->False().target->As()->BranchTo(func->EndTarget()); - i->Merge().target->As()->BranchTo(nullptr); + i->True()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); - func->StartTarget()->BranchTo(i); + func->StartTarget()->SetInstructions(utils::Vector{i}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" 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 20e5167980..a381412e87 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 @@ -22,11 +22,10 @@ namespace { TEST_F(SpvGeneratorImplTest, FunctionVar_NoInit) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); - func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty)}); + func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty), b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -44,14 +43,13 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, FunctionVar_WithInit) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); auto* v = b.Declare(ty); v->SetInitializer(b.Constant(42_i)); - func->StartTarget()->SetInstructions(utils::Vector{v}); + func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -71,12 +69,11 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, FunctionVar_Name) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); auto* v = b.Declare(ty); - func->StartTarget()->SetInstructions(utils::Vector{v}); + func->StartTarget()->SetInstructions(utils::Vector{v, b.Branch(func->EndTarget())}); mod.SetName(v, "myvar"); generator_.EmitFunction(func); @@ -96,7 +93,6 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); @@ -104,14 +100,11 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) { v->SetInitializer(b.Constant(42_i)); auto* i = b.CreateIf(b.Constant(true)); - i->False().target->As()->BranchTo(func->EndTarget()); - i->Merge().target->As()->BranchTo(func->EndTarget()); + i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())}); + i->False()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); + i->Merge()->SetInstructions(utils::Vector{b.Branch(func->EndTarget())}); - auto* true_block = i->True().target->As(); - true_block->SetInstructions(utils::Vector{v}); - true_block->BranchTo(i->Merge().target); - - func->StartTarget()->BranchTo(i); + func->StartTarget()->SetInstructions(utils::Vector{i}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -140,13 +133,12 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, FunctionVar_Load) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); auto* store_ty = mod.types.Get(); auto* ty = mod.types.Get(store_ty, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); auto* v = b.Declare(ty); - func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v)}); + func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v), b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -165,12 +157,12 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, FunctionVar_Store) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); auto* v = b.Declare(ty); - func->StartTarget()->SetInstructions(utils::Vector{v, b.Store(v, b.Constant(42_i))}); + func->StartTarget()->SetInstructions( + utils::Vector{v, b.Store(v, b.Constant(42_i)), b.Branch(func->EndTarget())}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"