[tint][ir] Add MultiInBlock that derives from Block
IncomingBranches and BlockParams are only used to track branches from
two different blocks. Create a new MultiInBlock class that derives
from Block, and moves the IncomingBranches and BlockParams to this
new derived class. Rename IncomingBranches to InboundSiblingBranches
to signify that the ControlInstruction is not included in this list.
The only places that MultiInBlock are used are on the merge target of
If, Loop and Switch, and the Body and Continuing blocks of a Loop.
The rest of the blocks have a single incoming branch, which is
implicitly known by the flow-control instruction.
Bug: tint:1718
Change-Id: I49b54950076412815457b1cc2566368e55b19bd2
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/136280
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index bd06461..77c4ade 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1269,6 +1269,8 @@
"ir/loop.h",
"ir/module.cc",
"ir/module.h",
+ "ir/multi_in_block.cc",
+ "ir/multi_in_block.h",
"ir/next_iteration.cc",
"ir/next_iteration.h",
"ir/operand_instruction.cc",
@@ -2330,6 +2332,7 @@
"ir/load_test.cc",
"ir/loop_test.cc",
"ir/module_test.cc",
+ "ir/multi_in_block_test.cc",
"ir/next_iteration_test.cc",
"ir/program_test_helper.h",
"ir/return_test.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 764be39..4e1a6b6 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -774,6 +774,8 @@
ir/loop.h
ir/module.cc
ir/module.h
+ ir/multi_in_block.cc
+ ir/multi_in_block.h
ir/next_iteration.cc
ir/next_iteration.h
ir/operand_instruction.cc
@@ -1531,6 +1533,7 @@
ir/load_test.cc
ir/loop_test.cc
ir/module_test.cc
+ ir/multi_in_block_test.cc
ir/next_iteration_test.cc
ir/program_test_helper.h
ir/return_test.cc
diff --git a/src/tint/ir/block.cc b/src/tint/ir/block.cc
index b545f34..8942ea9 100644
--- a/src/tint/ir/block.cc
+++ b/src/tint/ir/block.cc
@@ -22,22 +22,6 @@
Block::~Block() = default;
-void Block::SetParams(utils::VectorRef<const BlockParam*> params) {
- params_ = std::move(params);
-
- for (auto* param : params_) {
- TINT_ASSERT(IR, param != nullptr);
- }
-}
-
-void Block::AddInboundBranch(ir::Branch* node) {
- TINT_ASSERT(IR, node != nullptr);
-
- if (node) {
- inbound_branches_.Push(node);
- }
-}
-
void Block::Prepend(Instruction* inst) {
TINT_ASSERT_OR_RETURN(IR, inst);
TINT_ASSERT_OR_RETURN(IR, inst->Block() == nullptr);
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 1da6fb9..eaf8aca 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -17,7 +17,6 @@
#include <utility>
-#include "src/tint/ir/block_param.h"
#include "src/tint/ir/branch.h"
#include "src/tint/ir/instruction.h"
#include "src/tint/utils/vector.h"
@@ -131,21 +130,6 @@
/// @returns the number of instructions in the block
size_t Length() const { return instructions_.count; }
- /// Sets the params to the block
- /// @param params the params for the block
- void SetParams(utils::VectorRef<const BlockParam*> params);
- /// @return the parameters passed into the block
- utils::VectorRef<const BlockParam*> Params() const { return params_; }
- /// @returns the params to the block
- utils::Vector<const BlockParam*, 0>& Params() { return params_; }
-
- /// @returns the inbound branch list for the block
- utils::VectorRef<ir::Branch*> InboundBranches() const { return inbound_branches_; }
-
- /// Adds the given node to the inbound branches
- /// @param node the node to add
- void AddInboundBranch(ir::Branch* node);
-
/// @return the parent instruction that owns this block
ControlInstruction* Parent() const { return parent_; }
@@ -159,15 +143,6 @@
size_t count = 0;
} instructions_;
- utils::Vector<const BlockParam*, 0> params_;
-
- /// The list of branches into this node. This list maybe empty for several
- /// reasons:
- /// - 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<ir::Branch*, 2> inbound_branches_;
-
ControlInstruction* parent_ = nullptr;
};
diff --git a/src/tint/ir/block_test.cc b/src/tint/ir/block_test.cc
index 8a5308f..33721ce 100644
--- a/src/tint/ir/block_test.cc
+++ b/src/tint/ir/block_test.cc
@@ -14,7 +14,6 @@
#include "src/tint/ir/block.h"
#include "gtest/gtest-spi.h"
-#include "src/tint/ir/block_param.h"
#include "src/tint/ir/ir_test_helper.h"
namespace tint::ir {
@@ -69,18 +68,6 @@
EXPECT_TRUE(blk->HasBranchTarget());
}
-TEST_F(IR_BlockTest, HasBranchTarget_If) {
- auto* blk = b.CreateBlock();
- blk->Append(b.CreateIf(b.Constant(true)));
- EXPECT_TRUE(blk->HasBranchTarget());
-}
-
-TEST_F(IR_BlockTest, HasBranchTarget_Loop) {
- auto* blk = b.CreateBlock();
- blk->Append(b.CreateLoop());
- EXPECT_TRUE(blk->HasBranchTarget());
-}
-
TEST_F(IR_BlockTest, HasBranchTarget_NextIteration) {
auto* blk = b.CreateBlock();
auto* loop = b.CreateLoop();
@@ -96,12 +83,6 @@
EXPECT_TRUE(blk->HasBranchTarget());
}
-TEST_F(IR_BlockTest, HasBranchTarget_Switch) {
- auto* blk = b.CreateBlock();
- blk->Append(b.CreateSwitch(b.Constant(true)));
- EXPECT_TRUE(blk->HasBranchTarget());
-}
-
TEST_F(IR_BlockTest, SetInstructions) {
auto* inst1 = b.CreateLoop();
auto* inst2 = b.CreateLoop();
@@ -739,29 +720,5 @@
"internal compiler error");
}
-TEST_F(IR_BlockTest, Fail_NullBlockParam) {
- EXPECT_FATAL_FAILURE(
- {
- Module mod;
- Builder b{mod};
-
- auto* blk = b.CreateBlock();
- blk->SetParams(utils::Vector<const BlockParam*, 1>{nullptr});
- },
- "");
-}
-
-TEST_F(IR_BlockTest, Fail_NullInboundBranch) {
- EXPECT_FATAL_FAILURE(
- {
- Module mod;
- Builder b{mod};
-
- auto* blk = b.CreateBlock();
- blk->AddInboundBranch(nullptr);
- },
- "");
-}
-
} // namespace
} // namespace tint::ir
diff --git a/src/tint/ir/break_if.cc b/src/tint/ir/break_if.cc
index 644502b..4a18fb8 100644
--- a/src/tint/ir/break_if.cc
+++ b/src/tint/ir/break_if.cc
@@ -18,6 +18,7 @@
#include "src/tint/ir/block.h"
#include "src/tint/ir/loop.h"
+#include "src/tint/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::BreakIf);
@@ -32,8 +33,8 @@
AddOperand(condition);
if (loop_) {
- loop_->Body()->AddInboundBranch(this);
- loop_->Merge()->AddInboundBranch(this);
+ loop_->Body()->AddInboundSiblingBranch(this);
+ loop_->Merge()->AddInboundSiblingBranch(this);
}
AddOperands(std::move(args));
}
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index fdf11b9..ab46603 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -37,6 +37,10 @@
return ir.blocks.Create<Block>();
}
+MultiInBlock* Builder::CreateMultiInBlock() {
+ return ir.blocks.Create<MultiInBlock>();
+}
+
Function* Builder::CreateFunction(std::string_view name,
const type::Type* return_type,
Function::PipelineStage stage,
@@ -48,24 +52,23 @@
}
If* Builder::CreateIf(Value* condition) {
- return ir.values.Create<If>(condition, CreateBlock(), CreateBlock(), CreateBlock());
+ return ir.values.Create<If>(condition, CreateBlock(), CreateBlock(), CreateMultiInBlock());
}
Loop* Builder::CreateLoop() {
- return ir.values.Create<Loop>(CreateBlock(), CreateBlock(), CreateBlock(), CreateBlock());
+ return ir.values.Create<Loop>(CreateBlock(), CreateMultiInBlock(), CreateMultiInBlock(),
+ CreateMultiInBlock());
}
Switch* Builder::CreateSwitch(Value* condition) {
- return ir.values.Create<Switch>(condition, CreateBlock());
+ return ir.values.Create<Switch>(condition, CreateMultiInBlock());
}
Block* Builder::CreateCase(Switch* s, utils::VectorRef<Switch::CaseSelector> selectors) {
- s->Cases().Push(Switch::Case{std::move(selectors), CreateBlock()});
-
- Block* b = s->Cases().Back().Start();
- b->AddInboundBranch(s);
- b->SetParent(s);
- return b;
+ auto* block = CreateBlock();
+ s->Cases().Push(Switch::Case{std::move(selectors), block});
+ block->SetParent(s);
+ return block;
}
Binary* Builder::CreateBinary(enum Binary::Kind kind,
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 4849743..c4788e5 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -38,6 +38,7 @@
#include "src/tint/ir/load.h"
#include "src/tint/ir/loop.h"
#include "src/tint/ir/module.h"
+#include "src/tint/ir/multi_in_block.h"
#include "src/tint/ir/next_iteration.h"
#include "src/tint/ir/return.h"
#include "src/tint/ir/store.h"
@@ -67,9 +68,12 @@
/// Destructor
~Builder();
- /// @returns a new block flow node
+ /// @returns a new block
Block* CreateBlock();
+ /// @returns a new multi-in block
+ MultiInBlock* CreateMultiInBlock();
+
/// Creates a function flow node
/// @param name the function name
/// @param return_type the function return type
diff --git a/src/tint/ir/continue.cc b/src/tint/ir/continue.cc
index a7282ac..e1e8d84 100644
--- a/src/tint/ir/continue.cc
+++ b/src/tint/ir/continue.cc
@@ -18,6 +18,7 @@
#include "src/tint/ir/block.h"
#include "src/tint/ir/loop.h"
+#include "src/tint/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::Continue);
@@ -28,7 +29,7 @@
TINT_ASSERT(IR, loop_);
if (loop_) {
- loop_->Continuing()->AddInboundBranch(this);
+ loop_->Continuing()->AddInboundSiblingBranch(this);
}
AddOperands(std::move(args));
}
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 9f7aa8d..98d55b8 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -22,6 +22,7 @@
#include "src/tint/ir/binary.h"
#include "src/tint/ir/bitcast.h"
#include "src/tint/ir/block.h"
+#include "src/tint/ir/block_param.h"
#include "src/tint/ir/break_if.h"
#include "src/tint/ir/builtin.h"
#include "src/tint/ir/construct.h"
@@ -34,6 +35,7 @@
#include "src/tint/ir/if.h"
#include "src/tint/ir/load.h"
#include "src/tint/ir/loop.h"
+#include "src/tint/ir/multi_in_block.h"
#include "src/tint/ir/next_iteration.h"
#include "src/tint/ir/return.h"
#include "src/tint/ir/store.h"
@@ -137,10 +139,12 @@
void Disassembler::WalkInternal(const Block* blk) {
SourceMarker sm(this);
Indent() << "%b" << IdOf(blk) << " = block";
- if (!blk->Params().IsEmpty()) {
- out_ << " (";
- EmitValueList(blk->Params().Slice());
- out_ << ")";
+ if (auto* merge = blk->As<MultiInBlock>()) {
+ if (!merge->Params().IsEmpty()) {
+ out_ << " (";
+ EmitValueList(merge->Params().Slice());
+ out_ << ")";
+ }
}
out_ << " {";
diff --git a/src/tint/ir/exit_if.cc b/src/tint/ir/exit_if.cc
index 44a7002..391c786 100644
--- a/src/tint/ir/exit_if.cc
+++ b/src/tint/ir/exit_if.cc
@@ -16,8 +16,8 @@
#include <utility>
-#include "src/tint/ir/block.h"
#include "src/tint/ir/if.h"
+#include "src/tint/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitIf);
@@ -27,7 +27,7 @@
TINT_ASSERT(IR, if_);
if (if_) {
- if_->Merge()->AddInboundBranch(this);
+ if_->Merge()->AddInboundSiblingBranch(this);
}
AddOperands(std::move(args));
}
diff --git a/src/tint/ir/exit_loop.cc b/src/tint/ir/exit_loop.cc
index 729e466..865b25c 100644
--- a/src/tint/ir/exit_loop.cc
+++ b/src/tint/ir/exit_loop.cc
@@ -18,6 +18,7 @@
#include "src/tint/ir/block.h"
#include "src/tint/ir/loop.h"
+#include "src/tint/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitLoop);
@@ -28,7 +29,7 @@
TINT_ASSERT(IR, loop_);
if (loop_) {
- loop_->Merge()->AddInboundBranch(this);
+ loop_->Merge()->AddInboundSiblingBranch(this);
}
AddOperands(std::move(args));
}
diff --git a/src/tint/ir/exit_switch.cc b/src/tint/ir/exit_switch.cc
index 29c6a24..3372d6d 100644
--- a/src/tint/ir/exit_switch.cc
+++ b/src/tint/ir/exit_switch.cc
@@ -16,7 +16,7 @@
#include <utility>
-#include "src/tint/ir/block.h"
+#include "src/tint/ir/multi_in_block.h"
#include "src/tint/ir/switch.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitSwitch);
@@ -28,7 +28,7 @@
TINT_ASSERT(IR, switch_);
if (switch_) {
- switch_->Merge()->AddInboundBranch(this);
+ switch_->Merge()->AddInboundSiblingBranch(this);
}
AddOperands(std::move(args));
}
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index ad4d29f..2f01fb1 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -109,8 +109,8 @@
using ResultType = utils::Result<Module, diag::List>;
-bool IsConnected(const Block* b) {
- return b->InboundBranches().Length() > 0;
+bool IsConnected(const MultiInBlock* b) {
+ return b->InboundSiblingBranches().Length() > 0;
}
/// Impl is the private-implementation of FromProgram().
@@ -648,9 +648,6 @@
auto* loop_inst = builder_.CreateLoop();
current_block_->Append(loop_inst);
- // Loop branches directly to the body (no initializer)
- loop_inst->Body()->AddInboundBranch(loop_inst);
-
{
ControlStackScope scope(this, loop_inst);
current_block_ = loop_inst->Body();
@@ -695,9 +692,6 @@
auto* loop_inst = builder_.CreateLoop();
current_block_->Append(loop_inst);
- // Loop branches directly to the body (no initializer)
- loop_inst->Body()->AddInboundBranch(loop_inst);
-
// Continue is always empty, just go back to the start
current_block_ = loop_inst->Continuing();
SetBranch(builder_.NextIteration(loop_inst));
@@ -747,16 +741,10 @@
ControlStackScope scope(this, loop_inst);
if (stmt->initializer) {
- // Loop branches to the initializer
- loop_inst->Initializer()->AddInboundBranch(loop_inst);
-
// Emit the for initializer before branching to the body
current_block_ = loop_inst->Initializer();
EmitStatement(stmt->initializer);
SetBranch(builder_.NextIteration(loop_inst));
- } else {
- // If there's no initializer, then the loop branches directly to the body block
- loop_inst->Body()->AddInboundBranch(loop_inst);
}
current_block_ = loop_inst->Body();
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index e016481..679ec8a 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -19,6 +19,7 @@
#include "src/tint/ir/block.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/loop.h"
+#include "src/tint/ir/multi_in_block.h"
#include "src/tint/ir/program_test_helper.h"
#include "src/tint/ir/switch.h"
@@ -134,13 +135,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::If>(m);
+ auto* if_ = FindSingleValue<ir::If>(m);
ASSERT_EQ(1u, m.functions.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, if_->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -174,13 +173,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::If>(m);
+ auto* if_ = FindSingleValue<ir::If>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, if_->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -214,13 +211,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::If>(m);
+ auto* if_ = FindSingleValue<ir::If>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, if_->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -254,13 +249,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::If>(m);
+ auto* if_ = FindSingleValue<ir::If>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
- EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, if_->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -290,11 +283,6 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* if_flow = FindSingleValue<ir::If>(m);
- ASSERT_NE(if_flow, nullptr);
-
- auto* loop_flow = FindSingleValue<ir::Loop>(m);
- ASSERT_NE(loop_flow, nullptr);
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -338,13 +326,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::Loop>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->Body()->InboundBranches().Length());
- EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -374,18 +362,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* loop_flow = FindSingleValue<ir::Loop>(m);
-
- auto* if_flow = FindSingleValue<ir::If>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Body()->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(1u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -435,13 +418,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* loop_flow = FindSingleValue<ir::Loop>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Body()->InboundBranches().Length());
- EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
- EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -510,17 +493,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* loop_flow = FindSingleValue<ir::Loop>(m);
- auto* if_flow = FindSingleValue<ir::If>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, loop_flow->Body()->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(1u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -564,13 +543,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* loop_flow = FindSingleValue<ir::Loop>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Body()->InboundBranches().Length());
- EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
- EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -603,13 +582,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* loop_flow = FindSingleValue<ir::Loop>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Body()->InboundBranches().Length());
- EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
- EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -634,17 +613,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* loop_flow = FindSingleValue<ir::Loop>(m);
- auto* if_flow = FindSingleValue<ir::If>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, loop_flow->Body()->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(0u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(2u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -809,20 +784,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::Loop>(m);
-
- ASSERT_NE(flow->Body()->Branch(), nullptr);
- ASSERT_TRUE(flow->Body()->Branch()->Is<ir::If>());
- auto* if_flow = flow->Body()->Branch()->As<ir::If>();
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, flow->Body()->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(1u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -871,20 +839,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::Loop>(m);
-
- ASSERT_NE(flow->Body()->Branch(), nullptr);
- ASSERT_TRUE(flow->Body()->Branch()->Is<ir::If>());
- auto* if_flow = flow->Body()->Branch()->As<ir::If>();
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, flow->Body()->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(1u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -946,20 +907,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::Loop>(m);
-
- ASSERT_NE(flow->Body()->Branch(), nullptr);
- ASSERT_TRUE(flow->Body()->Branch()->Is<ir::If>());
- auto* if_flow = flow->Body()->Branch()->As<ir::If>();
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(2u, flow->Body()->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(1u, if_flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(2u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(2u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m), R"()");
}
@@ -972,14 +926,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::Loop>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->Initializer()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Body()->InboundBranches().Length());
- EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -1014,13 +967,13 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleValue<ir::Loop>(m);
+ auto* loop = FindSingleValue<ir::Loop>(m);
ASSERT_EQ(1u, m.functions.Length());
- EXPECT_EQ(1u, flow->Body()->InboundBranches().Length());
- EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, loop->Body()->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, loop->Continuing()->InboundSiblingBranches().Length());
+ EXPECT_EQ(1u, loop->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -1072,10 +1025,7 @@
ASSERT_EQ(1u, cases[2].selectors.Length());
EXPECT_TRUE(cases[2].selectors[0].IsDefault());
- 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(3u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(3u, flow->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -1135,8 +1085,7 @@
EXPECT_TRUE(cases[0].selectors[2].IsDefault());
- EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -1174,8 +1123,7 @@
ASSERT_EQ(1u, cases[0].selectors.Length());
EXPECT_TRUE(cases[0].selectors[0].IsDefault());
- EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
@@ -1220,9 +1168,7 @@
ASSERT_EQ(1u, cases[1].selectors.Length());
EXPECT_TRUE(cases[1].selectors[0].IsDefault());
- EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
- EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Merge()->InboundSiblingBranches().Length());
// This is 1 because the if is dead-code eliminated and the return doesn't happen.
EXPECT_EQ(Disassemble(m),
@@ -1260,7 +1206,6 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- ASSERT_EQ(FindSingleValue<ir::If>(m), nullptr);
auto* flow = FindSingleValue<ir::Switch>(m);
@@ -1276,9 +1221,7 @@
ASSERT_EQ(1u, cases[1].selectors.Length());
EXPECT_TRUE(cases[1].selectors[0].IsDefault());
- EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
- EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
- EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Merge()->InboundSiblingBranches().Length());
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc
index 29bcbb2..7f34986 100644
--- a/src/tint/ir/if.cc
+++ b/src/tint/ir/if.cc
@@ -16,11 +16,12 @@
TINT_INSTANTIATE_TYPEINFO(tint::ir::If);
-#include "src/tint/ir/block.h"
+#include "src/tint/ir/multi_in_block.h"
namespace tint::ir {
-If::If(Value* cond, ir::Block* t, ir::Block* f, ir::Block* m) : true_(t), false_(f), merge_(m) {
+If::If(Value* cond, ir::Block* t, ir::Block* f, ir::MultiInBlock* m)
+ : true_(t), false_(f), merge_(m) {
TINT_ASSERT(IR, cond);
TINT_ASSERT(IR, true_);
TINT_ASSERT(IR, false_);
@@ -28,11 +29,9 @@
AddOperand(cond);
if (true_) {
- true_->AddInboundBranch(this);
true_->SetParent(this);
}
if (false_) {
- false_->AddInboundBranch(this);
false_->SetParent(this);
}
if (merge_) {
diff --git a/src/tint/ir/if.h b/src/tint/ir/if.h
index 4566ef8..70d0288 100644
--- a/src/tint/ir/if.h
+++ b/src/tint/ir/if.h
@@ -17,6 +17,11 @@
#include "src/tint/ir/control_instruction.h"
+// Forward declarations
+namespace tint::ir {
+class MultiInBlock;
+} // namespace tint::ir
+
namespace tint::ir {
/// If instruction.
@@ -48,7 +53,7 @@
/// @param t the true block
/// @param f the false block
/// @param m the merge block
- explicit If(Value* cond, ir::Block* t, ir::Block* f, ir::Block* m);
+ If(Value* cond, ir::Block* t, ir::Block* f, ir::MultiInBlock* m);
~If() override;
/// @returns the branch arguments
@@ -70,14 +75,14 @@
ir::Block* False() { return false_; }
/// @returns the merge branch block
- const ir::Block* Merge() const { return merge_; }
+ const ir::MultiInBlock* Merge() const { return merge_; }
/// @returns the merge branch block
- ir::Block* Merge() { return merge_; }
+ ir::MultiInBlock* Merge() { return merge_; }
private:
ir::Block* true_ = nullptr;
ir::Block* false_ = nullptr;
- ir::Block* merge_ = nullptr;
+ ir::MultiInBlock* merge_ = nullptr;
};
} // namespace tint::ir
diff --git a/src/tint/ir/if_test.cc b/src/tint/ir/if_test.cc
index 4d5bcdf..032d77d 100644
--- a/src/tint/ir/if_test.cc
+++ b/src/tint/ir/if_test.cc
@@ -52,7 +52,7 @@
{
Module mod;
Builder b{mod};
- If if_(b.Constant(false), nullptr, b.CreateBlock(), b.CreateBlock());
+ If if_(b.Constant(false), nullptr, b.CreateBlock(), b.CreateMultiInBlock());
},
"");
}
@@ -62,12 +62,12 @@
{
Module mod;
Builder b{mod};
- If if_(b.Constant(false), b.CreateBlock(), nullptr, b.CreateBlock());
+ If if_(b.Constant(false), b.CreateBlock(), nullptr, b.CreateMultiInBlock());
},
"");
}
-TEST_F(IR_IfTest, Fail_NullMergeBlock) {
+TEST_F(IR_IfTest, Fail_NullMultiInBlock) {
EXPECT_FATAL_FAILURE(
{
Module mod;
diff --git a/src/tint/ir/loop.cc b/src/tint/ir/loop.cc
index 37c4f00..4e34f7d 100644
--- a/src/tint/ir/loop.cc
+++ b/src/tint/ir/loop.cc
@@ -16,13 +16,13 @@
#include <utility>
-#include "src/tint/ir/block.h"
+#include "src/tint/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::Loop);
namespace tint::ir {
-Loop::Loop(ir::Block* i, ir::Block* b, ir::Block* c, ir::Block* m)
+Loop::Loop(ir::Block* i, ir::MultiInBlock* b, ir::MultiInBlock* c, ir::MultiInBlock* m)
: initializer_(i), body_(b), continuing_(c), merge_(m) {
TINT_ASSERT(IR, initializer_);
TINT_ASSERT(IR, body_);
diff --git a/src/tint/ir/loop.h b/src/tint/ir/loop.h
index 50c69c7..88f27eb 100644
--- a/src/tint/ir/loop.h
+++ b/src/tint/ir/loop.h
@@ -17,6 +17,11 @@
#include "src/tint/ir/control_instruction.h"
+// Forward declarations
+namespace tint::ir {
+class MultiInBlock;
+} // namespace tint::ir
+
namespace tint::ir {
/// Loop instruction.
@@ -62,7 +67,7 @@
/// @param b the body block
/// @param c the continuing block
/// @param m the merge block
- Loop(ir::Block* i, ir::Block* b, ir::Block* c, ir::Block* m);
+ Loop(ir::Block* i, ir::MultiInBlock* b, ir::MultiInBlock* c, ir::MultiInBlock* m);
~Loop() override;
/// @returns the switch initializer block
@@ -75,25 +80,25 @@
bool HasInitializer() const;
/// @returns the switch start block
- const ir::Block* Body() const { return body_; }
+ const ir::MultiInBlock* Body() const { return body_; }
/// @returns the switch start block
- ir::Block* Body() { return body_; }
+ ir::MultiInBlock* Body() { return body_; }
/// @returns the switch continuing block
- const ir::Block* Continuing() const { return continuing_; }
+ const ir::MultiInBlock* Continuing() const { return continuing_; }
/// @returns the switch continuing block
- ir::Block* Continuing() { return continuing_; }
+ ir::MultiInBlock* Continuing() { return continuing_; }
/// @returns the switch merge branch
- const ir::Block* Merge() const { return merge_; }
+ const ir::MultiInBlock* Merge() const { return merge_; }
/// @returns the switch merge branch
- ir::Block* Merge() { return merge_; }
+ ir::MultiInBlock* Merge() { return merge_; }
private:
ir::Block* initializer_ = nullptr;
- ir::Block* body_ = nullptr;
- ir::Block* continuing_ = nullptr;
- ir::Block* merge_ = nullptr;
+ ir::MultiInBlock* body_ = nullptr;
+ ir::MultiInBlock* continuing_ = nullptr;
+ ir::MultiInBlock* merge_ = nullptr;
};
} // namespace tint::ir
diff --git a/src/tint/ir/loop_test.cc b/src/tint/ir/loop_test.cc
index e51a005..e063e7e 100644
--- a/src/tint/ir/loop_test.cc
+++ b/src/tint/ir/loop_test.cc
@@ -35,7 +35,8 @@
{
Module mod;
Builder b{mod};
- Loop loop(nullptr, b.CreateBlock(), b.CreateBlock(), b.CreateBlock());
+ Loop loop(nullptr, b.CreateMultiInBlock(), b.CreateMultiInBlock(),
+ b.CreateMultiInBlock());
},
"");
}
@@ -45,7 +46,7 @@
{
Module mod;
Builder b{mod};
- Loop loop(b.CreateBlock(), nullptr, b.CreateBlock(), b.CreateBlock());
+ Loop loop(b.CreateBlock(), nullptr, b.CreateMultiInBlock(), b.CreateMultiInBlock());
},
"");
}
@@ -55,17 +56,17 @@
{
Module mod;
Builder b{mod};
- Loop loop(b.CreateBlock(), b.CreateBlock(), nullptr, b.CreateBlock());
+ Loop loop(b.CreateBlock(), b.CreateMultiInBlock(), nullptr, b.CreateMultiInBlock());
},
"");
}
-TEST_F(IR_LoopTest, Fail_NullMergeBlock) {
+TEST_F(IR_LoopTest, Fail_NullMultiInBlock) {
EXPECT_FATAL_FAILURE(
{
Module mod;
Builder b{mod};
- Loop loop(b.CreateBlock(), b.CreateBlock(), b.CreateBlock(), nullptr);
+ Loop loop(b.CreateBlock(), b.CreateMultiInBlock(), b.CreateMultiInBlock(), nullptr);
},
"");
}
diff --git a/src/tint/ir/multi_in_block.cc b/src/tint/ir/multi_in_block.cc
new file mode 100644
index 0000000..9844b3a
--- /dev/null
+++ b/src/tint/ir/multi_in_block.cc
@@ -0,0 +1,41 @@
+// 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/multi_in_block.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::MultiInBlock);
+
+namespace tint::ir {
+
+MultiInBlock::MultiInBlock() : Base() {}
+
+MultiInBlock::~MultiInBlock() = default;
+
+void MultiInBlock::SetParams(utils::VectorRef<const BlockParam*> params) {
+ params_ = std::move(params);
+
+ for (auto* param : params_) {
+ TINT_ASSERT(IR, param != nullptr);
+ }
+}
+
+void MultiInBlock::AddInboundSiblingBranch(ir::Branch* node) {
+ TINT_ASSERT(IR, node != nullptr);
+
+ if (node) {
+ inbound_sibling_branches_.Push(node);
+ }
+}
+
+} // namespace tint::ir
diff --git a/src/tint/ir/multi_in_block.h b/src/tint/ir/multi_in_block.h
new file mode 100644
index 0000000..4a51901
--- /dev/null
+++ b/src/tint/ir/multi_in_block.h
@@ -0,0 +1,66 @@
+// Copyright 2022 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_IR_MULTI_IN_BLOCK_H_
+#define SRC_TINT_IR_MULTI_IN_BLOCK_H_
+
+#include <utility>
+
+#include "src/tint/ir/block.h"
+
+// Forward declarations
+namespace tint::ir {
+class BlockParam;
+}
+
+namespace tint::ir {
+
+/// A block that can be the target of multiple branches.
+/// MultiInBlocks maintain a list of inbound branches from branch instructions excluding ir::If,
+/// ir::Switch and ir::Loop which implicitly branch to the internal block.
+/// MultiInBlocks hold a number of BlockParam parameters, used to pass values from the branch source
+/// to this target.
+class MultiInBlock : public utils::Castable<MultiInBlock, Block> {
+ public:
+ /// Constructor
+ MultiInBlock();
+ ~MultiInBlock() override;
+
+ /// Sets the params to the block
+ /// @param params the params for the block
+ void SetParams(utils::VectorRef<const BlockParam*> params);
+
+ /// @return the parameters passed into the block
+ utils::VectorRef<const BlockParam*> Params() const { return params_; }
+
+ /// @returns the params to the block
+ utils::Vector<const BlockParam*, 2>& Params() { return params_; }
+
+ /// @returns branches made to this block by sibling blocks
+ utils::VectorRef<ir::Branch*> InboundSiblingBranches() const {
+ return inbound_sibling_branches_;
+ }
+
+ /// Adds the given branch to the list of branches made to this block by sibling blocks
+ /// @param branch the branch to add
+ void AddInboundSiblingBranch(ir::Branch* branch);
+
+ private:
+ utils::Vector<const BlockParam*, 2> params_;
+ utils::Vector<ir::Branch*, 2> inbound_sibling_branches_;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_MULTI_IN_BLOCK_H_
diff --git a/src/tint/ir/multi_in_block_test.cc b/src/tint/ir/multi_in_block_test.cc
new file mode 100644
index 0000000..b5d5442
--- /dev/null
+++ b/src/tint/ir/multi_in_block_test.cc
@@ -0,0 +1,51 @@
+// 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/multi_in_block.h"
+#include "gtest/gtest-spi.h"
+#include "src/tint/ir/block_param.h"
+#include "src/tint/ir/ir_test_helper.h"
+
+namespace tint::ir {
+namespace {
+
+using namespace tint::number_suffixes; // NOLINT
+using IR_MultiInBlockTest = IRTestHelper;
+
+TEST_F(IR_MultiInBlockTest, Fail_NullBlockParam) {
+ EXPECT_FATAL_FAILURE(
+ {
+ Module mod;
+ Builder b{mod};
+
+ auto* blk = b.CreateMultiInBlock();
+ blk->SetParams(utils::Vector<const BlockParam*, 1>{nullptr});
+ },
+ "");
+}
+
+TEST_F(IR_MultiInBlockTest, Fail_NullInboundBranch) {
+ EXPECT_FATAL_FAILURE(
+ {
+ Module mod;
+ Builder b{mod};
+
+ auto* blk = b.CreateMultiInBlock();
+ blk->AddInboundSiblingBranch(nullptr);
+ },
+ "");
+}
+
+} // namespace
+} // namespace tint::ir
diff --git a/src/tint/ir/next_iteration.cc b/src/tint/ir/next_iteration.cc
index b1bf620..dc1ccd2 100644
--- a/src/tint/ir/next_iteration.cc
+++ b/src/tint/ir/next_iteration.cc
@@ -16,8 +16,8 @@
#include <utility>
-#include "src/tint/ir/block.h"
#include "src/tint/ir/loop.h"
+#include "src/tint/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::NextIteration);
@@ -26,9 +26,8 @@
NextIteration::NextIteration(ir::Loop* loop, utils::VectorRef<Value*> args /* = utils::Empty */)
: loop_(loop) {
TINT_ASSERT(IR, loop_);
-
if (loop_) {
- loop_->Body()->AddInboundBranch(this);
+ loop_->Body()->AddInboundSiblingBranch(this);
}
AddOperands(std::move(args));
}
diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc
index 68137ac..8b70cfc 100644
--- a/src/tint/ir/switch.cc
+++ b/src/tint/ir/switch.cc
@@ -16,11 +16,11 @@
TINT_INSTANTIATE_TYPEINFO(tint::ir::Switch);
-#include "src/tint/ir/block.h"
+#include "src/tint/ir/multi_in_block.h"
namespace tint::ir {
-Switch::Switch(Value* cond, ir::Block* m) : merge_(m) {
+Switch::Switch(Value* cond, ir::MultiInBlock* m) : merge_(m) {
TINT_ASSERT(IR, cond);
TINT_ASSERT(IR, merge_);
diff --git a/src/tint/ir/switch.h b/src/tint/ir/switch.h
index b065074dc..56d6abe 100644
--- a/src/tint/ir/switch.h
+++ b/src/tint/ir/switch.h
@@ -20,6 +20,7 @@
// Forward declarations
namespace tint::ir {
class Constant;
+class MultiInBlock;
} // namespace tint::ir
namespace tint::ir {
@@ -70,13 +71,13 @@
/// Constructor
/// @param cond the condition
/// @param m the merge block
- explicit Switch(Value* cond, ir::Block* m);
+ explicit Switch(Value* cond, ir::MultiInBlock* m);
~Switch() override;
/// @returns the switch merge branch
- const ir::Block* Merge() const { return merge_; }
+ const ir::MultiInBlock* Merge() const { return merge_; }
/// @returns the switch merge branch
- ir::Block* Merge() { return merge_; }
+ ir::MultiInBlock* Merge() { return merge_; }
/// @returns the switch cases
utils::VectorRef<Case> Cases() const { return cases_; }
@@ -92,7 +93,7 @@
Value* Condition() { return operands_[0]; }
private:
- ir::Block* merge_ = nullptr;
+ ir::MultiInBlock* merge_ = nullptr;
utils::Vector<Case, 4> cases_;
};
diff --git a/src/tint/ir/switch_test.cc b/src/tint/ir/switch_test.cc
index cbcf11a..8dd0fb5 100644
--- a/src/tint/ir/switch_test.cc
+++ b/src/tint/ir/switch_test.cc
@@ -47,7 +47,7 @@
"");
}
-TEST_F(IR_SwitchTest, Fail_NullMergeBlock) {
+TEST_F(IR_SwitchTest, Fail_NullMultiInBlock) {
EXPECT_FATAL_FAILURE(
{
Module mod;
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index b226236..a041ca7 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -26,6 +26,7 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/load.h"
#include "src/tint/ir/module.h"
+#include "src/tint/ir/multi_in_block.h"
#include "src/tint/ir/return.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h"
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
index a173b49..b61dc21 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -22,6 +22,7 @@
#include "src/tint/ir/access.h"
#include "src/tint/ir/binary.h"
#include "src/tint/ir/block.h"
+#include "src/tint/ir/block_param.h"
#include "src/tint/ir/break_if.h"
#include "src/tint/ir/builtin.h"
#include "src/tint/ir/continue.h"
@@ -32,6 +33,7 @@
#include "src/tint/ir/load.h"
#include "src/tint/ir/loop.h"
#include "src/tint/ir/module.h"
+#include "src/tint/ir/multi_in_block.h"
#include "src/tint/ir/next_iteration.h"
#include "src/tint/ir/return.h"
#include "src/tint/ir/store.h"
@@ -426,20 +428,22 @@
return;
}
- // Emit all OpPhi nodes for incoming branches to block.
- EmitIncomingPhis(block);
+ if (auto* mib = block->As<ir::MultiInBlock>()) {
+ // Emit all OpPhi nodes for incoming branches to block.
+ EmitIncomingPhis(mib);
+ }
// Emit the block's statements.
EmitBlockInstructions(block);
}
-void GeneratorImplIr::EmitIncomingPhis(const ir::Block* block) {
+void GeneratorImplIr::EmitIncomingPhis(const ir::MultiInBlock* block) {
// Emit Phi nodes for all the incoming block parameters
for (size_t param_idx = 0; param_idx < block->Params().Length(); param_idx++) {
auto* param = block->Params()[param_idx];
OperandList ops{Type(param->Type()), Value(param)};
- for (auto* incoming : block->InboundBranches()) {
+ for (auto* incoming : block->InboundSiblingBranches()) {
auto* arg = incoming->Args()[param_idx];
ops.push_back(Value(arg));
ops.push_back(Label(incoming->Block()));
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.h b/src/tint/writer/spirv/ir/generator_impl_ir.h
index d3abeb6..1978d41 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.h
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.h
@@ -35,11 +35,12 @@
class BlockParam;
class Branch;
class Builtin;
-class If;
class Function;
+class If;
class Load;
class Loop;
class Module;
+class MultiInBlock;
class Store;
class Switch;
class UserCall;
@@ -119,7 +120,7 @@
/// Emit all OpPhi nodes for incoming branches to @p block.
/// @param block the block to emit the OpPhis for
- void EmitIncomingPhis(const ir::Block* block);
+ void EmitIncomingPhis(const ir::MultiInBlock* block);
/// Emit all instructions of @p block.
/// @param block the block's instructions to emit
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
index 807833f..41d1f93 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_loop_test.cc
@@ -352,7 +352,6 @@
auto* l = b.CreateLoop();
func->StartTarget()->Append(l);
- l->Initializer()->AddInboundBranch(l);
l->Initializer()->Append(b.NextIteration(l, utils::Vector{b.Constant(1_i)}));
auto* loop_param = b.BlockParam(b.ir.Types().i32());
@@ -405,7 +404,6 @@
auto* l = b.CreateLoop();
func->StartTarget()->Append(l);
- l->Initializer()->AddInboundBranch(l);
l->Initializer()->Append(b.NextIteration(l, utils::Vector{b.Constant(1_i), b.Constant(false)}));
auto* loop_param_a = b.BlockParam(b.ir.Types().i32());
