Import Tint changes from Dawn
Changes:
- a1c6c504cb67019e25f05939b6730aedb1540fe6 [ir] Check operands exist before adding/removing by dan sinclair <dsinclair@chromium.org>
- eac9ac23dee9491cfb62f90d18d6ff77346da2e6 [tint][ir][ToProgram] Test compound assignments by Ben Clayton <bclayton@google.com>
- e692fa7b6e76099a94974529a691a32672fe24bf [tint][ir][transform] Fix MergeReturn test by Ben Clayton <bclayton@google.com>
- b274678f4a243670cd5615ad801f647cf5aea283 [tint][ir] Rename Branch to Terminator by Ben Clayton <bclayton@google.com>
- 0fde633f706b6ce7f8dc40c5fd8c8111fbdc95de [tint][ir] Add comments for exit instructions by Ben Clayton <bclayton@google.com>
- 293e8ae22d0391c3568e1765d5a0117976a9e041 [tint][ir] Add implicit false-block 'undef' value comment by Ben Clayton <bclayton@google.com>
- 8c36b7035dadafe5409c2738ca8734d00150f3c5 [tint][ir] Reformat Block & ControlInstruction disassembly by Ben Clayton <bclayton@google.com>
- eb059e04284052f024287330a97aeaddb6dacf24 [tint][utils] Enable TINT_ASSERT_ITERATORS_NOT_INVALIDATED by Ben Clayton <bclayton@google.com>
- e3feea76729d52c050908e5f8588ee0b964112ff [tint][ir][to_program] Implement short circuit by Ben Clayton <bclayton@google.com>
- f8f852430c6b00a256686e249c03b596386bc8ee [tint][ir] Refactor ToProgram by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: a1c6c504cb67019e25f05939b6730aedb1540fe6
Change-Id: Id8686e57f9ed9544cf50a398c65e2e6ad8a300ed
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/138500
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 1d71177..ae14e47 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1234,8 +1234,6 @@
"ir/block.h",
"ir/block_param.cc",
"ir/block_param.h",
- "ir/branch.cc",
- "ir/branch.h",
"ir/break_if.cc",
"ir/break_if.h",
"ir/builder.cc",
@@ -1297,6 +1295,8 @@
"ir/switch.h",
"ir/swizzle.cc",
"ir/swizzle.h",
+ "ir/terminator.cc",
+ "ir/terminator.h",
"ir/transform/transform.cc",
"ir/transform/transform.h",
"ir/unary.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 5c0bf6f..0b4cd81 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -734,8 +734,6 @@
ir/block.h
ir/block_param.cc
ir/block_param.h
- ir/branch.cc
- ir/branch.h
ir/break_if.cc
ir/break_if.h
ir/builder.cc
@@ -799,6 +797,8 @@
ir/switch.h
ir/swizzle.cc
ir/swizzle.h
+ ir/terminator.cc
+ ir/terminator.h
ir/to_program.cc
ir/to_program.h
ir/unary.cc
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 60f6b31..8c9e06a 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -17,8 +17,8 @@
#include <utility>
-#include "src/tint/ir/branch.h"
#include "src/tint/ir/instruction.h"
+#include "src/tint/ir/terminator.h"
#include "src/tint/utils/vector.h"
// Forward declarations
@@ -29,25 +29,24 @@
namespace tint::ir {
/// A block of statements. The instructions in the block are a linear list of instructions to
-/// execute. The block will branch at the end. The only blocks which do not branch are the end
-/// blocks of functions.
+/// execute. The block will terminate with a Terminator instruction at the end.
class Block : public utils::Castable<Block> {
public:
/// Constructor
Block();
~Block() override;
- /// @returns true if this is block has a branch target set
- bool HasBranchTarget() {
- return instructions_.last != nullptr && instructions_.last->Is<ir::Branch>();
+ /// @returns true if this is block has a terminator instruction
+ bool HasTerminator() {
+ return instructions_.last != nullptr && instructions_.last->Is<ir::Terminator>();
}
- /// @return the node this block branches to or nullptr if the block doesn't branch
- ir::Branch* Branch() {
- if (!HasBranchTarget()) {
+ /// @return the terminator instruction for this block
+ ir::Terminator* Terminator() {
+ if (!HasTerminator()) {
return nullptr;
}
- return instructions_.last->As<ir::Branch>();
+ return instructions_.last->As<ir::Terminator>();
}
/// @returns the instructions in the block
diff --git a/src/tint/ir/block_test.cc b/src/tint/ir/block_test.cc
index e9523c5..9f5fbfa 100644
--- a/src/tint/ir/block_test.cc
+++ b/src/tint/ir/block_test.cc
@@ -22,65 +22,65 @@
using namespace tint::number_suffixes; // NOLINT
using IR_BlockTest = IRTestHelper;
-TEST_F(IR_BlockTest, HasBranchTarget_Empty) {
+TEST_F(IR_BlockTest, HasTerminator_Empty) {
auto* blk = b.Block();
- EXPECT_FALSE(blk->HasBranchTarget());
+ EXPECT_FALSE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_NoBranch) {
+TEST_F(IR_BlockTest, HasTerminator_None) {
auto* blk = b.Block();
blk->Append(b.Add(mod.Types().i32(), 1_u, 2_u));
- EXPECT_FALSE(blk->HasBranchTarget());
+ EXPECT_FALSE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_BreakIf) {
+TEST_F(IR_BlockTest, HasTerminator_BreakIf) {
auto* blk = b.Block();
auto* loop = b.Loop();
blk->Append(b.BreakIf(true, loop));
- EXPECT_TRUE(blk->HasBranchTarget());
+ EXPECT_TRUE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_Continue) {
+TEST_F(IR_BlockTest, HasTerminator_Continue) {
auto* blk = b.Block();
auto* loop = b.Loop();
blk->Append(b.Continue(loop));
- EXPECT_TRUE(blk->HasBranchTarget());
+ EXPECT_TRUE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_ExitIf) {
+TEST_F(IR_BlockTest, HasTerminator_ExitIf) {
auto* blk = b.Block();
auto* if_ = b.If(true);
blk->Append(b.ExitIf(if_));
- EXPECT_TRUE(blk->HasBranchTarget());
+ EXPECT_TRUE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_ExitLoop) {
+TEST_F(IR_BlockTest, HasTerminator_ExitLoop) {
auto* blk = b.Block();
auto* loop = b.Loop();
blk->Append(b.ExitLoop(loop));
- EXPECT_TRUE(blk->HasBranchTarget());
+ EXPECT_TRUE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_ExitSwitch) {
+TEST_F(IR_BlockTest, HasTerminator_ExitSwitch) {
auto* blk = b.Block();
auto* s = b.Switch(1_u);
blk->Append(b.ExitSwitch(s));
- EXPECT_TRUE(blk->HasBranchTarget());
+ EXPECT_TRUE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_NextIteration) {
+TEST_F(IR_BlockTest, HasTerminator_NextIteration) {
auto* blk = b.Block();
auto* loop = b.Loop();
blk->Append(b.NextIteration(loop));
- EXPECT_TRUE(blk->HasBranchTarget());
+ EXPECT_TRUE(blk->HasTerminator());
}
-TEST_F(IR_BlockTest, HasBranchTarget_Return) {
+TEST_F(IR_BlockTest, HasTerminator_Return) {
auto* f = b.Function("myFunc", mod.Types().void_());
auto* blk = b.Block();
blk->Append(b.Return(f));
- EXPECT_TRUE(blk->HasBranchTarget());
+ EXPECT_TRUE(blk->HasTerminator());
}
TEST_F(IR_BlockTest, Append) {
diff --git a/src/tint/ir/break_if.h b/src/tint/ir/break_if.h
index 59046b2..7078bc4 100644
--- a/src/tint/ir/break_if.h
+++ b/src/tint/ir/break_if.h
@@ -15,7 +15,7 @@
#ifndef SRC_TINT_IR_BREAK_IF_H_
#define SRC_TINT_IR_BREAK_IF_H_
-#include "src/tint/ir/branch.h"
+#include "src/tint/ir/terminator.h"
#include "src/tint/ir/value.h"
#include "src/tint/utils/castable.h"
@@ -27,7 +27,7 @@
namespace tint::ir {
/// A break-if iteration instruction.
-class BreakIf : public utils::Castable<BreakIf, Branch> {
+class BreakIf : public utils::Castable<BreakIf, Terminator> {
public:
/// The offset in Operands() for the condition
static constexpr size_t kConditionOperandOffset = 0;
@@ -38,11 +38,11 @@
/// Constructor
/// @param condition the break condition
/// @param loop the loop containing the break-if
- /// @param args the branch arguments
+ /// @param args the MultiInBlock arguments
BreakIf(Value* condition, ir::Loop* loop, utils::VectorRef<Value*> args = utils::Empty);
~BreakIf() override;
- /// @returns the branch arguments
+ /// @returns the MultiInBlock arguments
utils::Slice<Value* const> Args() override {
return operands_.Slice().Offset(kArgsOperandOffset);
}
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 9a2c817..8db4098 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -109,48 +109,48 @@
/// @returns a new multi-in block
ir::MultiInBlock* MultiInBlock();
- /// Creates a function flow node
+ /// Creates a function instruction
/// @param name the function name
/// @param return_type the function return type
/// @param stage the function stage
/// @param wg_size the workgroup_size
- /// @returns the flow node
+ /// @returns the instruction
ir::Function* Function(std::string_view name,
const type::Type* return_type,
Function::PipelineStage stage = Function::PipelineStage::kUndefined,
std::optional<std::array<uint32_t, 3>> wg_size = {});
- /// Creates an if flow node
+ /// Creates an if instruction
/// @param condition the if condition
- /// @returns the flow node
+ /// @returns the instruction
template <typename T>
ir::If* If(T&& condition) {
return Append(
ir.instructions.Create<ir::If>(Value(std::forward<T>(condition)), Block(), Block()));
}
- /// Creates a loop flow node
- /// @returns the flow node
+ /// Creates a loop instruction
+ /// @returns the instruction
ir::Loop* Loop();
- /// Creates a switch flow node
+ /// Creates a switch instruction
/// @param condition the switch condition
- /// @returns the flow node
+ /// @returns the instruction
template <typename T>
ir::Switch* Switch(T&& condition) {
return Append(ir.instructions.Create<ir::Switch>(Value(std::forward<T>(condition))));
}
- /// Creates a case flow node for the given case branch.
+ /// Creates a case for the switch @p s with the given selectors
/// @param s the switch to create the case into
/// @param selectors the case selectors for the case statement
- /// @returns the start block for the case flow node
+ /// @returns the start block for the case instruction
ir::Block* Case(ir::Switch* s, utils::VectorRef<Switch::CaseSelector> selectors);
- /// Creates a case flow node for the given case branch.
+ /// Creates a case for the switch @p s with the given selectors
/// @param s the switch to create the case into
/// @param selectors the case selectors for the case statement
- /// @returns the start block for the case flow node
+ /// @returns the start block for the case instruction
ir::Block* Case(ir::Switch* s, std::initializer_list<Switch::CaseSelector> selectors);
/// Creates a new ir::Constant
@@ -563,7 +563,7 @@
/// Creates a loop next iteration instruction
/// @param loop the loop being iterated
- /// @param args the branch arguments
+ /// @param args the arguments for the target MultiInBlock
/// @returns the instruction
template <typename... ARGS>
ir::NextIteration* NextIteration(ir::Loop* loop, ARGS&&... args) {
@@ -574,7 +574,7 @@
/// Creates a loop break-if instruction
/// @param condition the break condition
/// @param loop the loop being iterated
- /// @param args the branch arguments
+ /// @param args the arguments for the target MultiInBlock
/// @returns the instruction
template <typename CONDITION, typename... ARGS>
ir::BreakIf* BreakIf(CONDITION&& condition, ir::Loop* loop, ARGS&&... args) {
@@ -584,7 +584,7 @@
/// Creates a continue instruction
/// @param loop the loop being continued
- /// @param args the branch arguments
+ /// @param args the arguments for the target MultiInBlock
/// @returns the instruction
template <typename... ARGS>
ir::Continue* Continue(ir::Loop* loop, ARGS&&... args) {
@@ -594,7 +594,7 @@
/// Creates an exit switch instruction
/// @param sw the switch being exited
- /// @param args the branch arguments
+ /// @param args the arguments for the target MultiInBlock
/// @returns the instruction
template <typename... ARGS>
ir::ExitSwitch* ExitSwitch(ir::Switch* sw, ARGS&&... args) {
@@ -604,7 +604,7 @@
/// Creates an exit loop instruction
/// @param loop the loop being exited
- /// @param args the branch arguments
+ /// @param args the arguments for the target MultiInBlock
/// @returns the instruction
template <typename... ARGS>
ir::ExitLoop* ExitLoop(ir::Loop* loop, ARGS&&... args) {
@@ -614,7 +614,7 @@
/// Creates an exit if instruction
/// @param i the if being exited
- /// @param args the branch arguments
+ /// @param args the arguments for the target MultiInBlock
/// @returns the instruction
template <typename... ARGS>
ir::ExitIf* ExitIf(ir::If* i, ARGS&&... args) {
@@ -623,10 +623,10 @@
/// Creates an exit instruction for the given control instruction
/// @param inst the control instruction being exited
- /// @param args the branch arguments
+ /// @param args the arguments for the target MultiInBlock
/// @returns the exit instruction, or nullptr if the control instruction is not supported.
template <typename... ARGS>
- ir::Branch* Exit(ir::ControlInstruction* inst, ARGS&&... args) {
+ ir::Exit* Exit(ir::ControlInstruction* inst, ARGS&&... args) {
return tint::Switch(
inst, //
[&](ir::If* i) { return ExitIf(i, std::forward<ARGS>(args)...); },
diff --git a/src/tint/ir/continue.h b/src/tint/ir/continue.h
index 6beda97..2f3367c 100644
--- a/src/tint/ir/continue.h
+++ b/src/tint/ir/continue.h
@@ -15,7 +15,7 @@
#ifndef SRC_TINT_IR_CONTINUE_H_
#define SRC_TINT_IR_CONTINUE_H_
-#include "src/tint/ir/branch.h"
+#include "src/tint/ir/terminator.h"
#include "src/tint/utils/castable.h"
// Forward declarations
@@ -26,14 +26,14 @@
namespace tint::ir {
/// A continue instruction.
-class Continue : public utils::Castable<Continue, Branch> {
+class Continue : public utils::Castable<Continue, Terminator> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
/// Constructor
/// @param loop the loop owning the continue block
- /// @param args the branch arguments
+ /// @param args the arguments for the MultiInBlock
explicit Continue(ir::Loop* loop, utils::VectorRef<Value*> args = utils::Empty);
~Continue() override;
diff --git a/src/tint/ir/control_instruction.h b/src/tint/ir/control_instruction.h
index ebe5825..701e032 100644
--- a/src/tint/ir/control_instruction.h
+++ b/src/tint/ir/control_instruction.h
@@ -17,7 +17,6 @@
#include <utility>
-#include "src/tint/ir/branch.h"
#include "src/tint/ir/operand_instruction.h"
// Forward declarations
@@ -28,7 +27,7 @@
namespace tint::ir {
-/// Base class of instructions that perform branches to two or more blocks, owned by the
+/// Base class of instructions that perform control flow to two or more blocks, owned by the
/// ControlInstruction.
class ControlInstruction : public utils::Castable<ControlInstruction, OperandInstruction<1, 1>> {
public:
@@ -60,7 +59,7 @@
SetResults(utils::Vector{std::forward<ARGS>(values)...});
}
- /// @return All the exit branches for the flow control instruction
+ /// @return All the exits for the flow control instruction
const utils::Hashset<Exit*, 2>& Exits() const { return exits_; }
/// Adds the exit to the flow control instruction
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 27598f0..629146a 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -99,6 +99,23 @@
});
}
+std::string_view Disassembler::NameOf(If* inst) {
+ TINT_ASSERT(IR, inst);
+ return if_names_.GetOrCreate(inst, [&] { return "if_" + std::to_string(if_names_.Count()); });
+}
+
+std::string_view Disassembler::NameOf(Loop* inst) {
+ TINT_ASSERT(IR, inst);
+ return loop_names_.GetOrCreate(inst,
+ [&] { return "loop_" + std::to_string(loop_names_.Count()); });
+}
+
+std::string_view Disassembler::NameOf(Switch* inst) {
+ TINT_ASSERT(IR, inst);
+ return switch_names_.GetOrCreate(
+ inst, [&] { return "switch_" + std::to_string(switch_names_.Count()); });
+}
+
Source::Location Disassembler::MakeCurrentLocation() {
return Source::Location{current_output_line_, out_.tellp() - current_output_start_pos_ + 1};
}
@@ -111,9 +128,7 @@
}
if (mod_.root_block) {
- Indent() << "# Root block";
- EmitLine();
- EmitBlock(mod_.root_block);
+ EmitBlock(mod_.root_block, "root");
EmitLine();
}
@@ -123,7 +138,7 @@
return out_.str();
}
-void Disassembler::EmitBlock(Block* blk) {
+void Disassembler::EmitBlock(Block* blk, std::string_view comment /* = "" */) {
Indent();
SourceMarker sm(this);
@@ -138,6 +153,10 @@
sm.Store(blk);
out_ << " {";
+ if (!comment.empty()) {
+ out_ << " # " << comment;
+ }
+
EmitLine();
{
ScopedIndent si(indent_size_);
@@ -468,7 +487,7 @@
}
EmitLine();
},
- [&](Branch* b) { EmitBranch(b); },
+ [&](Terminator* b) { EmitTerminator(b); },
[&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; });
}
@@ -515,24 +534,29 @@
out_ << "]";
sm.Store(i);
+ out_ << " { # " << NameOf(i);
EmitLine();
if (has_true) {
ScopedIndent si(indent_size_);
- Indent() << "# True block";
- EmitLine();
-
- EmitBlock(i->True());
- EmitLine();
+ EmitBlock(i->True(), "true");
}
if (has_false) {
ScopedIndent si(indent_size_);
- Indent() << "# False block";
- EmitLine();
-
- EmitBlock(i->False());
+ EmitBlock(i->False(), "false");
+ } else if (i->HasResults()) {
+ ScopedIndent si(indent_size_);
+ Indent();
+ out_ << "# implicit false block: exit_if undef";
+ for (size_t v = 1; v < i->Results().Length(); v++) {
+ out_ << ", undef";
+ }
EmitLine();
}
+
+ Indent();
+ out_ << "}";
+ EmitLine();
}
void Disassembler::EmitLoop(Loop* l) {
@@ -551,31 +575,27 @@
out_ << "loop [" << utils::Join(parts, ", ") << "]";
sm.Store(l);
+ out_ << " { # " << NameOf(l);
EmitLine();
if (!l->Initializer()->IsEmpty()) {
ScopedIndent si(indent_size_);
- Indent() << "# Initializer block";
- EmitLine();
- EmitBlock(l->Initializer());
- EmitLine();
+ EmitBlock(l->Initializer(), "initializer");
}
if (!l->Body()->IsEmpty()) {
ScopedIndent si(indent_size_);
- Indent() << "# Body block";
- EmitLine();
- EmitBlock(l->Body());
- EmitLine();
+ EmitBlock(l->Body(), "body");
}
if (!l->Continuing()->IsEmpty()) {
ScopedIndent si(indent_size_);
- Indent() << "# Continuing block";
- EmitLine();
- EmitBlock(l->Continuing());
- EmitLine();
+ EmitBlock(l->Continuing(), "continuing");
}
+
+ Indent();
+ out_ << "}";
+ EmitLine();
}
void Disassembler::EmitSwitch(Switch* s) {
@@ -600,20 +620,20 @@
}
out_ << ", %b" << IdOf(c.Block()) << ")";
}
- out_ << "]";
+ out_ << "] { # " << NameOf(s);
EmitLine();
for (auto& c : s->Cases()) {
ScopedIndent si(indent_size_);
- Indent() << "# Case block";
- EmitLine();
-
- EmitBlock(c.Block());
- EmitLine();
+ EmitBlock(c.Block(), "case");
}
+
+ Indent();
+ out_ << "}";
+ EmitLine();
}
-void Disassembler::EmitBranch(Branch* b) {
+void Disassembler::EmitTerminator(Terminator* b) {
SourceMarker sm(this);
tint::Switch(
b, //
@@ -630,7 +650,7 @@
out_ << " %b" << IdOf(bi->Loop()->Body());
},
[&](Unreachable*) { out_ << "unreachable"; },
- [&](Default) { out_ << "Unknown branch " << b->TypeInfo().name; });
+ [&](Default) { out_ << "unknown terminator " << b->TypeInfo().name; });
if (!b->Args().IsEmpty()) {
out_ << " ";
@@ -638,6 +658,12 @@
}
sm.Store(b);
+ tint::Switch(
+ b, //
+ [&](ir::ExitIf* e) { out_ << " # " << NameOf(e->If()); }, //
+ [&](ir::ExitSwitch* e) { out_ << " # " << NameOf(e->Switch()); }, //
+ [&](ir::ExitLoop* e) { out_ << " # " << NameOf(e->Loop()); } //
+ );
EmitLine();
}
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index 71773f1..940f7f6 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -108,8 +108,11 @@
size_t IdOf(Block* blk);
std::string_view IdOf(Value* node);
+ std::string_view NameOf(If* inst);
+ std::string_view NameOf(Loop* inst);
+ std::string_view NameOf(Switch* inst);
- void EmitBlock(Block* blk);
+ void EmitBlock(Block* blk, std::string_view comment = "");
void EmitFunction(Function* func);
void EmitParamAttributes(FunctionParam* p);
void EmitReturnAttributes(Function* func);
@@ -123,7 +126,7 @@
void EmitArgs(Call* call);
void EmitBinary(Binary* b);
void EmitUnary(Unary* b);
- void EmitBranch(Branch* b);
+ void EmitTerminator(Terminator* b);
void EmitSwitch(Switch* s);
void EmitLoop(Loop* l);
void EmitIf(If* i);
@@ -148,6 +151,9 @@
utils::Hashmap<Block*, Source, 8> block_to_src_;
utils::Hashmap<Instruction*, Source, 8> instruction_to_src_;
utils::Hashmap<Usage, Source, 8, Usage::Hasher> operand_to_src_;
+ utils::Hashmap<If*, std::string, 8> if_names_;
+ utils::Hashmap<Loop*, std::string, 8> loop_names_;
+ utils::Hashmap<Switch*, std::string, 8> switch_names_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/exit.h b/src/tint/ir/exit.h
index f9734fb..b8b8db1 100644
--- a/src/tint/ir/exit.h
+++ b/src/tint/ir/exit.h
@@ -15,7 +15,7 @@
#ifndef SRC_TINT_IR_EXIT_H_
#define SRC_TINT_IR_EXIT_H_
-#include "src/tint/ir/branch.h"
+#include "src/tint/ir/terminator.h"
// Forward declarations
namespace tint::ir {
@@ -25,7 +25,7 @@
namespace tint::ir {
/// The base class for all exit terminators.
-class Exit : public utils::Castable<Exit, Branch> {
+class Exit : public utils::Castable<Exit, Terminator> {
public:
~Exit() override;
diff --git a/src/tint/ir/exit_if.h b/src/tint/ir/exit_if.h
index e3cf9e7..0b89d0a 100644
--- a/src/tint/ir/exit_if.h
+++ b/src/tint/ir/exit_if.h
@@ -33,7 +33,7 @@
/// Constructor
/// @param i the if being exited
- /// @param args the branch arguments
+ /// @param args the target MultiInBlock arguments
explicit ExitIf(ir::If* i, utils::VectorRef<Value*> args = utils::Empty);
~ExitIf() override;
diff --git a/src/tint/ir/exit_loop.h b/src/tint/ir/exit_loop.h
index 253655e..3931287 100644
--- a/src/tint/ir/exit_loop.h
+++ b/src/tint/ir/exit_loop.h
@@ -33,7 +33,7 @@
/// Constructor
/// @param loop the loop being exited
- /// @param args the branch arguments
+ /// @param args the target MultiInBlock arguments
explicit ExitLoop(ir::Loop* loop, utils::VectorRef<Value*> args = utils::Empty);
~ExitLoop() override;
diff --git a/src/tint/ir/exit_switch.h b/src/tint/ir/exit_switch.h
index 9567a59..6f32890 100644
--- a/src/tint/ir/exit_switch.h
+++ b/src/tint/ir/exit_switch.h
@@ -33,7 +33,7 @@
/// Constructor
/// @param sw the switch being exited
- /// @param args the branch arguments
+ /// @param args the target MultiInBlock arguments
explicit ExitSwitch(ir::Switch* sw, utils::VectorRef<Value*> args = utils::Empty);
~ExitSwitch() override;
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index c831d83..3ff8811 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -172,13 +172,13 @@
diagnostics_.add_error(tint::diag::System::IR, err, s);
}
- bool NeedBranch() { return current_block_ && !current_block_->HasBranchTarget(); }
+ bool NeedTerminator() { return current_block_ && !current_block_->HasTerminator(); }
- void SetBranch(Branch* br) {
+ void SetTerminator(Terminator* terminator) {
TINT_ASSERT(IR, current_block_);
- TINT_ASSERT(IR, !current_block_->HasBranchTarget());
+ TINT_ASSERT(IR, !current_block_->HasTerminator());
- current_block_->Append(br);
+ current_block_->Append(terminator);
current_block_ = nullptr;
}
@@ -211,7 +211,7 @@
// Folded away and doesn't appear in the IR.
},
[&](const ast::Variable* var) {
- // Setup the current flow node to be the root block for the module. The builder
+ // Setup the current block to be the root block for the module. The builder
// will handle creating it if it doesn't exist already.
TINT_SCOPED_ASSIGNMENT(current_block_, builder_.RootBlock());
EmitVariable(var);
@@ -420,10 +420,9 @@
TINT_SCOPED_ASSIGNMENT(current_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.
- if (NeedBranch()) {
- SetBranch(builder_.Return(current_function_));
+ // Add a terminator if one was not already created.
+ if (NeedTerminator()) {
+ SetTerminator(builder_.Return(current_function_));
}
TINT_ASSERT(IR, control_stack_.IsEmpty());
@@ -593,9 +592,9 @@
scopes_.Push();
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.
+ // Note, this doesn't need to emit a Block as the current block should be sufficient as the
+ // blocks all get flattened. Each flow control node will inject the basic blocks it
+ // requires.
EmitStatements(block->statements);
}
@@ -615,9 +614,9 @@
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->True());
EmitBlock(stmt->body);
- // If the true block did not branch, then emit an exit_if
- if (NeedBranch()) {
- SetBranch(builder_.ExitIf(if_inst));
+ // If the true block did not terminate, then emit an exit_if
+ if (NeedTerminator()) {
+ SetTerminator(builder_.ExitIf(if_inst));
}
}
@@ -625,9 +624,9 @@
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->False());
EmitStatement(stmt->else_statement);
- // If the false block did not branch, then emit an exit_if
- if (NeedBranch()) {
- SetBranch(builder_.ExitIf(if_inst));
+ // If the false block did not terminate, then emit an exit_if
+ if (NeedTerminator()) {
+ SetTerminator(builder_.ExitIf(if_inst));
}
}
}
@@ -650,8 +649,8 @@
EmitStatements(stmt->body->statements);
// The current block didn't `break`, `return` or `continue`, go to the continuing block.
- if (NeedBranch()) {
- SetBranch(builder_.Continue(loop_inst));
+ if (NeedTerminator()) {
+ SetTerminator(builder_.Continue(loop_inst));
}
}
@@ -660,9 +659,9 @@
if (stmt->continuing) {
EmitBlock(stmt->continuing);
}
- // Branch back to the start block if the continue target didn't branch out already
- if (NeedBranch()) {
- SetBranch(builder_.NextIteration(loop_inst));
+ // Branch back to the start block if the continue target didn't terminate already
+ if (NeedTerminator()) {
+ SetTerminator(builder_.NextIteration(loop_inst));
}
}
}
@@ -676,7 +675,7 @@
// Continue is always empty, just go back to the start
{
TINT_SCOPED_ASSIGNMENT(current_block_, loop_inst->Continuing());
- SetBranch(builder_.NextIteration(loop_inst));
+ SetTerminator(builder_.NextIteration(loop_inst));
}
{
@@ -694,16 +693,16 @@
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->True());
- SetBranch(builder_.ExitIf(if_inst));
+ SetTerminator(builder_.ExitIf(if_inst));
}
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->False());
- SetBranch(builder_.ExitLoop(loop_inst));
+ SetTerminator(builder_.ExitLoop(loop_inst));
}
- if (NeedBranch()) {
- SetBranch(builder_.Continue(loop_inst));
+ if (NeedTerminator()) {
+ SetTerminator(builder_.Continue(loop_inst));
}
}
}
@@ -724,8 +723,8 @@
// Emit the for initializer before branching to the loop body
EmitStatement(stmt->initializer);
- if (NeedBranch()) {
- SetBranch(builder_.NextIteration(loop_inst));
+ if (NeedTerminator()) {
+ SetTerminator(builder_.NextIteration(loop_inst));
}
}
@@ -744,24 +743,24 @@
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->True());
- SetBranch(builder_.ExitIf(if_inst));
+ SetTerminator(builder_.ExitIf(if_inst));
}
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->False());
- SetBranch(builder_.ExitLoop(loop_inst));
+ SetTerminator(builder_.ExitLoop(loop_inst));
}
}
EmitBlock(stmt->body);
- if (NeedBranch()) {
- SetBranch(builder_.Continue(loop_inst));
+ if (NeedTerminator()) {
+ SetTerminator(builder_.Continue(loop_inst));
}
if (stmt->continuing) {
TINT_SCOPED_ASSIGNMENT(current_block_, loop_inst->Continuing());
EmitStatement(stmt->continuing);
- SetBranch(builder_.NextIteration(loop_inst));
+ SetTerminator(builder_.NextIteration(loop_inst));
}
}
@@ -790,8 +789,8 @@
TINT_SCOPED_ASSIGNMENT(current_block_, builder_.Case(switch_inst, selectors));
EmitBlock(c->Body()->Declaration());
- if (NeedBranch()) {
- SetBranch(builder_.ExitSwitch(switch_inst));
+ if (NeedTerminator()) {
+ SetTerminator(builder_.ExitSwitch(switch_inst));
}
}
}
@@ -806,9 +805,9 @@
ret_value = ret.Get();
}
if (ret_value) {
- SetBranch(builder_.Return(current_function_, ret_value));
+ SetTerminator(builder_.Return(current_function_, ret_value));
} else {
- SetBranch(builder_.Return(current_function_));
+ SetTerminator(builder_.Return(current_function_));
}
}
@@ -817,9 +816,9 @@
TINT_ASSERT(IR, current_control);
if (auto* c = current_control->As<Loop>()) {
- SetBranch(builder_.ExitLoop(c));
+ SetTerminator(builder_.ExitLoop(c));
} else if (auto* s = current_control->As<Switch>()) {
- SetBranch(builder_.ExitSwitch(s));
+ SetTerminator(builder_.ExitSwitch(s));
} else {
TINT_UNREACHABLE(IR, diagnostics_);
}
@@ -830,7 +829,7 @@
TINT_ASSERT(IR, current_control);
if (auto* c = current_control->As<Loop>()) {
- SetBranch(builder_.Continue(c));
+ SetTerminator(builder_.Continue(c));
} else {
TINT_UNREACHABLE(IR, diagnostics_);
}
@@ -853,7 +852,7 @@
if (!cond) {
return;
}
- SetBranch(builder_.BreakIf(cond.Get(), current_control->As<ir::Loop>()));
+ SetTerminator(builder_.BreakIf(cond.Get(), current_control->As<ir::Loop>()));
}
struct AccessorInfo {
@@ -1174,16 +1173,16 @@
// if (lhs) {
// res = rhs;
// } else {
- // res = lhs;
+ // res = false;
// }
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->True());
auto rhs = EmitExpression(expr->rhs);
- SetBranch(builder_.ExitIf(if_inst, utils::Vector{rhs.Get()}));
+ SetTerminator(builder_.ExitIf(if_inst, rhs.Get()));
}
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->False());
- SetBranch(builder_.ExitIf(if_inst, utils::Vector{lhs.Get()}));
+ SetTerminator(builder_.ExitIf(if_inst, builder_.Constant(false)));
}
} else {
// res = lhs || rhs;
@@ -1191,18 +1190,18 @@
// transform into:
//
// if (lhs) {
- // res = lhs;
+ // res = true;
// } else {
// res = rhs;
// }
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->True());
- SetBranch(builder_.ExitIf(if_inst, utils::Vector{lhs.Get()}));
+ SetTerminator(builder_.ExitIf(if_inst, builder_.Constant(true)));
}
{
TINT_SCOPED_ASSIGNMENT(current_block_, if_inst->False());
auto rhs = EmitExpression(expr->rhs);
- SetBranch(builder_.ExitIf(if_inst, utils::Vector{rhs.Get()}));
+ SetTerminator(builder_.ExitIf(if_inst, rhs.Get()));
}
}
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 86570e4..9763942 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -56,8 +56,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -80,8 +79,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -127,8 +125,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, i32, read_write> = var
}
@@ -151,8 +148,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -198,8 +194,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -245,8 +240,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -292,8 +286,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -339,8 +332,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, bool, read_write> = var
}
@@ -386,8 +378,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, bool, read_write> = var
}
@@ -433,8 +424,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -466,23 +456,19 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
- %logical_and:bool = if %3 [t: %b3, f: %b4]
- # True block
- %b3 = block {
- exit_if false
+ %logical_and:bool = if %3 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ exit_if false # if_1
}
-
- # False block
- %b4 = block {
- exit_if %3
+ %b4 = block { # false
+ exit_if false # if_1
}
-
- if %logical_and [t: %b5]
- # True block
- %b5 = block {
- exit_if
+ }
+ if %logical_and [t: %b5] { # if_2
+ %b5 = block { # true
+ exit_if # if_2
}
-
+ }
ret
}
}
@@ -506,23 +492,19 @@
%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b2 {
%b2 = block {
%3:bool = call %my_func
- %logical_or:bool = if %3 [t: %b3, f: %b4]
- # True block
- %b3 = block {
- exit_if %3
+ %logical_or:bool = if %3 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
+ exit_if true # if_1
}
-
- # False block
- %b4 = block {
- exit_if true
+ %b4 = block { # false
+ exit_if true # if_1
}
-
- if %logical_or [t: %b5]
- # True block
- %b5 = block {
- exit_if
+ }
+ if %logical_or [t: %b5] { # if_2
+ %b5 = block { # true
+ exit_if # if_2
}
-
+ }
ret
}
}
@@ -698,8 +680,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -745,8 +726,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v1:ptr<private, u32, read_write> = var
}
@@ -780,22 +760,19 @@
%b2 = block {
%3:f32 = call %my_func
%4:bool = lt %3, 2.0f
- %tint_symbol:bool = if %4 [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ %tint_symbol:bool = if %4 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
%6:f32 = call %my_func
%7:f32 = call %my_func
%8:f32 = mul 2.29999995231628417969f, %7
%9:f32 = div %6, %8
%10:bool = gt 2.5f, %9
- exit_if %10
+ exit_if %10 # if_1
}
-
- # False block
- %b4 = block {
- exit_if %4
+ %b4 = block { # false
+ exit_if false # if_1
}
-
+ }
ret
}
}
diff --git a/src/tint/ir/from_program_builtin_test.cc b/src/tint/ir/from_program_builtin_test.cc
index 262c8bd..296ec79 100644
--- a/src/tint/ir/from_program_builtin_test.cc
+++ b/src/tint/ir/from_program_builtin_test.cc
@@ -33,8 +33,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%i:ptr<private, f32, read_write> = var, 1.0f
}
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 1148f56..3807e85 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -99,8 +99,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%i:ptr<private, i32, read_write> = var, 1i
}
@@ -121,8 +120,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%i:ptr<private, vec3<f32>, read_write> = var, vec3<f32>(0.0f)
}
@@ -137,8 +135,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%i:ptr<private, f32, read_write> = var, 1.0f
}
diff --git a/src/tint/ir/from_program_store_test.cc b/src/tint/ir/from_program_store_test.cc
index c58cd47..31e9792 100644
--- a/src/tint/ir/from_program_store_test.cc
+++ b/src/tint/ir/from_program_store_test.cc
@@ -34,8 +34,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%a:ptr<private, u32, read_write> = var
}
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index 1dd8f24..81738d6 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -26,9 +26,9 @@
namespace tint::ir {
namespace {
-/// Looks for the flow node with the given type T.
-/// If no flow node is found, then nullptr is returned.
-/// If multiple flow nodes are found with the type T, then an error is raised and the first is
+/// Looks for the instruction with the given type T.
+/// If no instruction is found, then nullptr is returned.
+/// If multiple instructions are found with the type T, then an error is raised and the first is
/// returned.
template <typename T>
T* FindSingleInstruction(Module& mod) {
@@ -141,17 +141,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- if true [t: %b2, f: %b3]
- # True block
- %b2 = block {
- exit_if
+ if true [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ exit_if # if_1
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
ret
}
}
@@ -172,12 +169,11 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- if true [t: %b2]
- # True block
- %b2 = block {
+ if true [t: %b2] { # if_1
+ %b2 = block { # true
ret
}
-
+ }
ret
}
}
@@ -198,17 +194,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- if true [t: %b2, f: %b3]
- # True block
- %b2 = block {
- exit_if
+ if true [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ exit_if # if_1
}
-
- # False block
- %b3 = block {
+ %b3 = block { # false
ret
}
-
+ }
ret
}
}
@@ -229,17 +222,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- if true [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ if true [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret
}
-
- # False block
- %b3 = block {
+ %b3 = block { # false
ret
}
-
+ }
ret
}
}
@@ -259,23 +249,19 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- if true [t: %b2]
- # True block
- %b2 = block {
- loop [b: %b3, c: %b4]
- # Body block
- %b3 = block {
- exit_loop
+ if true [t: %b2] { # if_1
+ %b2 = block { # true
+ loop [b: %b3, c: %b4] { # loop_1
+ %b3 = block { # body
+ exit_loop # loop_1
}
-
- # Continuing block
- %b4 = block {
+ %b4 = block { # continuing
next_iteration %b3
}
-
- exit_if
+ }
+ exit_if # if_1
}
-
+ }
ret
}
}
@@ -300,17 +286,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
- exit_loop
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
+ exit_loop # loop_1
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -336,23 +319,19 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
- if true [t: %b4]
- # True block
- %b4 = block {
- exit_loop
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
+ if true [t: %b4] { # if_1
+ %b4 = block { # true
+ exit_loop # loop_1
}
-
+ }
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -378,17 +357,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
break_if true %b2
}
-
+ }
ret
}
}
@@ -408,17 +384,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
break_if true %b2
}
-
+ }
ret
}
}
@@ -444,23 +417,19 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
- if true [t: %b4]
- # True block
- %b4 = block {
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
+ if true [t: %b4] { # if_1
+ %b4 = block { # true
ret
}
-
+ }
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -485,17 +454,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
ret
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -529,23 +495,19 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
ret
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
break_if true %b2
}
-
- if true [t: %b4]
- # True block
- %b4 = block {
+ }
+ if true [t: %b4] { # if_1
+ %b4 = block { # true
ret
}
-
+ }
ret
}
}
@@ -571,28 +533,22 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
- if true [t: %b4, f: %b5]
- # True block
- %b4 = block {
- exit_loop
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
+ if true [t: %b4, f: %b5] { # if_1
+ %b4 = block { # true
+ exit_loop # loop_1
}
-
- # False block
- %b5 = block {
- exit_loop
+ %b5 = block { # false
+ exit_loop # loop_1
}
-
+ }
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -619,68 +575,53 @@
EXPECT_EQ(Disassemble(m.Get()),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
- loop [b: %b4, c: %b5]
- # Body block
- %b4 = block {
- if true [t: %b6]
- # True block
- %b6 = block {
- exit_loop
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
+ loop [b: %b4, c: %b5] { # loop_2
+ %b4 = block { # body
+ if true [t: %b6] { # if_1
+ %b6 = block { # true
+ exit_loop # loop_2
}
-
- if true [t: %b7]
- # True block
- %b7 = block {
+ }
+ if true [t: %b7] { # if_2
+ %b7 = block { # true
continue %b5
}
-
+ }
continue %b5
}
-
- # Continuing block
- %b5 = block {
- loop [b: %b8, c: %b9]
- # Body block
- %b8 = block {
- exit_loop
+ %b5 = block { # continuing
+ loop [b: %b8, c: %b9] { # loop_3
+ %b8 = block { # body
+ exit_loop # loop_3
}
-
- # Continuing block
- %b9 = block {
+ %b9 = block { # continuing
next_iteration %b8
}
-
- loop [b: %b10, c: %b11]
- # Body block
- %b10 = block {
+ }
+ loop [b: %b10, c: %b11] { # loop_4
+ %b10 = block { # body
continue %b11
}
-
- # Continuing block
- %b11 = block {
+ %b11 = block { # continuing
break_if true %b10
}
-
+ }
next_iteration %b4
}
-
- if true [t: %b12]
- # True block
- %b12 = block {
- exit_loop
+ }
+ if true [t: %b12] { # if_3
+ %b12 = block { # true
+ exit_loop # loop_1
}
-
+ }
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -705,28 +646,22 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
- if false [t: %b4, f: %b5]
- # True block
- %b4 = block {
- exit_if
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
+ if false [t: %b4, f: %b5] { # if_1
+ %b4 = block { # true
+ exit_if # if_1
}
-
- # False block
- %b5 = block {
- exit_loop
+ %b5 = block { # false
+ exit_loop # loop_1
}
-
+ }
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -751,28 +686,22 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2, c: %b3]
- # Body block
- %b2 = block {
- if true [t: %b4, f: %b5]
- # True block
- %b4 = block {
- exit_if
+ loop [b: %b2, c: %b3] { # loop_1
+ %b2 = block { # body
+ if true [t: %b4, f: %b5] { # if_1
+ %b4 = block { # true
+ exit_if # if_1
}
-
- # False block
- %b5 = block {
- exit_loop
+ %b5 = block { # false
+ exit_loop # loop_1
}
-
+ }
continue %b3
}
-
- # Continuing block
- %b3 = block {
+ %b3 = block { # continuing
next_iteration %b2
}
-
+ }
ret
}
}
@@ -828,18 +757,15 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [i: %b2, b: %b3]
- # Initializer block
- %b2 = block {
+ loop [i: %b2, b: %b3] { # loop_1
+ %b2 = block { # initializer
%i:ptr<function, i32, read_write> = var
next_iteration %b3
}
-
- # Body block
- %b3 = block {
- exit_loop
+ %b3 = block { # body
+ exit_loop # loop_1
}
-
+ }
ret
}
}
@@ -864,12 +790,11 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- loop [b: %b2]
- # Body block
- %b2 = block {
- exit_loop
+ loop [b: %b2] { # loop_1
+ %b2 = block { # body
+ exit_loop # loop_1
}
-
+ }
ret
}
}
@@ -887,11 +812,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleInstruction<ir::Switch>(m);
+ auto* swtch = FindSingleInstruction<ir::Switch>(m);
ASSERT_EQ(1u, m.functions.Length());
- auto cases = flow->Cases();
+ auto cases = swtch->Cases();
ASSERT_EQ(3u, cases.Length());
ASSERT_EQ(1u, cases[0].selectors.Length());
@@ -910,22 +835,17 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4)]
- # Case block
- %b2 = block {
- exit_switch
+ switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4)] { # switch_1
+ %b2 = block { # case
+ exit_switch # switch_1
}
-
- # Case block
- %b3 = block {
- exit_switch
+ %b3 = block { # case
+ exit_switch # switch_1
}
-
- # Case block
- %b4 = block {
- exit_switch
+ %b4 = block { # case
+ exit_switch # switch_1
}
-
+ }
ret
}
}
@@ -944,11 +864,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleInstruction<ir::Switch>(m);
+ auto* swtch = FindSingleInstruction<ir::Switch>(m);
ASSERT_EQ(1u, m.functions.Length());
- auto cases = flow->Cases();
+ auto cases = swtch->Cases();
ASSERT_EQ(1u, cases.Length());
ASSERT_EQ(3u, cases[0].selectors.Length());
ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
@@ -964,12 +884,11 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- switch 1i [c: (0i 1i default, %b2)]
- # Case block
- %b2 = block {
- exit_switch
+ switch 1i [c: (0i 1i default, %b2)] { # switch_1
+ %b2 = block { # case
+ exit_switch # switch_1
}
-
+ }
ret
}
}
@@ -984,11 +903,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleInstruction<ir::Switch>(m);
+ auto* swtch = FindSingleInstruction<ir::Switch>(m);
ASSERT_EQ(1u, m.functions.Length());
- auto cases = flow->Cases();
+ auto cases = swtch->Cases();
ASSERT_EQ(1u, cases.Length());
ASSERT_EQ(1u, cases[0].selectors.Length());
EXPECT_TRUE(cases[0].selectors[0].IsDefault());
@@ -996,12 +915,11 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- switch 1i [c: (default, %b2)]
- # Case block
- %b2 = block {
- exit_switch
+ switch 1i [c: (default, %b2)] { # switch_1
+ %b2 = block { # case
+ exit_switch # switch_1
}
-
+ }
ret
}
}
@@ -1018,11 +936,11 @@
ASSERT_TRUE(res) << (!res ? res.Failure() : "");
auto m = res.Move();
- auto* flow = FindSingleInstruction<ir::Switch>(m);
+ auto* swtch = FindSingleInstruction<ir::Switch>(m);
ASSERT_EQ(1u, m.functions.Length());
- auto cases = flow->Cases();
+ auto cases = swtch->Cases();
ASSERT_EQ(2u, cases.Length());
ASSERT_EQ(1u, cases[0].selectors.Length());
ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
@@ -1037,17 +955,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- switch 1i [c: (0i, %b2), c: (default, %b3)]
- # Case block
- %b2 = block {
- exit_switch
+ switch 1i [c: (0i, %b2), c: (default, %b3)] { # switch_1
+ %b2 = block { # case
+ exit_switch # switch_1
}
-
- # Case block
- %b3 = block {
- exit_switch
+ %b3 = block { # case
+ exit_switch # switch_1
}
-
+ }
ret
}
}
@@ -1066,11 +981,11 @@
auto m = res.Move();
- auto* flow = FindSingleInstruction<ir::Switch>(m);
+ auto* swtch = FindSingleInstruction<ir::Switch>(m);
ASSERT_EQ(1u, m.functions.Length());
- auto cases = flow->Cases();
+ auto cases = swtch->Cases();
ASSERT_EQ(2u, cases.Length());
ASSERT_EQ(1u, cases[0].selectors.Length());
ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
@@ -1083,17 +998,14 @@
EXPECT_EQ(Disassemble(m),
R"(%test_function = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
- switch 1i [c: (0i, %b2), c: (default, %b3)]
- # Case block
- %b2 = block {
+ switch 1i [c: (0i, %b2), c: (default, %b3)] { # switch_1
+ %b2 = block { # case
ret
}
-
- # Case block
- %b3 = block {
+ %b3 = block { # case
ret
}
-
+ }
ret
}
}
diff --git a/src/tint/ir/from_program_unary_test.cc b/src/tint/ir/from_program_unary_test.cc
index 2a66d1c..976863e 100644
--- a/src/tint/ir/from_program_unary_test.cc
+++ b/src/tint/ir/from_program_unary_test.cc
@@ -103,8 +103,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v2:ptr<private, i32, read_write> = var
}
@@ -127,8 +126,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%v3:ptr<private, i32, read_write> = var
}
diff --git a/src/tint/ir/from_program_var_test.cc b/src/tint/ir/from_program_var_test.cc
index aedb0cb..ce75297 100644
--- a/src/tint/ir/from_program_var_test.cc
+++ b/src/tint/ir/from_program_var_test.cc
@@ -31,8 +31,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%a:ptr<private, u32, read_write> = var
}
@@ -46,8 +45,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%a:ptr<private, u32, read_write> = var, 2u
}
@@ -61,8 +59,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
-%b1 = block {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%b1 = block { # root
%a:ptr<storage, u32, read> = var @binding_point(2, 3)
}
diff --git a/src/tint/ir/if.h b/src/tint/ir/if.h
index 6b87463..216b54b 100644
--- a/src/tint/ir/if.h
+++ b/src/tint/ir/if.h
@@ -58,10 +58,10 @@
/// @returns the if condition
Value* Condition() { return operands_[kConditionOperandOffset]; }
- /// @returns the true branch block
+ /// @returns the true block
ir::Block* True() { return true_; }
- /// @returns the false branch block
+ /// @returns the false block
ir::Block* False() { return false_; }
private:
diff --git a/src/tint/ir/loop.cc b/src/tint/ir/loop.cc
index 4b931b4..a09699d 100644
--- a/src/tint/ir/loop.cc
+++ b/src/tint/ir/loop.cc
@@ -54,7 +54,7 @@
}
bool Loop::HasInitializer() {
- return initializer_->HasBranchTarget();
+ return initializer_->HasTerminator();
}
} // namespace tint::ir
diff --git a/src/tint/ir/loop.h b/src/tint/ir/loop.h
index f5e2ea4..1a00ac3 100644
--- a/src/tint/ir/loop.h
+++ b/src/tint/ir/loop.h
@@ -45,7 +45,7 @@
/// ┃ ▼ ┃
/// ┃ ┌─────────────────┐ ┃ BreakIf(false)
/// ExitLoop ┃ │ Continuing │━━┛
-/// │ (optional) │
+/// ┃ │ (optional) │
/// ┃ └─────────────────┘
/// ┃ ┃
/// ┃ ┃ BreakIf(true)
diff --git a/src/tint/ir/multi_in_block.cc b/src/tint/ir/multi_in_block.cc
index 94e7f59..53286ba 100644
--- a/src/tint/ir/multi_in_block.cc
+++ b/src/tint/ir/multi_in_block.cc
@@ -32,7 +32,7 @@
params_ = std::move(params);
}
-void MultiInBlock::AddInboundSiblingBranch(ir::Branch* node) {
+void MultiInBlock::AddInboundSiblingBranch(ir::Terminator* node) {
TINT_ASSERT(IR, node != nullptr);
if (node) {
diff --git a/src/tint/ir/multi_in_block.h b/src/tint/ir/multi_in_block.h
index e0d807f..c70bb9c 100644
--- a/src/tint/ir/multi_in_block.h
+++ b/src/tint/ir/multi_in_block.h
@@ -47,17 +47,17 @@
const utils::Vector<BlockParam*, 2>& Params() { return params_; }
/// @returns branches made to this block by sibling blocks
- const utils::VectorRef<ir::Branch*> InboundSiblingBranches() {
+ const utils::VectorRef<ir::Terminator*> InboundSiblingBranches() {
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);
+ void AddInboundSiblingBranch(ir::Terminator* branch);
private:
utils::Vector<BlockParam*, 2> params_;
- utils::Vector<ir::Branch*, 2> inbound_sibling_branches_;
+ utils::Vector<ir::Terminator*, 2> inbound_sibling_branches_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/next_iteration.h b/src/tint/ir/next_iteration.h
index 6b01634..4783d53 100644
--- a/src/tint/ir/next_iteration.h
+++ b/src/tint/ir/next_iteration.h
@@ -15,7 +15,7 @@
#ifndef SRC_TINT_IR_NEXT_ITERATION_H_
#define SRC_TINT_IR_NEXT_ITERATION_H_
-#include "src/tint/ir/branch.h"
+#include "src/tint/ir/terminator.h"
#include "src/tint/utils/castable.h"
// Forward declarations
@@ -26,14 +26,14 @@
namespace tint::ir {
/// A next iteration instruction.
-class NextIteration : public utils::Castable<NextIteration, Branch> {
+class NextIteration : public utils::Castable<NextIteration, Terminator> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
/// Constructor
/// @param loop the loop being iterated
- /// @param args the branch arguments
+ /// @param args the arguments for the MultiInBlock
explicit NextIteration(ir::Loop* loop, utils::VectorRef<Value*> args = utils::Empty);
~NextIteration() override;
diff --git a/src/tint/ir/operand_instruction.h b/src/tint/ir/operand_instruction.h
index 82957c0..040bfc9 100644
--- a/src/tint/ir/operand_instruction.h
+++ b/src/tint/ir/operand_instruction.h
@@ -57,13 +57,18 @@
ClearOperands();
operands_ = std::move(operands);
for (size_t i = 0; i < operands_.Length(); i++) {
- operands_[i]->AddUsage({this, static_cast<uint32_t>(i)});
+ if (operands_[i]) {
+ operands_[i]->AddUsage({this, static_cast<uint32_t>(i)});
+ }
}
}
/// Removes all operands from the instruction
void ClearOperands() {
for (uint32_t i = 0; i < operands_.Length(); i++) {
+ if (!operands_[i]) {
+ continue;
+ }
operands_[i]->RemoveUsage({this, i});
}
operands_.Clear();
diff --git a/src/tint/ir/operand_instruction_test.cc b/src/tint/ir/operand_instruction_test.cc
index 9d05a07..03c0435 100644
--- a/src/tint/ir/operand_instruction_test.cc
+++ b/src/tint/ir/operand_instruction_test.cc
@@ -18,7 +18,8 @@
namespace tint::ir {
namespace {
-using namespace tint::number_suffixes; // NOLINT
+using namespace tint::number_suffixes; // NOLINT
+using namespace tint::builtin::fluent_types; // NOLINT
using IR_OperandInstructionTest = IRTestHelper;
@@ -41,5 +42,24 @@
EXPECT_FALSE(inst->Result()->Alive());
}
+TEST_F(IR_OperandInstructionTest, ClearOperands_WithNullOperand) {
+ auto* block = b.Block();
+ // The var initializer is a nullptr
+ auto* inst = b.Var(ty.ptr<private_, f32>());
+ block->Append(inst);
+
+ inst->Destroy();
+ EXPECT_EQ(inst->Block(), nullptr);
+ EXPECT_FALSE(inst->Result()->Alive());
+}
+
+TEST_F(IR_OperandInstructionTest, SetOperands_WithNullOperand) {
+ auto* inst = b.Var(ty.ptr<private_, f32>());
+ utils::Vector<Value*, 1> ops;
+ ops.Push(nullptr);
+
+ inst->SetOperands(ops);
+}
+
} // namespace
} // namespace tint::ir
diff --git a/src/tint/ir/return.h b/src/tint/ir/return.h
index 403a215..14f1863 100644
--- a/src/tint/ir/return.h
+++ b/src/tint/ir/return.h
@@ -15,7 +15,7 @@
#ifndef SRC_TINT_IR_RETURN_H_
#define SRC_TINT_IR_RETURN_H_
-#include "src/tint/ir/branch.h"
+#include "src/tint/ir/terminator.h"
#include "src/tint/utils/castable.h"
// Forward declarations
@@ -26,7 +26,7 @@
namespace tint::ir {
/// A return instruction.
-class Return : public utils::Castable<Return, Branch> {
+class Return : public utils::Castable<Return, Terminator> {
public:
/// The offset in Operands() for the function being returned
static constexpr size_t kFunctionOperandOffset = 0;
@@ -57,7 +57,7 @@
/// @param val the new return value
void SetValue(ir::Value* val) { SetOperand(kArgOperandOffset, val); }
- /// @returns the branch arguments
+ /// @returns the return arguments
utils::Slice<ir::Value* const> Args() override {
return operands_.Slice().Offset(kArgOperandOffset);
}
diff --git a/src/tint/ir/branch.cc b/src/tint/ir/terminator.cc
similarity index 84%
rename from src/tint/ir/branch.cc
rename to src/tint/ir/terminator.cc
index a7ea495..ce76977 100644
--- a/src/tint/ir/branch.cc
+++ b/src/tint/ir/terminator.cc
@@ -12,14 +12,14 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/ir/branch.h"
+#include "src/tint/ir/terminator.h"
#include <utility>
-TINT_INSTANTIATE_TYPEINFO(tint::ir::Branch);
+TINT_INSTANTIATE_TYPEINFO(tint::ir::Terminator);
namespace tint::ir {
-Branch::~Branch() = default;
+Terminator::~Terminator() = default;
} // namespace tint::ir
diff --git a/src/tint/ir/branch.h b/src/tint/ir/terminator.h
similarity index 74%
rename from src/tint/ir/branch.h
rename to src/tint/ir/terminator.h
index 1069529..5698100 100644
--- a/src/tint/ir/branch.h
+++ b/src/tint/ir/terminator.h
@@ -12,8 +12,8 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#ifndef SRC_TINT_IR_BRANCH_H_
-#define SRC_TINT_IR_BRANCH_H_
+#ifndef SRC_TINT_IR_TERMINATOR_H_
+#define SRC_TINT_IR_TERMINATOR_H_
#include "src/tint/ir/operand_instruction.h"
#include "src/tint/ir/value.h"
@@ -26,15 +26,15 @@
namespace tint::ir {
-/// A branch instruction.
-class Branch : public utils::Castable<Branch, OperandInstruction<1, 0>> {
+/// The base class of all instructions that terminate a block.
+class Terminator : public utils::Castable<Terminator, OperandInstruction<1, 0>> {
public:
- ~Branch() override;
+ ~Terminator() override;
- /// @returns the branch arguments
+ /// @returns the terminator arguments
virtual utils::Slice<Value* const> Args() { return operands_.Slice(); }
};
} // namespace tint::ir
-#endif // SRC_TINT_IR_BRANCH_H_
+#endif // SRC_TINT_IR_TERMINATOR_H_
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index a286702..495f608 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -15,8 +15,10 @@
#include "src/tint/ir/to_program.h"
#include <string>
+#include <tuple>
#include <utility>
+#include "src/tint/ir/access.h"
#include "src/tint/ir/binary.h"
#include "src/tint/ir/block.h"
#include "src/tint/ir/call.h"
@@ -46,6 +48,7 @@
#include "src/tint/type/texture.h"
#include "src/tint/utils/hashmap.h"
#include "src/tint/utils/predicates.h"
+#include "src/tint/utils/reverse.h"
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/utils/transform.h"
#include "src/tint/utils/vector.h"
@@ -65,6 +68,10 @@
namespace {
+/// Empty struct used as a sentinel value to indicate that an ast::Value has been consumed by its
+/// single place of usage. Attempting to use this value a second time should result in an ICE.
+struct ConsumedValue {};
+
class State {
public:
explicit State(Module& m) : mod(m) {}
@@ -85,15 +92,25 @@
/// The target ProgramBuilder
ProgramBuilder b;
- /// A hashmap of value to symbol used in the emitted AST
- utils::Hashmap<Value*, Symbol, 32> value_names_;
+ using ValueBinding = std::variant<Symbol, const ast::Expression*, ConsumedValue>;
+
+ /// A hashmap of value to one of:
+ /// * Symbol - Name of 'let' (non-inlinable value), 'var' or parameter.
+ /// * ast::Expression* - single use, inlined expression.
+ /// * ConsumedValue - a special value used to indicate that the value has already been
+ /// consumed.
+ utils::Hashmap<Value*, ValueBinding, 32> bindings_;
/// The nesting depth of the currently generated AST
- /// 0 is module scope
- /// 1 is root-level function scope
+ /// 0 is module scope
+ /// 1 is root-level function scope
/// 2+ is within control flow
uint32_t nesting_depth_ = 0;
+ using StatementList = utils::Vector<const ast::Statement*,
+ decltype(ast::BlockStatement::statements)::static_length>;
+ StatementList* statements_ = nullptr;
+
/// The current switch case block
ir::Block* current_switch_case_ = nullptr;
@@ -103,12 +120,12 @@
// TODO(crbug.com/tint/1915): Properly implement this when we've fleshed out Function
static constexpr size_t N = decltype(ast::Function::params)::static_length;
auto params = utils::Transform<N>(fn->Params(), [&](FunctionParam* param) {
- auto name = AssignNameTo(param);
+ auto name = BindName(param);
auto ty = Type(param->Type());
return b.Param(name, ty);
});
- auto name = AssignNameTo(fn);
+ auto name = BindName(fn);
auto ret_ty = Type(fn->ReturnType());
auto* body = Block(fn->StartTarget());
utils::Vector<const ast::Attribute*, 1> attrs{};
@@ -118,175 +135,129 @@
}
const ast::BlockStatement* Block(ir::Block* block) {
- static constexpr size_t N = decltype(ast::BlockStatement::statements)::static_length;
- utils::Vector<const ast::Statement*, N> stmts;
-
// TODO(crbug.com/tint/1902): Handle block arguments.
+ return b.Block(Statements(block));
+ }
- for (auto* inst : *block) {
- if (auto stmt = Stmt(inst)) {
- stmts.Push(stmt);
+ StatementList Statements(ir::Block* block) {
+ StatementList stmts;
+ if (block) {
+ TINT_SCOPED_ASSIGNMENT(statements_, &stmts);
+ for (auto* inst : *block) {
+ Instruction(inst);
}
}
-
- return b.Block(std::move(stmts));
+ return stmts;
}
- ////////////////////////////////////////////////////////////////////////////////////////////////
- // Statements
- //
- // Statement methods may return nullptr, in the case of instructions that do not map to an AST
- // statement, or in the case of an error. These should simply be ignored.
- ////////////////////////////////////////////////////////////////////////////////////////////////
+ void Append(const ast::Statement* inst) { statements_->Push(inst); }
- /// @param inst the ir::Instruction
- /// @return an ast::Statement from @p inst, or nullptr if there was an error
- const ast::Statement* Stmt(ir::Instruction* inst) {
- return tint::Switch(
- inst, //
- [&](ir::Call* i) { return CallStmt(i); }, //
- [&](ir::Var* i) { return Var(i); }, //
- [&](ir::Load*) { return nullptr; }, //
- [&](ir::Store* i) { return Store(i); }, //
- [&](ir::If* if_) { return If(if_); }, //
- [&](ir::Switch* switch_) { return Switch(switch_); }, //
- [&](ir::Return* ret) { return Return(ret); }, //
- [&](ir::ExitSwitch* e) { return ExitSwitch(e); }, //
- [&](ir::ExitIf*) { return nullptr; }, //
- [&](ir::Instruction* i) { return ValueStmt(i->Result()); }, //
- [&](Default) {
- UNHANDLED_CASE(inst);
- return nullptr;
- });
+ void Instruction(ir::Instruction* inst) {
+ tint::Switch(
+ inst, //
+ [&](ir::Binary* u) { Binary(u); }, //
+ [&](ir::Call* i) { Call(i); }, //
+ [&](ir::ExitIf*) {}, //
+ [&](ir::ExitSwitch* i) { ExitSwitch(i); }, //
+ [&](ir::If* i) { If(i); }, //
+ [&](ir::Load* l) { Load(l); }, //
+ [&](ir::Return* i) { Return(i); }, //
+ [&](ir::Store* i) { Store(i); }, //
+ [&](ir::Switch* i) { Switch(i); }, //
+ [&](ir::Unary* u) { Unary(u); }, //
+ [&](ir::Var* i) { Var(i); }, //
+ [&](Default) { UNHANDLED_CASE(inst); });
}
- /// @param i the ir::If
- /// @return an ast::IfStatement from @p i, or nullptr if there was an error
- const ast::IfStatement* If(ir::If* i) {
+ void If(ir::If* if_) {
SCOPED_NESTING();
- auto* cond = Expr(i->Condition());
- auto* t = Block(i->True());
- if (TINT_UNLIKELY(!t)) {
- return nullptr;
+
+ auto true_stmts = Statements(if_->True());
+ auto false_stmts = Statements(if_->False());
+ if (IsShortCircuit(if_, true_stmts, false_stmts)) {
+ return;
}
- if (auto* false_blk = i->False(); false_blk && !false_blk->IsEmpty()) {
- bool maybe_elseif = (false_blk->Length() == 1) ||
- (false_blk->Length() == 2 && false_blk->Back()->Is<ir::Branch>());
- if (maybe_elseif) {
- if (auto* else_if = false_blk->Front()->As<ir::If>()) {
- auto* f = If(else_if);
- if (!f) {
- return nullptr;
- }
- return b.If(cond, t, b.Else(f));
+ auto* cond = Expr(if_->Condition());
+ auto* true_block = b.Block(std::move(true_stmts));
+
+ switch (false_stmts.Length()) {
+ case 0:
+ Append(b.If(cond, true_block));
+ return;
+ case 1:
+ if (auto* else_if = false_stmts.Front()->As<ast::IfStatement>()) {
+ Append(b.If(cond, true_block, b.Else(else_if)));
+ return;
}
- }
-
- auto* f = Block(i->False());
- if (!f) {
- return nullptr;
- }
- return b.If(cond, t, b.Else(f));
+ break;
}
- return b.If(cond, t);
+ auto* false_block = b.Block(std::move(false_stmts));
+ Append(b.If(cond, true_block, b.Else(false_block)));
}
- /// @param s the ir::Switch
- /// @return an ast::SwitchStatement from @p s, or nullptr if there was an error
- const ast::SwitchStatement* Switch(ir::Switch* s) {
+ void Switch(ir::Switch* s) {
SCOPED_NESTING();
auto* cond = Expr(s->Condition());
- if (!cond) {
- return nullptr;
- }
- auto cases =
- utils::Transform(s->Cases(), //
- [&](ir::Switch::Case c) -> const tint::ast::CaseStatement* {
- SCOPED_NESTING();
+ auto cases = utils::Transform(
+ s->Cases(), //
+ [&](ir::Switch::Case c) -> const tint::ast::CaseStatement* {
+ SCOPED_NESTING();
- const ast::BlockStatement* body = nullptr;
- {
- TINT_SCOPED_ASSIGNMENT(current_switch_case_, c.Block());
- body = Block(c.Block());
- }
- if (!body) {
- return nullptr;
- }
+ const ast::BlockStatement* body = nullptr;
+ {
+ TINT_SCOPED_ASSIGNMENT(current_switch_case_, c.Block());
+ body = Block(c.Block());
+ }
- auto selectors = utils::Transform(
- c.selectors, //
- [&](ir::Switch::CaseSelector cs) -> const ast::CaseSelector* {
- if (cs.IsDefault()) {
- return b.DefaultCaseSelector();
- }
- auto* expr = Expr(cs.val);
- if (!expr) {
- return nullptr;
- }
- return b.CaseSelector(expr);
- });
- if (selectors.Any(utils::IsNull)) {
- return nullptr;
- }
+ auto selectors = utils::Transform(c.selectors, //
+ [&](ir::Switch::CaseSelector cs) {
+ return cs.IsDefault()
+ ? b.DefaultCaseSelector()
+ : b.CaseSelector(Expr(cs.val));
+ });
+ return b.Case(std::move(selectors), body);
+ });
- return b.Case(std::move(selectors), body);
- });
- if (cases.Any(utils::IsNull)) {
- return nullptr;
- }
-
- return b.Switch(cond, std::move(cases));
+ Append(b.Switch(cond, std::move(cases)));
}
- const ast::BreakStatement* ExitSwitch(const ir::ExitSwitch* e) {
- if (current_switch_case_ && current_switch_case_->Branch() == e) {
- return nullptr; // No need to emit
+ void ExitSwitch(const ir::ExitSwitch* e) {
+ if (current_switch_case_ && current_switch_case_->Terminator() == e) {
+ return; // No need to emit
}
- return b.Break();
+ Append(b.Break());
}
- /// @param ret the ir::Return
- /// @return an ast::ReturnStatement from @p ret, or nullptr if there was an error
- const ast::ReturnStatement* Return(ir::Return* ret) {
+ void Return(ir::Return* ret) {
if (ret->Args().IsEmpty()) {
// Return 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 (nesting_depth_ > 1) {
- return b.Return();
+ Append(b.Return());
}
- return nullptr;
+ return;
}
// Return has arguments - this is the return value.
if (ret->Args().Length() != 1) {
TINT_ICE(IR, b.Diagnostics())
<< "expected 1 value for return, got " << ret->Args().Length();
- return b.Return();
+ return;
}
- auto* val = Expr(ret->Args().Front());
- if (TINT_UNLIKELY(!val)) {
- return b.Return();
- }
-
- return b.Return(val);
+ Append(b.Return(Expr(ret->Args().Front())));
}
- /// @param call the ir::Call
- /// @return an ast::CallStatement from @p call, or nullptr if there was an error
- const ast::CallStatement* CallStmt(ir::Call* call) { return b.CallStmt(Call(call)); }
-
- /// @param var the ir::Var
- /// @return an ast::VariableDeclStatement from @p var
- const ast::VariableDeclStatement* Var(ir::Var* var) {
- Symbol name = AssignNameTo(var);
- auto* ptr = var->Result()->Type()->As<type::Pointer>();
+ void Var(ir::Var* var) {
+ auto* val = var->Result();
+ Symbol name = BindName(val);
+ auto* ptr = As<type::Pointer>(val->Type());
auto ty = Type(ptr->StoreType());
const ast::Expression* init = nullptr;
if (var->Initializer()) {
@@ -294,104 +265,163 @@
}
switch (ptr->AddressSpace()) {
case builtin::AddressSpace::kFunction:
- return b.Decl(b.Var(name, ty, init));
+ Append(b.Decl(b.Var(name, ty, init)));
+ return;
case builtin::AddressSpace::kStorage:
- return b.Decl(b.Var(name, ty, init, ptr->Access(), ptr->AddressSpace()));
+ Append(b.Decl(b.Var(name, ty, init, ptr->Access(), ptr->AddressSpace())));
+ return;
default:
- return b.Decl(b.Var(name, ty, init, ptr->AddressSpace()));
+ Append(b.Decl(b.Var(name, ty, init, ptr->AddressSpace())));
+ return;
}
}
- /// @param store the ir::Store
- /// @return an ast::AssignmentStatement from @p call
- const ast::AssignmentStatement* Store(ir::Store* store) {
- auto* expr = Expr(store->From());
- return b.Assign(AssignNameTo(store->To()), expr);
+ void Store(ir::Store* store) {
+ auto* dst = Expr(store->To());
+ auto* src = Expr(store->From());
+ Append(b.Assign(dst, src));
}
- /// @param val the ir::Value
- /// @return an ast::Statement from @p val, or nullptr if the value does not produce a statement.
- const ast::Statement* ValueStmt(ir::Value* val) {
- // As we're visiting this value's declaration it shouldn't already have a name reserved.
- TINT_ASSERT(IR, !value_names_.Contains(val));
-
- // Determine whether the value should be placed into a let, or inlined in its single place
- // of usage. Currently a value is inlined if it has a single usage and is unnamed.
- // TODO(crbug.com/tint/1902): This logic needs to check that the sequence of side -
- // effecting expressions is not changed by inlining the expression. This needs fixing.
- bool create_let = val->Usages().Count() > 1 || mod.NameOf(val).IsValid();
- if (create_let) {
- auto* init = Expr(val); // Must come before giving the value a name
- auto name = AssignNameTo(val);
- return b.Decl(b.Let(name, init));
- }
- return nullptr; // Value will be inlined at its place of usage.
- }
-
- ////////////////////////////////////////////////////////////////////////////////////////////////
- // Expressions
- //
- // The the case of an error:
- // * The expression generating methods must return a non-null ast expression pointer, which may
- // not be semantically legal, but is enough to populate the AST.
- // * A diagnostic error must be added to the ast::ProgramBuilder.
- // This prevents littering the ToProgram logic with expensive error checking code.
- ////////////////////////////////////////////////////////////////////////////////////////////////
-
- /// @param val the value
- /// @returns the ast::Expression from the values source instruction
- const ast::Expression* Expr(ir::Value* val) {
- if (auto name = value_names_.Get(val)) {
- return b.Expr(name.value());
- }
-
- return tint::Switch(
- val, //
- [&](ir::Constant* c) { return ConstExpr(c); }, //
- [&](ir::InstructionResult* r) { return Expr(r->Source()); }, //
- [&](Default) -> const ast::Expression* {
- UNHANDLED_CASE(val);
- return b.Expr("<error>");
- });
- }
-
- /// @param val the ir::Expression
- /// @return an ast::Expression from @p val.
- /// @note May be a semantically-invalid placeholder expression on error.
- const ast::Expression* Expr(ir::Instruction* val) {
- if (auto name = value_names_.Get(val->Result())) {
- return b.Expr(name.value());
- }
-
- return tint::Switch(
- val, //
- [&](ir::Load* l) { return LoadExpr(l); }, //
- [&](ir::Unary* u) { return UnaryExpr(u); }, //
- [&](ir::Binary* u) { return BinaryExpr(u); }, //
- [&](Default) {
- UNHANDLED_CASE(val);
- return b.Expr("<error>");
- });
- }
-
- /// @param call the ir::Call
- /// @return an ast::CallExpression from @p call.
- /// @note May be a semantically-invalid placeholder expression on error.
- const ast::CallExpression* Call(ir::Call* call) {
+ void Call(ir::Call* call) {
auto args = utils::Transform<2>(call->Args(), [&](ir::Value* arg) { return Expr(arg); });
- return tint::Switch(
+ tint::Switch(
call, //
- [&](ir::UserCall* c) { return b.Call(AssignNameTo(c->Func()), std::move(args)); },
- [&](Default) {
- UNHANDLED_CASE(call);
- return b.Call("<error>");
+ [&](ir::UserCall* c) {
+ auto* expr = b.Call(BindName(c->Func()), std::move(args));
+ if (!call->HasResults() || call->Result()->Usages().IsEmpty()) {
+ Append(b.CallStmt(expr));
+ return;
+ }
+ Bind(c->Result(), expr);
+ },
+ [&](Default) { UNHANDLED_CASE(call); });
+ }
+
+ void Load(ir::Load* l) { Bind(l->Result(), Expr(l->From())); }
+
+ void Unary(ir::Unary* u) {
+ const ast::Expression* expr = nullptr;
+ switch (u->Kind()) {
+ case ir::Unary::Kind::kComplement:
+ expr = b.Complement(Expr(u->Val()));
+ break;
+ case ir::Unary::Kind::kNegation:
+ expr = b.Negation(Expr(u->Val()));
+ break;
+ }
+ Bind(u->Result(), expr);
+ }
+
+ void Binary(ir::Binary* e) {
+ if (e->Kind() == ir::Binary::Kind::kEqual) {
+ auto* rhs = e->RHS()->As<ir::Constant>();
+ if (rhs && rhs->Type()->Is<type::Bool>() && rhs->Value()->ValueAs<bool>() == false) {
+ // expr == false
+ Bind(e->Result(), b.Not(Expr(e->LHS())));
+ return;
+ }
+ }
+ auto* lhs = Expr(e->LHS());
+ auto* rhs = Expr(e->RHS());
+ const ast::Expression* expr = nullptr;
+ switch (e->Kind()) {
+ case ir::Binary::Kind::kAdd:
+ expr = b.Add(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kSubtract:
+ expr = b.Sub(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kMultiply:
+ expr = b.Mul(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kDivide:
+ expr = b.Div(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kModulo:
+ expr = b.Mod(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kAnd:
+ expr = b.And(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kOr:
+ expr = b.Or(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kXor:
+ expr = b.Xor(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kEqual:
+ expr = b.Equal(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kNotEqual:
+ expr = b.NotEqual(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kLessThan:
+ expr = b.LessThan(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kGreaterThan:
+ expr = b.GreaterThan(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kLessThanEqual:
+ expr = b.LessThanEqual(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kGreaterThanEqual:
+ expr = b.GreaterThanEqual(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kShiftLeft:
+ expr = b.Shl(lhs, rhs);
+ break;
+ case ir::Binary::Kind::kShiftRight:
+ expr = b.Shr(lhs, rhs);
+ break;
+ }
+ Bind(e->Result(), expr);
+ }
+
+ TINT_BEGIN_DISABLE_WARNING(UNREACHABLE_CODE);
+
+ const ast::Expression* Expr(ir::Value* value) {
+ return tint::Switch(
+ value, //
+ [&](ir::Constant* c) { return Constant(c); }, //
+ [&](Default) -> const ast::Expression* {
+ auto lookup = bindings_.Find(value);
+ if (TINT_UNLIKELY(!lookup)) {
+ TINT_ICE(IR, b.Diagnostics())
+ << "Expr(" << (value ? value->TypeInfo().name : "null")
+ << ") value has no expression";
+ return b.Expr("<error>");
+ }
+ return std::visit(
+ [&](auto&& got) -> const ast::Expression* {
+ using T = std::decay_t<decltype(got)>;
+
+ if constexpr (std::is_same_v<T, Symbol>) {
+ return b.Expr(got); // var, let or parameter.
+ }
+
+ if constexpr (std::is_same_v<T, const ast::Expression*>) {
+ // Single use (inlined) expression.
+ // Mark the bindings_ map entry as consumed.
+ *lookup = ConsumedValue{};
+ return got;
+ }
+
+ if constexpr (std::is_same_v<T, ConsumedValue>) {
+ TINT_ICE(IR, b.Diagnostics()) << "Expr(" << value->TypeInfo().name
+ << ") called twice on the same value";
+ } else {
+ TINT_ICE(IR, b.Diagnostics())
+ << "Expr(" << value->TypeInfo().name << ") has unhandled value";
+ }
+ return b.Expr("<error>");
+ },
+ *lookup);
});
}
- /// @param c the ir::Constant
- /// @return an ast::Expression from @p c.
- /// @note May be a semantically-invalid placeholder expression on error.
- const ast::Expression* ConstExpr(ir::Constant* c) {
+ TINT_END_DISABLE_WARNING(UNREACHABLE_CODE);
+
+ const ast::Expression* Constant(ir::Constant* c) {
return tint::Switch(
c->Type(), //
[&](const type::I32*) { return b.Expr(c->Value()->ValueAs<i32>()); },
@@ -405,74 +435,6 @@
});
}
- /// @param l the ir::Load
- /// @return an ast::Expression from @p l.
- /// @note May be a semantically-invalid placeholder expression on error.
- const ast::Expression* LoadExpr(ir::Load* l) { return Expr(l->From()); }
-
- /// @param u the ir::Unary
- /// @return an ast::UnaryOpExpression from @p u.
- /// @note May be a semantically-invalid placeholder expression on error.
- const ast::Expression* UnaryExpr(ir::Unary* u) {
- switch (u->Kind()) {
- case ir::Unary::Kind::kComplement:
- return b.Complement(Expr(u->Val()));
- case ir::Unary::Kind::kNegation:
- return b.Negation(Expr(u->Val()));
- }
- return b.Expr("<error>");
- }
-
- /// @param e the ir::Binary
- /// @return an ast::BinaryOpExpression from @p e.
- /// @note May be a semantically-invalid placeholder expression on error.
- const ast::Expression* BinaryExpr(ir::Binary* e) {
- if (e->Kind() == ir::Binary::Kind::kEqual) {
- auto* rhs = e->RHS()->As<ir::Constant>();
- if (rhs && rhs->Type()->Is<type::Bool>() && rhs->Value()->ValueAs<bool>() == false) {
- // expr == false
- return b.Not(Expr(e->LHS()));
- }
- }
- auto* lhs = Expr(e->LHS());
- auto* rhs = Expr(e->RHS());
- switch (e->Kind()) {
- case ir::Binary::Kind::kAdd:
- return b.Add(lhs, rhs);
- case ir::Binary::Kind::kSubtract:
- return b.Sub(lhs, rhs);
- case ir::Binary::Kind::kMultiply:
- return b.Mul(lhs, rhs);
- case ir::Binary::Kind::kDivide:
- return b.Div(lhs, rhs);
- case ir::Binary::Kind::kModulo:
- return b.Mod(lhs, rhs);
- case ir::Binary::Kind::kAnd:
- return b.And(lhs, rhs);
- case ir::Binary::Kind::kOr:
- return b.Or(lhs, rhs);
- case ir::Binary::Kind::kXor:
- return b.Xor(lhs, rhs);
- case ir::Binary::Kind::kEqual:
- return b.Equal(lhs, rhs);
- case ir::Binary::Kind::kNotEqual:
- return b.NotEqual(lhs, rhs);
- case ir::Binary::Kind::kLessThan:
- return b.LessThan(lhs, rhs);
- case ir::Binary::Kind::kGreaterThan:
- return b.GreaterThan(lhs, rhs);
- case ir::Binary::Kind::kLessThanEqual:
- return b.LessThanEqual(lhs, rhs);
- case ir::Binary::Kind::kGreaterThanEqual:
- return b.GreaterThanEqual(lhs, rhs);
- case ir::Binary::Kind::kShiftLeft:
- return b.Shl(lhs, rhs);
- case ir::Binary::Kind::kShiftRight:
- return b.Shr(lhs, rhs);
- }
- return b.Expr("<error>");
- }
-
////////////////////////////////////////////////////////////////////////////////////////////////
// Types
//
@@ -563,25 +525,124 @@
}
////////////////////////////////////////////////////////////////////////////////////////////////
- // Helpers
+ // Bindings
////////////////////////////////////////////////////////////////////////////////////////////////
- /// Creates and returns a new, unique name for the instructions result value, or returns the
- /// previously created name. Must not be called with a multi-result instruction.
- /// @return the instruction values name
- Symbol AssignNameTo(Instruction* inst) { return AssignNameTo(inst->Result()); }
-
/// Creates and returns a new, unique name for the given value, or returns the previously
/// created name.
/// @return the value's name
- Symbol AssignNameTo(Value* value) {
+ Symbol BindName(Value* value, std::string_view suggested = {}) {
TINT_ASSERT(IR, value);
- return value_names_.GetOrCreate(value, [&] {
- if (auto sym = mod.NameOf(value)) {
- return b.Symbols().New(sym.Name());
+ auto& existing = bindings_.GetOrCreate(value, [&] {
+ if (!suggested.empty()) {
+ return b.Symbols().New(suggested);
}
- return b.Symbols().New("v" + std::to_string(value_names_.Count()));
+ if (auto sym = mod.NameOf(value)) {
+ return b.Symbols().New(sym.NameView());
+ }
+ return b.Symbols().New("v");
});
+ if (auto* name = std::get_if<Symbol>(&existing); TINT_LIKELY(name)) {
+ return *name;
+ }
+
+ TINT_ICE(IR, b.Diagnostics()) << "BindName(" << value->TypeInfo().name
+ << ") called on value that has non-name binding";
+ return {};
+ }
+
+ template <typename T>
+ void Bind(ir::Value* value, const T* expr) {
+ TINT_ASSERT(IR, value);
+ if (CanInline(value)) {
+ // Value will be inlined at its place of usage.
+ bool added = bindings_.Add(value, expr);
+ if (TINT_UNLIKELY(!added)) {
+ TINT_ICE(IR, b.Diagnostics())
+ << "Bind(" << value->TypeInfo().name << ") called twice for same node";
+ }
+ } else {
+ Append(b.Decl(b.Let(BindName(value), expr)));
+ }
+ }
+
+ /// @returns true if the if the value can be inlined into its single place
+ /// of usage. Currently a value is inlined if it has a single usage and is unnamed.
+ /// TODO(crbug.com/tint/1902): This logic needs to check that the sequence of side-effecting
+ /// expressions is not changed by inlining the expression. This needs fixing.
+ bool CanInline(Value* val) { return val->Usages().Count() == 1 && !mod.NameOf(val).IsValid(); }
+
+ ////////////////////////////////////////////////////////////////////////////////////////////////
+ // Helpers
+ ////////////////////////////////////////////////////////////////////////////////////////////////
+ bool IsShortCircuit(ir::If* i,
+ const StatementList& true_stmts,
+ const StatementList& false_stmts) {
+ if (!i->HasResults()) {
+ return false;
+ }
+ auto* result = i->Result();
+ if (!result->Type()->Is<type::Bool>()) {
+ return false; // Wrong result type
+ }
+ if (i->Exits().Count() != 2) {
+ return false; // Doesn't have two exits
+ }
+ if (!true_stmts.IsEmpty() || !false_stmts.IsEmpty()) {
+ return false; // True or False blocks contain statements
+ }
+
+ auto* cond = i->Condition();
+ auto* true_val = i->True()->Back()->Operands().Front();
+ auto* false_val = i->False()->Back()->Operands().Front();
+ if (IsConstant(false_val, false)) {
+ // %res = if %cond {
+ // block { # true
+ // exit_if %true_val;
+ // }
+ // block { # false
+ // exit_if false;
+ // }
+ // }
+ //
+ // transform into:
+ //
+ // res = cond && true_val;
+ //
+ auto* lhs = Expr(cond);
+ auto* rhs = Expr(true_val);
+ Bind(result, b.LogicalAnd(lhs, rhs));
+ return true;
+ }
+ if (IsConstant(true_val, true)) {
+ // %res = if %cond {
+ // block { # true
+ // exit_if true;
+ // }
+ // block { # false
+ // exit_if %false_val;
+ // }
+ // }
+ //
+ // transform into:
+ //
+ // res = cond || false_val;
+ //
+ auto* lhs = Expr(cond);
+ auto* rhs = Expr(false_val);
+ Bind(result, b.LogicalOr(lhs, rhs));
+ return true;
+ }
+ return false;
+ }
+
+ bool IsConstant(ir::Value* val, bool value) {
+ if (auto* c = val->As<ir::Constant>()) {
+ if (c->Type()->Is<type::Bool>()) {
+ return c->Value()->ValueAs<bool>() == value;
+ }
+ }
+ return false;
}
};
diff --git a/src/tint/ir/to_program_roundtrip_test.cc b/src/tint/ir/to_program_roundtrip_test.cc
index 587d999..2d91267 100644
--- a/src/tint/ir/to_program_roundtrip_test.cc
+++ b/src/tint/ir/to_program_roundtrip_test.cc
@@ -39,12 +39,14 @@
auto ir_module = FromProgram(&input_program);
ASSERT_TRUE(ir_module);
+ tint::ir::Disassembler d{ir_module.Get()};
+ auto disassembly = d.Disassemble();
+
auto output_program = ToProgram(ir_module.Get());
if (!output_program.IsValid()) {
- tint::ir::Disassembler d{ir_module.Get()};
FAIL() << output_program.Diagnostics().str() << std::endl //
<< "IR:" << std::endl //
- << d.Disassemble() << std::endl //
+ << disassembly << std::endl //
<< "AST:" << std::endl //
<< Program::printer(&output_program) << std::endl;
}
@@ -56,10 +58,7 @@
auto expected = expected_wgsl.empty() ? input : utils::TrimSpace(expected_wgsl);
auto got = utils::TrimSpace(output.wgsl);
- if (expected != got) {
- tint::ir::Disassembler d{ir_module.Get()};
- EXPECT_EQ(expected, got) << "IR:" << std::endl << d.Disassemble();
- }
+ EXPECT_EQ(expected, got) << "IR:" << std::endl << disassembly;
}
void Test(std::string_view wgsl) { Test(wgsl, wgsl); }
@@ -266,8 +265,7 @@
// Short-circuiting binary ops
////////////////////////////////////////////////////////////////////////////////
-// TODO(crbug.com/tint/1902): Pattern detect this
-TEST_F(IRToProgramRoundtripTest, DISABLED_BinaryOp_LogicalAnd) {
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Param_2) {
Test(R"(
fn f(a : bool, b : bool) -> bool {
return (a && b);
@@ -275,11 +273,334 @@
)");
}
-// TODO(crbug.com/tint/1902): Pattern detect this
-TEST_F(IRToProgramRoundtripTest, DISABLED_BinaryOp_LogicalOr) {
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Param_3_ab_c) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ return ((a && b) && c);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Param_3_a_bc) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ return ((a && b) && c);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Let_2) {
Test(R"(
fn f(a : bool, b : bool) -> bool {
- return (a && b);
+ let l = (a && b);
+ return l;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Let_3_ab_c) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ let l = ((a && b) && c);
+ return l;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Let_3_a_bc) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ let l = (a && (b && c));
+ return l;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Call_2) {
+ Test(R"(
+fn a() -> bool {
+ return true;
+}
+
+fn b() -> bool {
+ return true;
+}
+
+fn f() -> bool {
+ return (a() && b());
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Call_3_ab_c) {
+ Test(R"(
+fn a() -> bool {
+ return true;
+}
+
+fn b() -> bool {
+ return true;
+}
+
+fn c() -> bool {
+ return true;
+}
+
+fn f() -> bool {
+ return ((a() && b()) && c());
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalAnd_Call_3_a_bc) {
+ Test(R"(
+fn a() -> bool {
+ return true;
+}
+
+fn b() -> bool {
+ return true;
+}
+
+fn c() -> bool {
+ return true;
+}
+
+fn f() -> bool {
+ return (a() && (b() && c()));
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Param_2) {
+ Test(R"(
+fn f(a : bool, b : bool) -> bool {
+ return (a || b);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Param_3_ab_c) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ return ((a || b) || c);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Param_3_a_bc) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ return (a || (b || c));
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Let_2) {
+ Test(R"(
+fn f(a : bool, b : bool) -> bool {
+ let l = (a || b);
+ return l;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Let_3_ab_c) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ let l = ((a || b) || c);
+ return l;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Let_3_a_bc) {
+ Test(R"(
+fn f(a : bool, b : bool, c : bool) -> bool {
+ let l = (a || (b || c));
+ return l;
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Call_2) {
+ Test(R"(
+fn a() -> bool {
+ return true;
+}
+
+fn b() -> bool {
+ return true;
+}
+
+fn f() -> bool {
+ return (a() || b());
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Call_3_ab_c) {
+ Test(R"(
+fn a() -> bool {
+ return true;
+}
+
+fn b() -> bool {
+ return true;
+}
+
+fn c() -> bool {
+ return true;
+}
+
+fn f() -> bool {
+ return ((a() || b()) || c());
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalOr_Call_3_a_bc) {
+ Test(R"(
+fn a() -> bool {
+ return true;
+}
+
+fn b() -> bool {
+ return true;
+}
+
+fn c() -> bool {
+ return true;
+}
+
+fn f() -> bool {
+ return (a() || (b() || c()));
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, BinaryOp_LogicalMixed) {
+ Test(R"(
+fn b() -> bool {
+ return true;
+}
+
+fn d() -> bool {
+ return true;
+}
+
+fn f(a : bool, c : bool) -> bool {
+ let l = ((a || b()) && (c || d()));
+ return l;
+}
+)");
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// Compound assignment
+////////////////////////////////////////////////////////////////////////////////
+
+TEST_F(IRToProgramRoundtripTest, CompoundAssign_Increment) {
+ Test(R"(
+fn f() {
+ var v : i32;
+ v++;
+}
+)",
+ R"(
+fn f() {
+ var v : i32;
+ v = (v + 1i);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, CompoundAssign_Decrement) {
+ Test(R"(
+fn f() {
+ var v : i32;
+ v++;
+}
+)",
+ R"(
+fn f() {
+ var v : i32;
+ v = (v + 1i);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, CompoundAssign_Add) {
+ Test(R"(
+fn f() {
+ var v : i32;
+ v += 8i;
+}
+)",
+ R"(
+fn f() {
+ var v : i32;
+ v = (v + 8i);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, CompoundAssign_Subtract) {
+ Test(R"(
+fn f() {
+ var v : i32;
+ v -= 8i;
+}
+)",
+ R"(
+fn f() {
+ var v : i32;
+ v = (v - 8i);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, CompoundAssign_Multiply) {
+ Test(R"(
+fn f() {
+ var v : i32;
+ v *= 8i;
+}
+)",
+ R"(
+fn f() {
+ var v : i32;
+ v = (v * 8i);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, CompoundAssign_Divide) {
+ Test(R"(
+fn f() {
+ var v : i32;
+ v /= 8i;
+}
+)",
+ R"(
+fn f() {
+ var v : i32;
+ v = (v / 8i);
+}
+)");
+}
+
+TEST_F(IRToProgramRoundtripTest, CompoundAssign_Xor) {
+ Test(R"(
+fn f() {
+ var v : i32;
+ v ^= 8i;
+}
+)",
+ R"(
+fn f() {
+ var v : i32;
+ v = (v ^ 8i);
}
)");
}
@@ -450,6 +771,26 @@
)");
}
+TEST_F(IRToProgramRoundtripTest, If_Else_Chain) {
+ Test(R"(
+fn x(i : i32) -> bool {
+ return true;
+}
+
+fn f(a : bool, b : bool, c : bool, d : bool) {
+ if (a) {
+ x(0i);
+ } else if (b) {
+ x(1i);
+ } else if (c) {
+ x(2i);
+ } else {
+ x(3i);
+ }
+}
+)");
+}
+
////////////////////////////////////////////////////////////////////////////////
// Switch
////////////////////////////////////////////////////////////////////////////////
diff --git a/src/tint/ir/transform/block_decorated_structs_test.cc b/src/tint/ir/transform/block_decorated_structs_test.cc
index 92e3922..2211f99 100644
--- a/src/tint/ir/transform/block_decorated_structs_test.cc
+++ b/src/tint/ir/transform/block_decorated_structs_test.cc
@@ -64,8 +64,7 @@
tint_symbol:i32 @offset(0)
}
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<uniform, tint_symbol_1, read_write> = var @binding_point(0, 0)
}
@@ -98,8 +97,7 @@
tint_symbol:i32 @offset(0)
}
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
}
@@ -136,8 +134,7 @@
tint_symbol:array<i32> @offset(0)
}
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
}
@@ -191,8 +188,7 @@
arr:array<i32> @offset(4)
}
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<storage, tint_symbol, read_write> = var @binding_point(0, 0)
}
@@ -240,8 +236,7 @@
tint_symbol:MyStruct @offset(0)
}
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
%2:ptr<private, MyStruct, read_write> = var
}
@@ -293,8 +288,7 @@
tint_symbol_4:i32 @offset(0)
}
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<storage, tint_symbol_1, read_write> = var @binding_point(0, 0)
%2:ptr<storage, tint_symbol_3, read_write> = var @binding_point(0, 1)
%3:ptr<storage, tint_symbol_5, read_write> = var @binding_point(0, 2)
diff --git a/src/tint/ir/transform/merge_return.cc b/src/tint/ir/transform/merge_return.cc
index 0b95a3a..4817d04 100644
--- a/src/tint/ir/transform/merge_return.cc
+++ b/src/tint/ir/transform/merge_return.cc
@@ -91,7 +91,7 @@
}
// Look to see if the function ends with a return
- fn_return = tint::As<Return>(fn->StartTarget()->Branch());
+ fn_return = tint::As<Return>(fn->StartTarget()->Terminator());
// Process the function's block.
// This will traverse into control instructions that hold returns, and apply the necessary
@@ -139,7 +139,7 @@
if (inst->Is<Unreachable>()) {
// Unreachable can become reachable once returns are turned into exits.
- // As this is the terminator for the branch, simply stop processing the
+ // As this is the terminator for the block, simply stop processing the
// instructions. A appropriate terminator will be created for this block below.
inst->Remove();
break;
@@ -170,8 +170,8 @@
// new_value_with_type returns a new RuntimeValue with the same type as 'v'
auto new_value_with_type = [&](Value* v) { return b.InstructionResult(v->Type()); };
- if (inner_if->True()->HasBranchTarget()) {
- if (auto* exit_if = inner_if->True()->Branch()->As<ExitIf>()) {
+ if (inner_if->True()->HasTerminator()) {
+ if (auto* exit_if = inner_if->True()->Terminator()->As<ExitIf>()) {
// Ensure the associated 'if' is updated.
exit_if->SetIf(inner_if);
@@ -190,7 +190,7 @@
// Loop over the 'if' instructions, starting with the inner-most, and add any missing
// terminating instructions to the blocks holding the 'if'.
for (auto* i = inner_if; i; i = tint::As<If>(i->Block()->Parent())) {
- if (!i->Block()->HasBranchTarget()) {
+ if (!i->Block()->HasTerminator()) {
// Append the exit instruction to the block holding the 'if'.
utils::Vector<InstructionResult*, 8> exit_args = i->Results();
if (!i->HasResults()) {
diff --git a/src/tint/ir/transform/merge_return_test.cc b/src/tint/ir/transform/merge_return_test.cc
index cf2a2cc..bf68eae 100644
--- a/src/tint/ir/transform/merge_return_test.cc
+++ b/src/tint/ir/transform/merge_return_test.cc
@@ -73,19 +73,16 @@
auto* src = R"(
%foo = func(%2:i32):i32 -> %b1 {
%b1 = block {
- %3:i32 = if %4 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ %3:i32 = if %4 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
%5:i32 = add %2, 1i
- exit_if %5
+ exit_if %5 # if_1
}
-
- # False block
- %b3 = block {
+ %b3 = block { # false
%6:i32 = add %2, 2i
- exit_if %6
+ exit_if %6 # if_1
}
-
+ }
ret %3
}
}
@@ -125,26 +122,23 @@
auto* src = R"(
%foo = func(%2:i32):i32 -> %b1 {
%b1 = block {
- switch %2 [c: (default, %b2)]
- # Case block
- %b2 = block {
- exit_switch
+ switch %2 [c: (default, %b2)] { # switch_1
+ %b2 = block { # case
+ exit_switch # switch_1
}
-
- loop []
- %3:i32 = if %4 [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ }
+ loop [] { # loop_1
+ }
+ %3:i32 = if %4 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
%5:i32 = add %2, 1i
- exit_if %5
+ exit_if %5 # if_1
}
-
- # False block
- %b4 = block {
+ %b4 = block { # false
%6:i32 = add %2, 2i
- exit_if %6
+ exit_if %6 # if_1
}
-
+ }
ret %3
}
}
@@ -177,17 +171,14 @@
auto* src = R"(
%foo = func(%2:bool):void -> %b1 {
%b1 = block {
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
ret
}
}
@@ -197,17 +188,14 @@
auto* expect = R"(
%foo = func(%2:bool):void -> %b1 {
%b1 = block {
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
- exit_if
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ exit_if # if_1
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
ret
}
}
@@ -239,17 +227,14 @@
auto* src = R"(
%foo = func(%2:bool):void -> %b1 {
%b1 = block {
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
ret
}
}
@@ -259,17 +244,14 @@
auto* expect = R"(
%foo = func(%2:bool):void -> %b1 {
%b1 = block {
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
- exit_if
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ exit_if # if_1
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
ret
}
}
@@ -299,17 +281,14 @@
auto* src = R"(
%foo = func(%2:bool):i32 -> %b1 {
%b1 = block {
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret 1i
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
ret 2i
}
}
@@ -321,27 +300,23 @@
%b1 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
store %continue_execution, false
store %return_value, 1i
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
%5:bool = load %continue_execution
- if %5 [t: %b4]
- # True block
- %b4 = block {
+ if %5 [t: %b4] { # if_2
+ %b4 = block { # true
store %return_value, 2i
- exit_if
+ exit_if # if_2
}
-
+ }
%6:i32 = load %return_value
ret %6
}
@@ -373,17 +348,14 @@
auto* src = R"(
%foo = func(%2:bool):i32 -> %b1 {
%b1 = block {
- %3:i32 = if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ %3:i32 = if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret 1i
}
-
- # False block
- %b3 = block {
- exit_if 2i
+ %b3 = block { # false
+ exit_if 2i # if_1
}
-
+ }
ret %3
}
}
@@ -395,27 +367,23 @@
%b1 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- %5:i32 = if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ %5:i32 = if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
store %continue_execution, false
store %return_value, 1i
- exit_if undef
+ exit_if undef # if_1
}
-
- # False block
- %b3 = block {
- exit_if 2i
+ %b3 = block { # false
+ exit_if 2i # if_1
}
-
+ }
%6:bool = load %continue_execution
- if %6 [t: %b4]
- # True block
- %b4 = block {
+ if %6 [t: %b4] { # if_2
+ %b4 = block { # true
store %return_value, %5
- exit_if
+ exit_if # if_2
}
-
+ }
%7:i32 = load %return_value
ret %7
}
@@ -440,24 +408,21 @@
auto tb = b.With(ifelse->True());
tb.Return(func, 1_i);
auto fb = b.With(ifelse->False());
- fb.ExitIf(ifelse, b.Constant(8_i));
+ fb.ExitIf(ifelse, nullptr);
sb.Return(func, ifelse->Result(0));
auto* src = R"(
%foo = func(%2:bool):i32 -> %b1 {
%b1 = block {
- %3:i32 = if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ %3:i32 = if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret 1i
}
-
- # False block
- %b3 = block {
- exit_if 8i
+ %b3 = block { # false
+ exit_if undef # if_1
}
-
+ }
ret %3
}
}
@@ -469,27 +434,23 @@
%b1 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- %5:i32 = if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ %5:i32 = if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
store %continue_execution, false
store %return_value, 1i
- exit_if undef
+ exit_if undef # if_1
}
-
- # False block
- %b3 = block {
- exit_if 8i
+ %b3 = block { # false
+ exit_if undef # if_1
}
-
+ }
%6:bool = load %continue_execution
- if %6 [t: %b4]
- # True block
- %b4 = block {
+ if %6 [t: %b4] { # if_2
+ %b4 = block { # true
store %return_value, %5
- exit_if
+ exit_if # if_2
}
-
+ }
%7:i32 = load %return_value
ret %7
}
@@ -520,17 +481,14 @@
auto* src = R"(
%foo = func(%2:bool):void -> %b1 {
%b1 = block {
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret
}
-
- # False block
- %b3 = block {
+ %b3 = block { # false
ret
}
-
+ }
unreachable
}
}
@@ -541,17 +499,14 @@
auto* expect = R"(
%foo = func(%2:bool):void -> %b1 {
%b1 = block {
- if %2 [t: %b2, f: %b3]
- # True block
- %b2 = block {
- exit_if
+ if %2 [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
+ exit_if # if_1
}
-
- # False block
- %b3 = block {
- exit_if
+ %b3 = block { # false
+ exit_if # if_1
}
-
+ }
ret
}
}
@@ -583,24 +538,20 @@
sb.Return(func);
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:bool):void -> %b2 {
%b2 = block {
- if %3 [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %3 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
ret
}
-
- # False block
- %b4 = block {
- exit_if
+ %b4 = block { # false
+ exit_if # if_1
}
-
+ }
store %1, 42i
ret
}
@@ -610,34 +561,29 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:bool):void -> %b2 {
%b2 = block {
%continue_execution:ptr<function, bool, read_write> = var, true
- if %3 [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %3 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
store %continue_execution, false
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b4 = block {
- exit_if
+ %b4 = block { # false
+ exit_if # if_1
}
-
+ }
%5:bool = load %continue_execution
- if %5 [t: %b5]
- # True block
- %b5 = block {
+ if %5 [t: %b5] { # if_2
+ %b5 = block { # true
store %1, 42i
- exit_if
+ exit_if # if_2
}
-
+ }
ret
}
}
@@ -671,24 +617,20 @@
fb.ExitIf(ifelse);
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:bool):void -> %b2 {
%b2 = block {
- if %3 [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %3 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
ret
}
-
- # False block
- %b4 = block {
- exit_if
+ %b4 = block { # false
+ exit_if # if_1
}
-
+ }
store %1, 42i
ret
}
@@ -698,34 +640,29 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:bool):void -> %b2 {
%b2 = block {
%continue_execution:ptr<function, bool, read_write> = var, true
- if %3 [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %3 [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
store %continue_execution, false
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b4 = block {
- exit_if
+ %b4 = block { # false
+ exit_if # if_1
}
-
+ }
%5:bool = load %continue_execution
- if %5 [t: %b5]
- # True block
- %b5 = block {
+ if %5 [t: %b5] { # if_2
+ %b5 = block { # true
store %1, 42i
- exit_if
+ exit_if # if_2
}
-
+ }
ret
}
}
@@ -778,48 +715,38 @@
middle_true.Return(func, 2_i);
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%condA:bool, %condB:bool, %condC:bool):i32 -> %b2 {
%b2 = block {
- if %condA [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %condA [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
ret 3i
}
-
- # False block
- %b4 = block {
- if %condB [t: %b5, f: %b6]
- # True block
- %b5 = block {
- if %condC [t: %b7, f: %b8]
- # True block
- %b7 = block {
+ %b4 = block { # false
+ if %condB [t: %b5, f: %b6] { # if_2
+ %b5 = block { # true
+ if %condC [t: %b7, f: %b8] { # if_3
+ %b7 = block { # true
ret 1i
}
-
- # False block
- %b8 = block {
- exit_if
+ %b8 = block { # false
+ exit_if # if_3
}
-
+ }
store %1, 1i
ret 2i
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_2
}
-
+ }
store %1, 2i
- exit_if
+ exit_if # if_1
}
-
+ }
store %1, 3i
%6:i32 = add 5i, 6i
ret %6
@@ -830,8 +757,7 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
@@ -839,71 +765,59 @@
%b2 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- if %condA [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %condA [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
store %continue_execution, false
store %return_value, 3i
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b4 = block {
- if %condB [t: %b5, f: %b6]
- # True block
- %b5 = block {
- if %condC [t: %b7, f: %b8]
- # True block
- %b7 = block {
+ %b4 = block { # false
+ if %condB [t: %b5, f: %b6] { # if_2
+ %b5 = block { # true
+ if %condC [t: %b7, f: %b8] { # if_3
+ %b7 = block { # true
store %continue_execution, false
store %return_value, 1i
- exit_if
+ exit_if # if_3
}
-
- # False block
- %b8 = block {
- exit_if
+ %b8 = block { # false
+ exit_if # if_3
}
-
+ }
%8:bool = load %continue_execution
- if %8 [t: %b9]
- # True block
- %b9 = block {
+ if %8 [t: %b9] { # if_4
+ %b9 = block { # true
store %1, 1i
store %continue_execution, false
store %return_value, 2i
- exit_if
+ exit_if # if_4
}
-
- exit_if
+ }
+ exit_if # if_2
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_2
}
-
+ }
%9:bool = load %continue_execution
- if %9 [t: %b10]
- # True block
- %b10 = block {
+ if %9 [t: %b10] { # if_5
+ %b10 = block { # true
store %1, 2i
- exit_if
+ exit_if # if_5
}
-
- exit_if
+ }
+ exit_if # if_1
}
-
+ }
%10:bool = load %continue_execution
- if %10 [t: %b11]
- # True block
- %b11 = block {
+ if %10 [t: %b11] { # if_6
+ %b11 = block { # true
store %1, 3i
%11:i32 = add 5i, 6i
store %return_value, %11
- exit_if
+ exit_if # if_6
}
-
+ }
%12:i32 = load %return_value
ret %12
}
@@ -954,46 +868,36 @@
middle_true.ExitIf(ifelse_middle);
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%condA:bool, %condB:bool, %condC:bool):i32 -> %b2 {
%b2 = block {
- if %condA [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %condA [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
ret 3i
}
-
- # False block
- %b4 = block {
- if %condB [t: %b5, f: %b6]
- # True block
- %b5 = block {
- if %condC [t: %b7, f: %b8]
- # True block
- %b7 = block {
+ %b4 = block { # false
+ if %condB [t: %b5, f: %b6] { # if_2
+ %b5 = block { # true
+ if %condC [t: %b7, f: %b8] { # if_3
+ %b7 = block { # true
ret 1i
}
-
- # False block
- %b8 = block {
- exit_if
+ %b8 = block { # false
+ exit_if # if_3
}
-
- exit_if
+ }
+ exit_if # if_2
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_2
}
-
- exit_if
+ }
+ exit_if # if_1
}
-
+ }
ret 3i
}
}
@@ -1002,8 +906,7 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
@@ -1011,51 +914,41 @@
%b2 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- if %condA [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ if %condA [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
store %continue_execution, false
store %return_value, 3i
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b4 = block {
- if %condB [t: %b5, f: %b6]
- # True block
- %b5 = block {
- if %condC [t: %b7, f: %b8]
- # True block
- %b7 = block {
+ %b4 = block { # false
+ if %condB [t: %b5, f: %b6] { # if_2
+ %b5 = block { # true
+ if %condC [t: %b7, f: %b8] { # if_3
+ %b7 = block { # true
store %continue_execution, false
store %return_value, 1i
- exit_if
+ exit_if # if_3
}
-
- # False block
- %b8 = block {
- exit_if
+ %b8 = block { # false
+ exit_if # if_3
}
-
- exit_if
+ }
+ exit_if # if_2
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_2
}
-
- exit_if
+ }
+ exit_if # if_1
}
-
+ }
%8:bool = load %continue_execution
- if %8 [t: %b9]
- # True block
- %b9 = block {
+ if %8 [t: %b9] { # if_4
+ %b9 = block { # true
store %return_value, 3i
- exit_if
+ exit_if # if_4
}
-
+ }
%9:i32 = load %return_value
ret %9
}
@@ -1108,49 +1001,39 @@
middle_true.ExitIf(ifelse_middle, middle_true.Add(ty.i32(), 42_i, 1_i));
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%condA:bool, %condB:bool, %condC:bool):i32 -> %b2 {
%b2 = block {
- %6:i32 = if %condA [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ %6:i32 = if %condA [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
ret 3i
}
-
- # False block
- %b4 = block {
- %7:i32 = if %condB [t: %b5, f: %b6]
- # True block
- %b5 = block {
- if %condC [t: %b7, f: %b8]
- # True block
- %b7 = block {
+ %b4 = block { # false
+ %7:i32 = if %condB [t: %b5, f: %b6] { # if_2
+ %b5 = block { # true
+ if %condC [t: %b7, f: %b8] { # if_3
+ %b7 = block { # true
ret 1i
}
-
- # False block
- %b8 = block {
- exit_if
+ %b8 = block { # false
+ exit_if # if_3
}
-
+ }
%8:i32 = add 42i, 1i
- exit_if %8
+ exit_if %8 # if_2
}
-
- # False block
- %b6 = block {
+ %b6 = block { # false
%9:i32 = add 43i, 2i
- exit_if %9
+ exit_if %9 # if_2
}
-
+ }
%10:i32 = add %7, 1i
- exit_if %10
+ exit_if %10 # if_1
}
-
+ }
%11:i32 = add %6, 1i
ret %11
}
@@ -1160,8 +1043,7 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
@@ -1169,69 +1051,59 @@
%b2 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- %8:i32 = if %condA [t: %b3, f: %b4]
- # True block
- %b3 = block {
+ %8:i32 = if %condA [t: %b3, f: %b4] { # if_1
+ %b3 = block { # true
store %continue_execution, false
store %return_value, 3i
- exit_if undef
+ exit_if undef # if_1
}
-
- # False block
- %b4 = block {
- %9:i32 = if %condB [t: %b5, f: %b6]
- # True block
- %b5 = block {
- if %condC [t: %b7, f: %b8]
- # True block
- %b7 = block {
+ %b4 = block { # false
+ %9:i32 = if %condB [t: %b5, f: %b6] { # if_2
+ %b5 = block { # true
+ if %condC [t: %b7, f: %b8] { # if_3
+ %b7 = block { # true
store %continue_execution, false
store %return_value, 1i
- exit_if
+ exit_if # if_3
}
-
- # False block
- %b8 = block {
- exit_if
+ %b8 = block { # false
+ exit_if # if_3
}
-
+ }
%10:bool = load %continue_execution
- %11:i32 = if %10 [t: %b9]
- # True block
- %b9 = block {
+ %11:i32 = if %10 [t: %b9] { # if_4
+ %b9 = block { # true
%12:i32 = add 42i, 1i
- exit_if %12
+ exit_if %12 # if_4
}
-
- exit_if %11
+ # implicit false block: exit_if undef
+ }
+ exit_if %11 # if_2
}
-
- # False block
- %b6 = block {
+ %b6 = block { # false
%13:i32 = add 43i, 2i
- exit_if %13
+ exit_if %13 # if_2
}
-
+ }
%14:bool = load %continue_execution
- %15:i32 = if %14 [t: %b10]
- # True block
- %b10 = block {
+ %15:i32 = if %14 [t: %b10] { # if_5
+ %b10 = block { # true
%16:i32 = add %9, 1i
- exit_if %16
+ exit_if %16 # if_5
}
-
- exit_if %15
+ # implicit false block: exit_if undef
+ }
+ exit_if %15 # if_1
}
-
+ }
%17:bool = load %continue_execution
- if %17 [t: %b11]
- # True block
- %b11 = block {
+ if %17 [t: %b11] { # if_6
+ %b11 = block { # true
%18:i32 = add %8, 1i
store %return_value, %18
- exit_if
+ exit_if # if_6
}
-
+ }
%19:i32 = load %return_value
ret %19
}
@@ -1257,12 +1129,11 @@
auto* src = R"(
%foo = func():i32 -> %b1 {
%b1 = block {
- loop [b: %b2]
- # Body block
- %b2 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block { # body
ret 42i
}
-
+ }
unreachable
}
}
@@ -1273,13 +1144,12 @@
%foo = func():i32 -> %b1 {
%b1 = block {
%return_value:ptr<function, i32, read_write> = var
- loop [b: %b2]
- # Body block
- %b2 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block { # body
store %return_value, 42i
- exit_loop
+ exit_loop # loop_1
}
-
+ }
%3:i32 = load %return_value
ret %3
}
@@ -1322,37 +1192,30 @@
sb.Return(func, 43_i);
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:bool):i32 -> %b2 {
%b2 = block {
- loop [b: %b3, c: %b4]
- # Body block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ loop [b: %b3, c: %b4] { # loop_1
+ %b3 = block { # body
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
ret 42i
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
store %1, 2i
continue %b4
}
-
- # Continuing block
- %b4 = block {
+ %b4 = block { # continuing
store %1, 1i
break_if true %b3
}
-
+ }
store %1, 3i
ret 43i
}
@@ -1361,8 +1224,7 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
@@ -1370,48 +1232,40 @@
%b2 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- loop [b: %b3, c: %b4]
- # Body block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ loop [b: %b3, c: %b4] { # loop_1
+ %b3 = block { # body
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
store %continue_execution, false
store %return_value, 42i
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
%6:bool = load %continue_execution
- if %6 [t: %b7]
- # True block
- %b7 = block {
+ if %6 [t: %b7] { # if_2
+ %b7 = block { # true
store %1, 2i
continue %b4
}
-
- exit_loop
+ }
+ exit_loop # loop_1
}
-
- # Continuing block
- %b4 = block {
+ %b4 = block { # continuing
store %1, 1i
break_if true %b3
}
-
+ }
%7:bool = load %continue_execution
- if %7 [t: %b8]
- # True block
- %b8 = block {
+ if %7 [t: %b8] { # if_3
+ %b8 = block { # true
store %1, 3i
store %return_value, 43i
- exit_if
+ exit_if # if_3
}
-
+ }
%8:i32 = load %return_value
ret %8
}
@@ -1453,37 +1307,30 @@
sb.Unreachable();
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:bool):i32 -> %b2 {
%b2 = block {
- loop [b: %b3, c: %b4]
- # Body block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ loop [b: %b3, c: %b4] { # loop_1
+ %b3 = block { # body
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
ret 42i
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
store %1, 2i
continue %b4
}
-
- # Continuing block
- %b4 = block {
+ %b4 = block { # continuing
store %1, 1i
next_iteration %b3
}
-
+ }
unreachable
}
}
@@ -1491,8 +1338,7 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
@@ -1500,39 +1346,32 @@
%b2 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- loop [b: %b3, c: %b4]
- # Body block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ loop [b: %b3, c: %b4] { # loop_1
+ %b3 = block { # body
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
store %continue_execution, false
store %return_value, 42i
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
%6:bool = load %continue_execution
- if %6 [t: %b7]
- # True block
- %b7 = block {
+ if %6 [t: %b7] { # if_2
+ %b7 = block { # true
store %1, 2i
continue %b4
}
-
- exit_loop
+ }
+ exit_loop # loop_1
}
-
- # Continuing block
- %b4 = block {
+ %b4 = block { # continuing
store %1, 1i
next_iteration %b3
}
-
+ }
%7:i32 = load %return_value
ret %7
}
@@ -1576,37 +1415,30 @@
sb.Return(func, loop->Result(0));
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:bool):i32 -> %b2 {
%b2 = block {
- loop [b: %b3, c: %b4]
- # Body block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ loop [b: %b3, c: %b4] { # loop_1
+ %b3 = block { # body
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
ret 42i
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
store %1, 2i
continue %b4
}
-
- # Continuing block
- %b4 = block {
+ %b4 = block { # continuing
store %1, 1i
break_if true %b3 4i
}
-
+ }
store %1, 3i
ret %4
}
@@ -1615,8 +1447,7 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
@@ -1624,48 +1455,40 @@
%b2 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- loop [b: %b3, c: %b4]
- # Body block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ loop [b: %b3, c: %b4] { # loop_1
+ %b3 = block { # body
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
store %continue_execution, false
store %return_value, 42i
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
%6:bool = load %continue_execution
- if %6 [t: %b7]
- # True block
- %b7 = block {
+ if %6 [t: %b7] { # if_2
+ %b7 = block { # true
store %1, 2i
continue %b4
}
-
- exit_loop
+ }
+ exit_loop # loop_1
}
-
- # Continuing block
- %b4 = block {
+ %b4 = block { # continuing
store %1, 1i
break_if true %b3 4i
}
-
+ }
%7:bool = load %continue_execution
- if %7 [t: %b8]
- # True block
- %b8 = block {
+ if %7 [t: %b8] { # if_3
+ %b8 = block { # true
store %1, 3i
store %return_value, %8
- exit_if
+ exit_if # if_3
}
-
+ }
%9:i32 = load %return_value
ret %9
}
@@ -1696,17 +1519,14 @@
auto* src = R"(
%foo = func(%2:i32):i32 -> %b1 {
%b1 = block {
- switch %2 [c: (1i, %b2), c: (default, %b3)]
- # Case block
- %b2 = block {
+ switch %2 [c: (1i, %b2), c: (default, %b3)] { # switch_1
+ %b2 = block { # case
ret 42i
}
-
- # Case block
- %b3 = block {
- exit_switch
+ %b3 = block { # case
+ exit_switch # switch_1
}
-
+ }
ret 0i
}
}
@@ -1718,27 +1538,23 @@
%b1 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- switch %2 [c: (1i, %b2), c: (default, %b3)]
- # Case block
- %b2 = block {
+ switch %2 [c: (1i, %b2), c: (default, %b3)] { # switch_1
+ %b2 = block { # case
store %continue_execution, false
store %return_value, 42i
- exit_switch
+ exit_switch # switch_1
}
-
- # Case block
- %b3 = block {
- exit_switch
+ %b3 = block { # case
+ exit_switch # switch_1
}
-
+ }
%5:bool = load %continue_execution
- if %5 [t: %b4]
- # True block
- %b4 = block {
+ if %5 [t: %b4] { # if_1
+ %b4 = block { # true
store %return_value, 0i
- exit_if
+ exit_if # if_1
}
-
+ }
%6:i32 = load %return_value
ret %6
}
@@ -1779,36 +1595,29 @@
sb.Return(func, 0_i);
auto* src = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
%foo = func(%3:i32):i32 -> %b2 {
%b2 = block {
- switch %3 [c: (1i, %b3), c: (default, %b4)]
- # Case block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ switch %3 [c: (1i, %b3), c: (default, %b4)] { # switch_1
+ %b3 = block { # case
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
ret 42i
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
store %1, 2i
- exit_switch
+ exit_switch # switch_1
}
-
- # Case block
- %b4 = block {
- exit_switch
+ %b4 = block { # case
+ exit_switch # switch_1
}
-
+ }
ret 0i
}
}
@@ -1816,8 +1625,7 @@
EXPECT_EQ(src, str());
auto* expect = R"(
-# Root block
-%b1 = block {
+%b1 = block { # root
%1:ptr<private, i32, read_write> = var
}
@@ -1825,46 +1633,38 @@
%b2 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- switch %3 [c: (1i, %b3), c: (default, %b4)]
- # Case block
- %b3 = block {
- if %3 [t: %b5, f: %b6]
- # True block
- %b5 = block {
+ switch %3 [c: (1i, %b3), c: (default, %b4)] { # switch_1
+ %b3 = block { # case
+ if %3 [t: %b5, f: %b6] { # if_1
+ %b5 = block { # true
store %continue_execution, false
store %return_value, 42i
- exit_if
+ exit_if # if_1
}
-
- # False block
- %b6 = block {
- exit_if
+ %b6 = block { # false
+ exit_if # if_1
}
-
+ }
%6:bool = load %continue_execution
- if %6 [t: %b7]
- # True block
- %b7 = block {
+ if %6 [t: %b7] { # if_2
+ %b7 = block { # true
store %1, 2i
- exit_switch
+ exit_switch # switch_1
}
-
- exit_switch
+ }
+ exit_switch # switch_1
}
-
- # Case block
- %b4 = block {
- exit_switch
+ %b4 = block { # case
+ exit_switch # switch_1
}
-
+ }
%7:bool = load %continue_execution
- if %7 [t: %b8]
- # True block
- %b8 = block {
+ if %7 [t: %b8] { # if_3
+ %b8 = block { # true
store %return_value, 0i
- exit_if
+ exit_if # if_3
}
-
+ }
%8:i32 = load %return_value
ret %8
}
@@ -1900,27 +1700,20 @@
auto* src = R"(
%foo = func(%2:i32):i32 -> %b1 {
%b1 = block {
- switch %2 [c: (1i, %b2), c: (2i, %b3), c: (3i, %b4), c: (default, %b5)]
- # Case block
- %b2 = block {
+ switch %2 [c: (1i, %b2), c: (2i, %b3), c: (3i, %b4), c: (default, %b5)] { # switch_1
+ %b2 = block { # case
ret 42i
}
-
- # Case block
- %b3 = block {
+ %b3 = block { # case
ret 99i
}
-
- # Case block
- %b4 = block {
- exit_switch 1i
+ %b4 = block { # case
+ exit_switch 1i # switch_1
}
-
- # Case block
- %b5 = block {
- exit_switch 0i
+ %b5 = block { # case
+ exit_switch 0i # switch_1
}
-
+ }
ret %3
}
}
@@ -1932,39 +1725,31 @@
%b1 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- switch %2 [c: (1i, %b2), c: (2i, %b3), c: (3i, %b4), c: (default, %b5)]
- # Case block
- %b2 = block {
+ switch %2 [c: (1i, %b2), c: (2i, %b3), c: (3i, %b4), c: (default, %b5)] { # switch_1
+ %b2 = block { # case
store %continue_execution, false
store %return_value, 42i
- exit_switch undef
+ exit_switch undef # switch_1
}
-
- # Case block
- %b3 = block {
+ %b3 = block { # case
store %continue_execution, false
store %return_value, 99i
- exit_switch undef
+ exit_switch undef # switch_1
}
-
- # Case block
- %b4 = block {
- exit_switch 1i
+ %b4 = block { # case
+ exit_switch 1i # switch_1
}
-
- # Case block
- %b5 = block {
- exit_switch 0i
+ %b5 = block { # case
+ exit_switch 0i # switch_1
}
-
+ }
%5:bool = load %continue_execution
- if %5 [t: %b6]
- # True block
- %b6 = block {
+ if %5 [t: %b6] { # if_1
+ %b6 = block { # true
store %return_value, %6
- exit_if
+ exit_if # if_1
}
-
+ }
%7:i32 = load %return_value
ret %7
}
@@ -1995,18 +1780,16 @@
auto* src = R"(
%foo = func():void -> %b1 {
%b1 = block {
- loop [b: %b2]
- # Body block
- %b2 = block {
- if true [t: %b3]
- # True block
- %b3 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block { # body
+ if true [t: %b3] { # if_1
+ %b3 = block { # true
ret
}
-
+ }
continue %b4
}
-
+ }
unreachable
}
}
@@ -2017,26 +1800,23 @@
%foo = func():void -> %b1 {
%b1 = block {
%continue_execution:ptr<function, bool, read_write> = var, true
- loop [b: %b2]
- # Body block
- %b2 = block {
- if true [t: %b3]
- # True block
- %b3 = block {
+ loop [b: %b2] { # loop_1
+ %b2 = block { # body
+ if true [t: %b3] { # if_1
+ %b3 = block { # true
store %continue_execution, false
- exit_if
+ exit_if # if_1
}
-
+ }
%3:bool = load %continue_execution
- if %3 [t: %b4]
- # True block
- %b4 = block {
+ if %3 [t: %b4] { # if_2
+ %b4 = block { # true
continue %b5
}
-
- exit_loop
+ }
+ exit_loop # loop_1
}
-
+ }
ret
}
}
@@ -2065,18 +1845,16 @@
auto* src = R"(
%foo = func():i32 -> %b1 {
%b1 = block {
- if true [t: %b2]
- # True block
- %b2 = block {
- if true [t: %b3]
- # True block
- %b3 = block {
+ if true [t: %b2] { # if_1
+ %b2 = block { # true
+ if true [t: %b3] { # if_2
+ %b3 = block { # true
ret 1i
}
-
+ }
ret 2i
}
-
+ }
ret 3i
}
}
@@ -2088,37 +1866,33 @@
%b1 = block {
%return_value:ptr<function, i32, read_write> = var
%continue_execution:ptr<function, bool, read_write> = var, true
- if true [t: %b2]
- # True block
- %b2 = block {
- if true [t: %b3]
- # True block
- %b3 = block {
+ if true [t: %b2] { # if_1
+ %b2 = block { # true
+ if true [t: %b3] { # if_2
+ %b3 = block { # true
store %continue_execution, false
store %return_value, 1i
- exit_if
+ exit_if # if_2
}
-
+ }
%4:bool = load %continue_execution
- if %4 [t: %b4]
- # True block
- %b4 = block {
+ if %4 [t: %b4] { # if_3
+ %b4 = block { # true
store %continue_execution, false
store %return_value, 2i
- exit_if
+ exit_if # if_3
}
-
- exit_if
+ }
+ exit_if # if_1
}
-
+ }
%5:bool = load %continue_execution
- if %5 [t: %b5]
- # True block
- %b5 = block {
+ if %5 [t: %b5] { # if_4
+ %b5 = block { # true
store %return_value, 3i
- exit_if
+ exit_if # if_4
}
-
+ }
%6:i32 = load %return_value
ret %6
}
diff --git a/src/tint/ir/unreachable.h b/src/tint/ir/unreachable.h
index 51fdc24..507d6cb 100644
--- a/src/tint/ir/unreachable.h
+++ b/src/tint/ir/unreachable.h
@@ -15,12 +15,12 @@
#ifndef SRC_TINT_IR_UNREACHABLE_H_
#define SRC_TINT_IR_UNREACHABLE_H_
-#include "src/tint/ir/branch.h"
+#include "src/tint/ir/terminator.h"
namespace tint::ir {
/// An unreachable instruction in the IR.
-class Unreachable : public utils::Castable<Unreachable, Branch> {
+class Unreachable : public utils::Castable<Unreachable, Terminator> {
public:
~Unreachable() override;
};
diff --git a/src/tint/ir/validate.cc b/src/tint/ir/validate.cc
index 7d4a15c..6687111 100644
--- a/src/tint/ir/validate.cc
+++ b/src/tint/ir/validate.cc
@@ -166,13 +166,13 @@
void CheckBlock(Block* blk) {
TINT_SCOPED_ASSIGNMENT(current_block_, blk);
- if (!blk->HasBranchTarget()) {
- AddError(blk, "block: does not end in a branch");
+ if (!blk->HasTerminator()) {
+ AddError(blk, "block: does not end in a terminator instruction");
}
for (auto* inst : *blk) {
- if (inst->Is<ir::Branch>() && inst != blk->Branch()) {
- AddError(inst, "block: branch which isn't the final instruction");
+ if (inst->Is<ir::Terminator>() && inst != blk->Terminator()) {
+ AddError(inst, "block: terminator which isn't the final instruction");
continue;
}
@@ -182,19 +182,19 @@
void CheckInstruction(Instruction* inst) {
tint::Switch(
- inst, //
- [&](Access* a) { CheckAccess(a); }, //
- [&](Binary*) {}, //
- [&](Branch* b) { CheckBranch(b); }, //
- [&](Call* c) { CheckCall(c); }, //
- [&](If* if_) { CheckIf(if_); }, //
- [&](Load*) {}, //
- [&](Loop*) {}, //
- [&](Store*) {}, //
- [&](Switch*) {}, //
- [&](Swizzle*) {}, //
- [&](Unary*) {}, //
- [&](Var*) {}, //
+ inst, //
+ [&](Access* a) { CheckAccess(a); }, //
+ [&](Binary*) {}, //
+ [&](Call* c) { CheckCall(c); }, //
+ [&](If* if_) { CheckIf(if_); }, //
+ [&](Load*) {}, //
+ [&](Loop*) {}, //
+ [&](Store*) {}, //
+ [&](Switch*) {}, //
+ [&](Swizzle*) {}, //
+ [&](Terminator* b) { CheckTerminator(b); }, //
+ [&](Unary*) {}, //
+ [&](Var*) {}, //
[&](Default) {
AddError(std::string("missing validation of: ") + inst->TypeInfo().name);
});
@@ -282,7 +282,7 @@
}
}
- void CheckBranch(ir::Branch* b) {
+ void CheckTerminator(ir::Terminator* b) {
tint::Switch(
b, //
[&](ir::BreakIf*) {}, //
@@ -298,7 +298,7 @@
},
[&](ir::Unreachable*) {}, //
[&](Default) {
- AddError(std::string("missing validation of branch: ") + b->TypeInfo().name);
+ AddError(std::string("missing validation of terminator: ") + b->TypeInfo().name);
});
}
diff --git a/src/tint/ir/validate_test.cc b/src/tint/ir/validate_test.cc
index c05674a..080625c 100644
--- a/src/tint/ir/validate_test.cc
+++ b/src/tint/ir/validate_test.cc
@@ -46,23 +46,21 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().str(), R"(:3:3 error: root block: invalid instruction: tint::ir::Loop
- loop [b: %b2]
+ EXPECT_EQ(res.Failure().str(), R"(:2:3 error: root block: invalid instruction: tint::ir::Loop
+ loop [b: %b2] { # loop_1
^^^^^^^^^^^^^
-:2:1 note: In block
-%b1 = block {
+:1:1 note: In block
+%b1 = block { # root
^^^^^^^^^^^
note: # Disassembly
-# Root block
-%b1 = block {
- loop [b: %b2]
- # Body block
- %b2 = block {
+%b1 = block { # root
+ loop [b: %b2] { # loop_1
+ %b2 = block { # body
continue %b3
}
-
+ }
}
)");
@@ -79,13 +77,13 @@
EXPECT_TRUE(res) << res.Failure().str();
}
-TEST_F(IR_ValidateTest, Block_NoBranchAtEnd) {
+TEST_F(IR_ValidateTest, Block_NoTerminator) {
auto* f = b.Function("my_func", ty.void_());
mod.functions.Push(f);
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().str(), R"(:2:3 error: block: does not end in a branch
+ EXPECT_EQ(res.Failure().str(), R"(:2:3 error: block: does not end in a terminator instruction
%b1 = block {
^^^^^^^^^^^
@@ -463,7 +461,7 @@
)");
}
-TEST_F(IR_ValidateTest, Block_BranchInMiddle) {
+TEST_F(IR_ValidateTest, Block_TerminatorInMiddle) {
auto* f = b.Function("my_func", ty.void_());
mod.functions.Push(f);
@@ -473,7 +471,8 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
- EXPECT_EQ(res.Failure().str(), R"(:3:5 error: block: branch which isn't the final instruction
+ EXPECT_EQ(res.Failure().str(),
+ R"(:3:5 error: block: terminator which isn't the final instruction
ret
^^^
@@ -505,7 +504,7 @@
auto res = ir::Validate(mod);
ASSERT_FALSE(res);
EXPECT_EQ(res.Failure().str(), R"(:3:8 error: if: condition must be a `bool` type
- if 1i [t: %b2, f: %b3]
+ if 1i [t: %b2, f: %b3] { # if_1
^^
:2:3 note: In block
@@ -515,17 +514,14 @@
note: # Disassembly
%my_func = func():void -> %b1 {
%b1 = block {
- if 1i [t: %b2, f: %b3]
- # True block
- %b2 = block {
+ if 1i [t: %b2, f: %b3] { # if_1
+ %b2 = block { # true
ret
}
-
- # False block
- %b3 = block {
+ %b3 = block { # false
ret
}
-
+ }
ret
}
}
diff --git a/src/tint/utils/hashmap_base.h b/src/tint/utils/hashmap_base.h
index 73cb0d7..d2aaf1d 100644
--- a/src/tint/utils/hashmap_base.h
+++ b/src/tint/utils/hashmap_base.h
@@ -25,9 +25,7 @@
#include "src/tint/utils/hash.h"
#include "src/tint/utils/vector.h"
-#ifndef NDEBUG
#define TINT_ASSERT_ITERATORS_NOT_INVALIDATED
-#endif
namespace tint::utils {
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
index 1a821da..fc52fe7 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -39,6 +39,7 @@
#include "src/tint/ir/return.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h"
+#include "src/tint/ir/terminator.h"
#include "src/tint/ir/transform/add_empty_entry_point.h"
#include "src/tint/ir/transform/block_decorated_structs.h"
#include "src/tint/ir/transform/merge_return.h"
@@ -479,7 +480,7 @@
[&](ir::UserCall* c) { EmitUserCall(c); }, //
[&](ir::Var* v) { EmitVar(v); }, //
[&](ir::If* i) { EmitIf(i); }, //
- [&](ir::Branch* b) { EmitBranch(b); }, //
+ [&](ir::Terminator* t) { EmitTerminator(t); }, //
[&](Default) {
TINT_ICE(Writer, diagnostics_)
<< "unimplemented instruction: " << inst->TypeInfo().name;
@@ -492,14 +493,14 @@
}
}
-void GeneratorImplIr::EmitBranch(ir::Branch* b) {
+void GeneratorImplIr::EmitTerminator(ir::Terminator* t) {
tint::Switch( //
- b, //
+ t, //
[&](ir::Return*) {
- if (!b->Args().IsEmpty()) {
- TINT_ASSERT(Writer, b->Args().Length() == 1u);
+ if (!t->Args().IsEmpty()) {
+ TINT_ASSERT(Writer, t->Args().Length() == 1u);
OperandList operands;
- operands.push_back(Value(b->Args()[0]));
+ operands.push_back(Value(t->Args()[0]));
current_function_.push_inst(spv::Op::OpReturnValue, operands);
} else {
current_function_.push_inst(spv::Op::OpReturn, {});
@@ -528,7 +529,7 @@
[&](ir::Unreachable*) { current_function_.push_inst(spv::Op::OpUnreachable, {}); },
[&](Default) {
- TINT_ICE(Writer, diagnostics_) << "unimplemented branch: " << b->TypeInfo().name;
+ TINT_ICE(Writer, diagnostics_) << "unimplemented branch: " << t->TypeInfo().name;
});
}
@@ -547,11 +548,11 @@
uint32_t true_label = merge_label;
uint32_t false_label = merge_label;
if (true_block->Length() > 1 || i->HasResults() ||
- (true_block->HasBranchTarget() && !true_block->Branch()->Is<ir::ExitIf>())) {
+ (true_block->HasTerminator() && !true_block->Terminator()->Is<ir::ExitIf>())) {
true_label = Label(true_block);
}
if (false_block->Length() > 1 || i->HasResults() ||
- (false_block->HasBranchTarget() && !false_block->Branch()->Is<ir::ExitIf>())) {
+ (false_block->HasTerminator() && !false_block->Terminator()->Is<ir::ExitIf>())) {
false_label = Label(false_block);
}
@@ -837,7 +838,7 @@
EmitBlockInstructions(loop->Body());
// Emit the loop continuing block.
- if (loop->Continuing()->HasBranchTarget()) {
+ if (loop->Continuing()->HasTerminator()) {
EmitBlock(loop->Continuing());
} else {
// We still need to emit a continuing block with a back-edge, even if it is unreachable.
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.h b/src/tint/writer/spirv/ir/generator_impl_ir.h
index 1a2ec8e..357c732 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.h
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.h
@@ -33,7 +33,6 @@
class Binary;
class Block;
class BlockParam;
-class Branch;
class BuiltinCall;
class Construct;
class ControlInstruction;
@@ -48,6 +47,7 @@
class MultiInBlock;
class Store;
class Switch;
+class Terminator;
class UserCall;
class Value;
class Var;
@@ -184,9 +184,9 @@
/// @param var the var instruction to emit
void EmitVar(ir::Var* var);
- /// Emit a branch instruction.
- /// @param b the branch instruction to emit
- void EmitBranch(ir::Branch* b);
+ /// Emit a terminator instruction.
+ /// @param term the terminator instruction to emit
+ void EmitTerminator(ir::Terminator* term);
/// Emit the OpPhis for the given flow control instruction.
/// @param inst the flow control instruction
