[ir] Add the Exit instructions.

This CL adds the ExitIf, ExitLoop and ExitSwitch instructions. The dump
to Dot has been removed as the graph is substantially different and it
needs a full re-write if we want to draw the graph.

Bug: tint:1718
Change-Id: I5ff4282abaa7542575d4f8b4b8640a3ed4d5c68f
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134464
Reviewed-by: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
This commit is contained in:
dan sinclair 2023-05-26 11:42:22 +00:00 committed by Dawn LUCI CQ
parent e982520e70
commit bdbbffbdfb
32 changed files with 481 additions and 578 deletions

View File

@ -1230,12 +1230,16 @@ if (tint_build_ir) {
"ir/continue.h", "ir/continue.h",
"ir/convert.cc", "ir/convert.cc",
"ir/convert.h", "ir/convert.h",
"ir/debug.cc",
"ir/debug.h",
"ir/disassembler.cc", "ir/disassembler.cc",
"ir/disassembler.h", "ir/disassembler.h",
"ir/discard.cc", "ir/discard.cc",
"ir/discard.h", "ir/discard.h",
"ir/exit_if.cc",
"ir/exit_if.h",
"ir/exit_loop.cc",
"ir/exit_loop.h",
"ir/exit_switch.cc",
"ir/exit_switch.h",
"ir/function.cc", "ir/function.cc",
"ir/function.h", "ir/function.h",
"ir/function_param.cc", "ir/function_param.cc",
@ -1252,8 +1256,6 @@ if (tint_build_ir) {
"ir/module.h", "ir/module.h",
"ir/return.cc", "ir/return.cc",
"ir/return.h", "ir/return.h",
"ir/root_terminator.cc",
"ir/root_terminator.h",
"ir/store.cc", "ir/store.cc",
"ir/store.h", "ir/store.h",
"ir/switch.cc", "ir/switch.cc",

View File

@ -738,12 +738,16 @@ if(${TINT_BUILD_IR})
ir/continue.h ir/continue.h
ir/convert.cc ir/convert.cc
ir/convert.h ir/convert.h
ir/debug.cc
ir/debug.h
ir/disassembler.cc ir/disassembler.cc
ir/disassembler.h ir/disassembler.h
ir/discard.cc ir/discard.cc
ir/discard.h ir/discard.h
ir/exit_if.cc
ir/exit_if.h
ir/exit_loop.cc
ir/exit_loop.h
ir/exit_switch.cc
ir/exit_switch.h
ir/from_program.cc ir/from_program.cc
ir/from_program.h ir/from_program.h
ir/function.cc ir/function.cc
@ -762,8 +766,6 @@ if(${TINT_BUILD_IR})
ir/module.h ir/module.h
ir/return.cc ir/return.cc
ir/return.h ir/return.h
ir/root_terminator.cc
ir/root_terminator.h
ir/store.cc ir/store.cc
ir/store.h ir/store.h
ir/switch.cc ir/switch.cc

View File

@ -49,7 +49,6 @@
#include "tint/tint.h" #include "tint/tint.h"
#if TINT_BUILD_IR #if TINT_BUILD_IR
#include "src/tint/ir/debug.h" // nogncheck
#include "src/tint/ir/disassembler.h" // nogncheck #include "src/tint/ir/disassembler.h" // nogncheck
#include "src/tint/ir/from_program.h" // nogncheck #include "src/tint/ir/from_program.h" // nogncheck
#include "src/tint/ir/module.h" // nogncheck #include "src/tint/ir/module.h" // nogncheck
@ -110,7 +109,6 @@ struct Options {
#if TINT_BUILD_IR #if TINT_BUILD_IR
bool dump_ir = false; bool dump_ir = false;
bool dump_ir_graph = false;
bool use_ir = false; bool use_ir = false;
#endif // TINT_BUILD_IR #endif // TINT_BUILD_IR
@ -374,8 +372,6 @@ bool ParseArgs(const std::vector<std::string>& args, Options* opts) {
#if TINT_BUILD_IR #if TINT_BUILD_IR
} else if (arg == "--dump-ir") { } else if (arg == "--dump-ir") {
opts->dump_ir = true; opts->dump_ir = true;
} else if (arg == "--dump-ir-graph") {
opts->dump_ir_graph = true;
} else if (arg == "--use-ir") { } else if (arg == "--use-ir") {
opts->use_ir = true; opts->use_ir = true;
#endif // TINT_BUILD_IR #endif // TINT_BUILD_IR
@ -1072,7 +1068,7 @@ int main(int argc, const char** argv) {
#endif // TINT_BUILD_SYNTAX_TREE_WRITER #endif // TINT_BUILD_SYNTAX_TREE_WRITER
#if TINT_BUILD_IR #if TINT_BUILD_IR
if (options.dump_ir || options.dump_ir_graph) { if (options.dump_ir) {
auto result = tint::ir::FromProgram(program.get()); auto result = tint::ir::FromProgram(program.get());
if (!result) { if (!result) {
std::cerr << "Failed to build IR from program: " << result.Failure() << std::endl; std::cerr << "Failed to build IR from program: " << result.Failure() << std::endl;
@ -1082,10 +1078,6 @@ int main(int argc, const char** argv) {
tint::ir::Disassembler d(mod); tint::ir::Disassembler d(mod);
std::cout << d.Disassemble() << std::endl; std::cout << d.Disassemble() << std::endl;
} }
if (options.dump_ir_graph) {
auto graph = tint::ir::Debug::AsDotGraph(&mod);
WriteFile("tint.dot", "w", graph);
}
} }
} }
#endif // TINT_BUILD_IR #endif // TINT_BUILD_IR

View File

@ -46,18 +46,6 @@ class Block : public utils::Castable<Block> {
return instructions_.Back()->As<ir::Branch>(); return instructions_.Back()->As<ir::Branch>();
} }
/// @param target the block to see if we trampoline too
/// @returns if this block just branches to the provided target.
bool IsTrampoline(const Block* target) const {
if (instructions_.Length() != 1) {
return false;
}
if (auto* inst = instructions_.Front()->As<ir::Branch>()) {
return inst->To() == target;
}
return false;
}
/// Sets the instructions in the block /// Sets the instructions in the block
/// @param instructions the instructions to set /// @param instructions the instructions to set
void SetInstructions(utils::VectorRef<const Instruction*> instructions) { void SetInstructions(utils::VectorRef<const Instruction*> instructions) {

View File

@ -28,13 +28,6 @@ Branch::Branch(utils::VectorRef<Value*> args) : args_(std::move(args)) {
} }
} }
Branch::Branch(Block* to, utils::VectorRef<Value*> args) : Branch(args) {
to_ = to;
TINT_ASSERT(IR, to_);
to_->AddInboundBranch(this);
}
Branch::~Branch() = default; Branch::~Branch() = default;
} // namespace tint::ir } // namespace tint::ir

View File

@ -29,15 +29,8 @@ namespace tint::ir {
/// A branch instruction. /// A branch instruction.
class Branch : public utils::Castable<Branch, Instruction> { class Branch : public utils::Castable<Branch, Instruction> {
public: public:
/// Constructor
/// @param to the block to branch too
/// @param args the branch arguments
explicit Branch(Block* to, utils::VectorRef<Value*> args = {});
~Branch() override; ~Branch() override;
/// @returns the block being branched too.
const Block* To() const { return to_; }
/// @returns the branch arguments /// @returns the branch arguments
utils::VectorRef<Value*> Args() const { return args_; } utils::VectorRef<Value*> Args() const { return args_; }
@ -47,7 +40,6 @@ class Branch : public utils::Castable<Branch, Instruction> {
explicit Branch(utils::VectorRef<Value*> args); explicit Branch(utils::VectorRef<Value*> args);
private: private:
Block* to_ = nullptr;
utils::Vector<Value*, 2> args_; utils::Vector<Value*, 2> args_;
}; };

View File

@ -37,10 +37,6 @@ Block* Builder::CreateBlock() {
return ir.blocks.Create<Block>(); return ir.blocks.Create<Block>();
} }
RootTerminator* Builder::CreateRootTerminator() {
return ir.blocks.Create<RootTerminator>();
}
Function* Builder::CreateFunction(std::string_view name, Function* Builder::CreateFunction(std::string_view name,
const type::Type* return_type, const type::Type* return_type,
Function::PipelineStage stage, Function::PipelineStage stage,
@ -205,10 +201,6 @@ ir::Var* Builder::Declare(const type::Type* type) {
return ir.values.Create<ir::Var>(type); return ir.values.Create<ir::Var>(type);
} }
ir::Branch* Builder::Branch(Block* to, utils::VectorRef<Value*> args) {
return ir.values.Create<ir::Branch>(to, args);
}
ir::Return* Builder::Return(Function* func, utils::VectorRef<Value*> args) { ir::Return* Builder::Return(Function* func, utils::VectorRef<Value*> args) {
return ir.values.Create<ir::Return>(func, args); return ir.values.Create<ir::Return>(func, args);
} }
@ -220,6 +212,17 @@ ir::BreakIf* Builder::BreakIf(Value* condition, Loop* loop) {
ir::Continue* Builder::Continue(Loop* loop) { ir::Continue* Builder::Continue(Loop* loop) {
return ir.values.Create<ir::Continue>(loop); return ir.values.Create<ir::Continue>(loop);
} }
ir::ExitSwitch* Builder::ExitSwitch(Switch* sw) {
return ir.values.Create<ir::ExitSwitch>(sw);
}
ir::ExitLoop* Builder::ExitLoop(Loop* loop) {
return ir.values.Create<ir::ExitLoop>(loop);
}
ir::ExitIf* Builder::ExitIf(If* i, utils::VectorRef<Value*> args) {
return ir.values.Create<ir::ExitIf>(i, args);
}
ir::BlockParam* Builder::BlockParam(const type::Type* type) { ir::BlockParam* Builder::BlockParam(const type::Type* type) {
return ir.values.Create<ir::BlockParam>(type); return ir.values.Create<ir::BlockParam>(type);

View File

@ -28,6 +28,9 @@
#include "src/tint/ir/continue.h" #include "src/tint/ir/continue.h"
#include "src/tint/ir/convert.h" #include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h" #include "src/tint/ir/discard.h"
#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/exit_loop.h"
#include "src/tint/ir/exit_switch.h"
#include "src/tint/ir/function.h" #include "src/tint/ir/function.h"
#include "src/tint/ir/function_param.h" #include "src/tint/ir/function_param.h"
#include "src/tint/ir/if.h" #include "src/tint/ir/if.h"
@ -35,7 +38,6 @@
#include "src/tint/ir/loop.h" #include "src/tint/ir/loop.h"
#include "src/tint/ir/module.h" #include "src/tint/ir/module.h"
#include "src/tint/ir/return.h" #include "src/tint/ir/return.h"
#include "src/tint/ir/root_terminator.h"
#include "src/tint/ir/store.h" #include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h" #include "src/tint/ir/switch.h"
#include "src/tint/ir/unary.h" #include "src/tint/ir/unary.h"
@ -64,9 +66,6 @@ class Builder {
/// @returns a new block flow node /// @returns a new block flow node
Block* CreateBlock(); Block* CreateBlock();
/// @returns a new root terminator flow node
RootTerminator* CreateRootTerminator();
/// Creates a function flow node /// Creates a function flow node
/// @param name the function name /// @param name the function name
/// @param return_type the function return type /// @param return_type the function return type
@ -349,11 +348,21 @@ class Builder {
/// @returns the instruction /// @returns the instruction
ir::Continue* Continue(Loop* loop); ir::Continue* Continue(Loop* loop);
/// Creates a branch declaration /// Creates an exit switch instruction
/// @param to the node being branched too /// @param sw the switch being exited
/// @returns the instruction
ir::ExitSwitch* ExitSwitch(Switch* sw);
/// Creates an exit loop instruction
/// @param loop the loop being exited
/// @returns the instruction
ir::ExitLoop* ExitLoop(Loop* loop);
/// Creates an exit if instruction
/// @param i the if being exited
/// @param args the branch arguments /// @param args the branch arguments
/// @returns the instruction /// @returns the instruction
ir::Branch* Branch(Block* to, utils::VectorRef<Value*> args = {}); ir::ExitIf* ExitIf(If* i, utils::VectorRef<Value*> args = {});
/// Creates a new `BlockParam` /// Creates a new `BlockParam`
/// @param type the parameter type /// @param type the parameter type

View File

@ -1,95 +0,0 @@
// Copyright 2022 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#include "src/tint/ir/debug.h"
#include <unordered_map>
#include <unordered_set>
#include "src/tint/ir/block.h"
#include "src/tint/ir/continue.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/loop.h"
#include "src/tint/ir/return.h"
#include "src/tint/ir/switch.h"
#include "src/tint/switch.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
// static
std::string Debug::AsDotGraph(const Module* mod) {
size_t block_count = 0;
std::unordered_set<const Block*> visited;
std::unordered_set<const Block*> merge_blocks;
std::unordered_map<const Block*, std::string> block_to_name;
utils::StringStream out;
auto name_for = [&](const Block* blk) -> std::string {
if (block_to_name.count(blk) > 0) {
return block_to_name[blk];
}
std::string name = "blk_" + std::to_string(block_count);
block_count += 1;
block_to_name[blk] = name;
return name;
};
std::function<void(const Block*)> Graph = [&](const Block* blk) {
if (visited.count(blk) > 0) {
return;
}
visited.insert(blk);
tint::Switch(blk, //
[&](const ir::Block* b) {
if (block_to_name.count(b) == 0) {
out << name_for(b) << R"( [label="block"])" << std::endl;
}
out << name_for(b) << " -> " << name_for(b->Branch()->To());
// Dashed lines to merge blocks
if (merge_blocks.count(b->Branch()->To()) != 0) {
out << " [style=dashed]";
}
out << std::endl;
if (b->Branch()->Is<ir::Return>()) {
return;
} else if (auto* cont = b->Branch()->As<ir::Continue>()) {
Graph(cont->Loop()->Continuing());
} else {
Graph(b->Branch()->To());
}
});
};
out << "digraph G {" << std::endl;
for (const auto* func : mod->functions) {
// Cluster each function to label and draw a box around it.
out << "subgraph cluster_" << mod->NameOf(func).Name() << " {" << std::endl;
out << R"(label=")" << mod->NameOf(func).Name() << R"(")" << std::endl;
out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl;
Graph(func->StartTarget());
out << "}" << std::endl;
}
out << "}";
return out.str();
}
} // namespace tint::ir

View File

@ -27,11 +27,13 @@
#include "src/tint/ir/continue.h" #include "src/tint/ir/continue.h"
#include "src/tint/ir/convert.h" #include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h" #include "src/tint/ir/discard.h"
#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/exit_loop.h"
#include "src/tint/ir/exit_switch.h"
#include "src/tint/ir/if.h" #include "src/tint/ir/if.h"
#include "src/tint/ir/load.h" #include "src/tint/ir/load.h"
#include "src/tint/ir/loop.h" #include "src/tint/ir/loop.h"
#include "src/tint/ir/return.h" #include "src/tint/ir/return.h"
#include "src/tint/ir/root_terminator.h"
#include "src/tint/ir/store.h" #include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h" #include "src/tint/ir/switch.h"
#include "src/tint/ir/user_call.h" #include "src/tint/ir/user_call.h"
@ -70,7 +72,6 @@ void Disassembler::EmitBlockInstructions(const Block* b) {
for (const auto* inst : b->Instructions()) { for (const auto* inst : b->Instructions()) {
Indent(); Indent();
EmitInstruction(inst); EmitInstruction(inst);
out_ << std::endl;
} }
} }
@ -92,8 +93,8 @@ std::string_view Disassembler::IdOf(const Value* value) {
std::string Disassembler::Disassemble() { std::string Disassembler::Disassemble() {
if (mod_.root_block) { if (mod_.root_block) {
Indent() << "# Root block" << std::endl; Indent() << "# Root block" << std::endl;
Walk(mod_.root_block); WalkInternal(mod_.root_block);
Walk(mod_.root_block->Branch()->To()); out_ << std::endl;
} }
for (auto* func : mod_.functions) { for (auto* func : mod_.functions) {
@ -108,41 +109,33 @@ void Disassembler::Walk(const Block* blk) {
} }
visited_.Add(blk); visited_.Add(blk);
tint::Switch( // If this block is dead, nothing to do
blk, if (!blk->HasBranchTarget()) {
[&](const ir::RootTerminator* t) { return;
TINT_ASSERT(IR, !in_function_); }
Indent() << "%b" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
},
[&](const ir::Block* b) {
// If this block is dead, nothing to do
if (!b->HasBranchTarget()) {
return;
}
Indent() << "%b" << IdOf(b) << " = block"; WalkInternal(blk);
if (!b->Params().IsEmpty()) { }
out_ << " (";
for (auto* p : b->Params()) {
if (p != b->Params().Front()) {
out_ << ", ";
}
EmitValue(p);
}
out_ << ")";
}
out_ << " {" << std::endl; void Disassembler::WalkInternal(const Block* blk) {
{ Indent() << "%b" << IdOf(blk) << " = block";
ScopedIndent si(indent_size_); if (!blk->Params().IsEmpty()) {
EmitBlockInstructions(b); out_ << " (";
for (auto* p : blk->Params()) {
if (p != blk->Params().Front()) {
out_ << ", ";
} }
Indent() << "}" << std::endl; EmitValue(p);
}
out_ << ")";
}
if (!b->Branch()->Is<ir::Return>()) { out_ << " {" << std::endl;
out_ << std::endl; {
} ScopedIndent si(indent_size_);
}); EmitBlockInstructions(blk);
}
Indent() << "}" << std::endl;
} }
void Disassembler::EmitFunction(const Function* func) { void Disassembler::EmitFunction(const Function* func) {
@ -258,33 +251,39 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
EmitValueWithType(b); EmitValueWithType(b);
out_ << " = bitcast "; out_ << " = bitcast ";
EmitArgs(b); EmitArgs(b);
out_ << std::endl;
}, },
[&](const ir::Discard*) { out_ << "discard"; }, [&](const ir::Discard*) { out_ << "discard" << std::endl; },
[&](const ir::Builtin* b) { [&](const ir::Builtin* b) {
EmitValueWithType(b); EmitValueWithType(b);
out_ << " = " << builtin::str(b->Func()) << " "; out_ << " = " << builtin::str(b->Func()) << " ";
EmitArgs(b); EmitArgs(b);
out_ << std::endl;
}, },
[&](const ir::Construct* c) { [&](const ir::Construct* c) {
EmitValueWithType(c); EmitValueWithType(c);
out_ << " = construct "; out_ << " = construct ";
EmitArgs(c); EmitArgs(c);
out_ << std::endl;
}, },
[&](const ir::Convert* c) { [&](const ir::Convert* c) {
EmitValueWithType(c); EmitValueWithType(c);
out_ << " = convert " << c->FromType()->FriendlyName() << ", "; out_ << " = convert " << c->FromType()->FriendlyName() << ", ";
EmitArgs(c); EmitArgs(c);
out_ << std::endl;
}, },
[&](const ir::Load* l) { [&](const ir::Load* l) {
EmitValueWithType(l); EmitValueWithType(l);
out_ << " = load "; out_ << " = load ";
EmitValue(l->From()); EmitValue(l->From());
out_ << std::endl;
}, },
[&](const ir::Store* s) { [&](const ir::Store* s) {
out_ << "store "; out_ << "store ";
EmitValue(s->To()); EmitValue(s->To());
out_ << ", "; out_ << ", ";
EmitValue(s->From()); EmitValue(s->From());
out_ << std::endl;
}, },
[&](const ir::UserCall* uc) { [&](const ir::UserCall* uc) {
EmitValueWithType(uc); EmitValueWithType(uc);
@ -293,6 +292,7 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
out_ << ", "; out_ << ", ";
} }
EmitArgs(uc); EmitArgs(uc);
out_ << std::endl;
}, },
[&](const ir::Var* v) { [&](const ir::Var* v) {
EmitValueWithType(v); EmitValueWithType(v);
@ -301,6 +301,7 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
out_ << ", "; out_ << ", ";
EmitValue(v->Initializer()); EmitValue(v->Initializer());
} }
out_ << std::endl;
}, },
[&](const ir::Branch* b) { EmitBranch(b); }, [&](const ir::Branch* b) { EmitBranch(b); },
[&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; }); [&](Default) { out_ << "Unknown instruction: " << inst->TypeInfo().name; });
@ -332,15 +333,18 @@ void Disassembler::EmitIf(const If* i) {
ScopedIndent si(indent_size_); ScopedIndent si(indent_size_);
Indent() << "# True block" << std::endl; Indent() << "# True block" << std::endl;
Walk(i->True()); Walk(i->True());
out_ << std::endl;
} }
if (has_false) { if (has_false) {
ScopedIndent si(indent_size_); ScopedIndent si(indent_size_);
Indent() << "# False block" << std::endl; Indent() << "# False block" << std::endl;
Walk(i->False()); Walk(i->False());
out_ << std::endl;
} }
if (i->Merge()->HasBranchTarget()) { if (i->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl; Indent() << "# Merge block" << std::endl;
Walk(i->Merge()); Walk(i->Merge());
out_ << std::endl;
} }
} }
@ -358,16 +362,19 @@ void Disassembler::EmitLoop(const Loop* l) {
{ {
ScopedIndent si(indent_size_); ScopedIndent si(indent_size_);
Walk(l->Start()); Walk(l->Start());
out_ << std::endl;
} }
if (l->Continuing()->HasBranchTarget()) { if (l->Continuing()->HasBranchTarget()) {
ScopedIndent si(indent_size_); ScopedIndent si(indent_size_);
Indent() << "# Continuing block" << std::endl; Indent() << "# Continuing block" << std::endl;
Walk(l->Continuing()); Walk(l->Continuing());
out_ << std::endl;
} }
if (l->Merge()->HasBranchTarget()) { if (l->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl; Indent() << "# Merge block" << std::endl;
Walk(l->Merge()); Walk(l->Merge());
out_ << std::endl;
} }
} }
@ -402,29 +409,31 @@ void Disassembler::EmitSwitch(const Switch* s) {
ScopedIndent si(indent_size_); ScopedIndent si(indent_size_);
Indent() << "# Case block" << std::endl; Indent() << "# Case block" << std::endl;
Walk(c.Start()); Walk(c.Start());
out_ << std::endl;
} }
if (s->Merge()->HasBranchTarget()) { if (s->Merge()->HasBranchTarget()) {
Indent() << "# Merge block" << std::endl; Indent() << "# Merge block" << std::endl;
Walk(s->Merge()); Walk(s->Merge());
out_ << std::endl;
} }
} }
void Disassembler::EmitBranch(const Branch* b) { void Disassembler::EmitBranch(const Branch* b) {
std::string suffix = ""; tint::Switch(
if (b->Is<ir::Return>()) { b, //
out_ << "ret"; [&](const ir::Return*) { out_ << "ret"; },
} else if (auto* cont = b->As<ir::Continue>()) { [&](const ir::Continue* cont) {
out_ << "continue %b" << IdOf(cont->Loop()->Continuing()); out_ << "continue %b" << IdOf(cont->Loop()->Continuing());
} else if (auto* bi = b->As<ir::BreakIf>()) { },
out_ << "break_if "; [&](const ir::ExitIf* ei) { out_ << "exit_if %b" << IdOf(ei->If()->Merge()); },
EmitValue(bi->Condition()); [&](const ir::ExitSwitch* es) { out_ << "exit_switch %b" << IdOf(es->Switch()->Merge()); },
out_ << " %b" << IdOf(bi->Loop()->Start()); [&](const ir::ExitLoop* el) { out_ << "exit_loop %b" << IdOf(el->Loop()->Merge()); },
} else { [&](const ir::BreakIf* bi) {
out_ << "br %b" << IdOf(b->To()); out_ << "break_if ";
if (b->To()->Is<RootTerminator>()) { EmitValue(bi->Condition());
suffix = "root_end"; out_ << " %b" << IdOf(bi->Loop()->Start());
} },
} [&](Default) { out_ << "Unknown branch " << b->TypeInfo().name; });
if (!b->Args().IsEmpty()) { if (!b->Args().IsEmpty()) {
out_ << " "; out_ << " ";
@ -435,9 +444,7 @@ void Disassembler::EmitBranch(const Branch* b) {
EmitValue(v); EmitValue(v);
} }
} }
if (!suffix.empty()) { out_ << std::endl;
out_ << " # " << suffix;
}
} }
void Disassembler::EmitArgs(const Call* call) { void Disassembler::EmitArgs(const Call* call) {
@ -508,6 +515,7 @@ void Disassembler::EmitBinary(const Binary* b) {
EmitValue(b->LHS()); EmitValue(b->LHS());
out_ << ", "; out_ << ", ";
EmitValue(b->RHS()); EmitValue(b->RHS());
out_ << std::endl;
} }
void Disassembler::EmitUnary(const Unary* u) { void Disassembler::EmitUnary(const Unary* u) {
@ -523,6 +531,7 @@ void Disassembler::EmitUnary(const Unary* u) {
} }
out_ << " "; out_ << " ";
EmitValue(u->Val()); EmitValue(u->Val());
out_ << std::endl;
} }
} // namespace tint::ir } // namespace tint::ir

View File

@ -57,6 +57,7 @@ class Disassembler {
std::string_view IdOf(const Value* node); std::string_view IdOf(const Value* node);
void Walk(const Block* blk); void Walk(const Block* blk);
void WalkInternal(const Block* blk);
void EmitFunction(const Function* func); void EmitFunction(const Function* func);
void EmitInstruction(const Instruction* inst); void EmitInstruction(const Instruction* inst);
void EmitValueWithType(const Value* val); void EmitValueWithType(const Value* val);

View File

@ -1,4 +1,4 @@
// Copyright 2022 The Tint Authors. // Copyright 2023 The Tint Authors.
// //
// Licensed under the Apache License, Version 2.0 (the "License"); // Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License. // you may not use this file except in compliance with the License.
@ -12,14 +12,20 @@
// See the License for the specific language governing permissions and // See the License for the specific language governing permissions and
// limitations under the License. // limitations under the License.
#include "src/tint/ir/root_terminator.h" #include "src/tint/ir/exit_if.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::RootTerminator); #include "src/tint/ir/if.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitIf);
namespace tint::ir { namespace tint::ir {
RootTerminator::RootTerminator() : Base() {} ExitIf::ExitIf(ir::If* i, utils::VectorRef<Value*> args) : Base(args), if_(i) {
TINT_ASSERT(IR, if_);
if_->AddUsage(this);
if_->Merge()->AddInboundBranch(this);
}
RootTerminator::~RootTerminator() = default; ExitIf::~ExitIf() = default;
} // namespace tint::ir } // namespace tint::ir

46
src/tint/ir/exit_if.h Normal file
View File

@ -0,0 +1,46 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef SRC_TINT_IR_EXIT_IF_H_
#define SRC_TINT_IR_EXIT_IF_H_
#include "src/tint/ir/branch.h"
#include "src/tint/utils/castable.h"
// Forward declarations
namespace tint::ir {
class If;
} // namespace tint::ir
namespace tint::ir {
/// A exit if instruction.
class ExitIf : public utils::Castable<ExitIf, Branch> {
public:
/// Constructor
/// @param i the if being exited
/// @param args the branch arguments
explicit ExitIf(ir::If* i, utils::VectorRef<Value*> args = {});
~ExitIf() override;
/// @returns the if being exited
const ir::If* If() const { return if_; }
private:
ir::If* if_ = nullptr;
};
} // namespace tint::ir
#endif // SRC_TINT_IR_EXIT_IF_H_

View File

@ -1,4 +1,4 @@
// Copyright 2022 The Tint Authors. // Copyright 2023 The Tint Authors.
// //
// Licensed under the Apache License, Version 2.0 (the "License"); // Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License. // you may not use this file except in compliance with the License.
@ -12,21 +12,20 @@
// See the License for the specific language governing permissions and // See the License for the specific language governing permissions and
// limitations under the License. // limitations under the License.
#ifndef SRC_TINT_IR_ROOT_TERMINATOR_H_ #include "src/tint/ir/exit_loop.h"
#define SRC_TINT_IR_ROOT_TERMINATOR_H_
#include "src/tint/ir/block.h" #include "src/tint/ir/loop.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitLoop);
namespace tint::ir { namespace tint::ir {
/// Block used as the end of a root block. There are no instructions in this block. ExitLoop::ExitLoop(ir::Loop* loop) : Base(utils::Empty), loop_(loop) {
class RootTerminator : public utils::Castable<RootTerminator, Block> { TINT_ASSERT(IR, loop_);
public: loop_->AddUsage(this);
/// Constructor loop_->Merge()->AddInboundBranch(this);
RootTerminator(); }
~RootTerminator() override;
}; ExitLoop::~ExitLoop() = default;
} // namespace tint::ir } // namespace tint::ir
#endif // SRC_TINT_IR_ROOT_TERMINATOR_H_

45
src/tint/ir/exit_loop.h Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef SRC_TINT_IR_EXIT_LOOP_H_
#define SRC_TINT_IR_EXIT_LOOP_H_
#include "src/tint/ir/branch.h"
#include "src/tint/utils/castable.h"
// Forward declarations
namespace tint::ir {
class Loop;
} // namespace tint::ir
namespace tint::ir {
/// A exit loop instruction.
class ExitLoop : public utils::Castable<ExitLoop, Branch> {
public:
/// Constructor
/// @param loop the loop being exited
explicit ExitLoop(ir::Loop* loop);
~ExitLoop() override;
/// @returns the loop being exited
const ir::Loop* Loop() const { return loop_; }
private:
ir::Loop* loop_ = nullptr;
};
} // namespace tint::ir
#endif // SRC_TINT_IR_EXIT_LOOP_H_

View File

@ -1,4 +1,4 @@
// Copyright 2022 The Tint Authors. // Copyright 2023 The Tint Authors.
// //
// Licensed under the Apache License, Version 2.0 (the "License"); // Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License. // you may not use this file except in compliance with the License.
@ -12,29 +12,20 @@
// See the License for the specific language governing permissions and // See the License for the specific language governing permissions and
// limitations under the License. // limitations under the License.
#ifndef SRC_TINT_IR_DEBUG_H_ #include "src/tint/ir/exit_switch.h"
#define SRC_TINT_IR_DEBUG_H_
#include <string> #include "src/tint/ir/switch.h"
#include "src/tint/ir/module.h" TINT_INSTANTIATE_TYPEINFO(tint::ir::ExitSwitch);
namespace tint::ir { namespace tint::ir {
/// Helper class to debug IR. ExitSwitch::ExitSwitch(ir::Switch* sw) : Base(utils::Empty), switch_(sw) {
class Debug { TINT_ASSERT(IR, switch_);
public: switch_->AddUsage(this);
/// Returns the module as a dot graph switch_->Merge()->AddInboundBranch(this);
/// @param mod the module to emit }
/// @returns the dot graph for the given module
static std::string AsDotGraph(const Module* mod);
/// Returns the module as a string ExitSwitch::~ExitSwitch() = default;
/// @param mod the module to emit
/// @returns the string representation of the module
static std::string AsString(const Module* mod);
};
} // namespace tint::ir } // namespace tint::ir
#endif // SRC_TINT_IR_DEBUG_H_

45
src/tint/ir/exit_switch.h Normal file
View File

@ -0,0 +1,45 @@
// Copyright 2023 The Tint Authors.
//
// Licensed under the Apache License, Version 2.0 (the "License");
// you may not use this file except in compliance with the License.
// You may obtain a copy of the License at
//
// http://www.apache.org/licenses/LICENSE-2.0
//
// Unless required by applicable law or agreed to in writing, software
// distributed under the License is distributed on an "AS IS" BASIS,
// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
// See the License for the specific language governing permissions and
// limitations under the License.
#ifndef SRC_TINT_IR_EXIT_SWITCH_H_
#define SRC_TINT_IR_EXIT_SWITCH_H_
#include "src/tint/ir/branch.h"
#include "src/tint/utils/castable.h"
// Forward declarations
namespace tint::ir {
class Switch;
} // namespace tint::ir
namespace tint::ir {
/// A exit switch instruction.
class ExitSwitch : public utils::Castable<ExitSwitch, Branch> {
public:
/// Constructor
/// @param sw the switch being exited
explicit ExitSwitch(ir::Switch* sw);
~ExitSwitch() override;
/// @returns the switch being exited
const ir::Switch* Switch() const { return switch_; }
private:
ir::Switch* switch_ = nullptr;
};
} // namespace tint::ir
#endif // SRC_TINT_IR_EXIT_SWITCH_H_

View File

@ -62,6 +62,9 @@
#include "src/tint/ast/while_statement.h" #include "src/tint/ast/while_statement.h"
#include "src/tint/ir/block_param.h" #include "src/tint/ir/block_param.h"
#include "src/tint/ir/builder.h" #include "src/tint/ir/builder.h"
#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/exit_loop.h"
#include "src/tint/ir/exit_switch.h"
#include "src/tint/ir/function.h" #include "src/tint/ir/function.h"
#include "src/tint/ir/if.h" #include "src/tint/ir/if.h"
#include "src/tint/ir/loop.h" #include "src/tint/ir/loop.h"
@ -98,11 +101,8 @@ namespace {
using ResultType = utils::Result<Module, diag::List>; using ResultType = utils::Result<Module, diag::List>;
// For an `if` and `switch` block, the merge has a registered incoming branch instruction of the bool IsConnected(const Block* b) {
// `if` and `switch. So, to determine if the merge is connected to any of the branches that happend return b->InboundBranches().Length() > 0;
// in the `if` or `switch` we need a `count` value that is larger then 1.
bool IsConnected(const Block* b, uint32_t count) {
return b->InboundBranches().Length() > count;
} }
/// Impl is the private-implementation of FromProgram(). /// Impl is the private-implementation of FromProgram().
@ -176,21 +176,6 @@ class Impl {
current_flow_block_ = nullptr; current_flow_block_ = nullptr;
} }
void BranchTo(Block* node, utils::VectorRef<Value*> args = {}) {
TINT_ASSERT(IR, current_flow_block_);
TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
current_flow_block_->Instructions().Push(builder_.Branch(node, args));
current_flow_block_ = nullptr;
}
void BranchToIfNeeded(Block* node) {
if (!NeedBranch()) {
return;
}
BranchTo(node);
}
Branch* FindEnclosingControl(ControlFlags flags) { Branch* FindEnclosingControl(ControlFlags flags) {
for (auto it = control_stack_.rbegin(); it != control_stack_.rend(); ++it) { for (auto it = control_stack_.rbegin(); it != control_stack_.rend(); ++it) {
if ((*it)->Is<Loop>()) { if ((*it)->Is<Loop>()) {
@ -238,11 +223,6 @@ class Impl {
}); });
} }
// Add the root terminator if needed
if (mod.root_block) {
mod.root_block->Instructions().Push(builder_.Branch(builder_.CreateRootTerminator()));
}
if (diagnostics_.contains_errors()) { if (diagnostics_.contains_errors()) {
return ResultType(std::move(diagnostics_)); return ResultType(std::move(diagnostics_));
} }
@ -541,7 +521,9 @@ class Impl {
EmitBlock(stmt->body); EmitBlock(stmt->body);
// If the true branch did not execute control flow, then go to the Merge().target // If the true branch did not execute control flow, then go to the Merge().target
BranchToIfNeeded(if_inst->Merge()); if (NeedBranch()) {
SetBranch(builder_.ExitIf(if_inst));
}
current_flow_block_ = if_inst->False(); current_flow_block_ = if_inst->False();
if (stmt->else_statement) { if (stmt->else_statement) {
@ -549,14 +531,16 @@ class Impl {
} }
// If the false branch did not execute control flow, then go to the Merge().target // If the false branch did not execute control flow, then go to the Merge().target
BranchToIfNeeded(if_inst->Merge()); if (NeedBranch()) {
SetBranch(builder_.ExitIf(if_inst));
}
} }
current_flow_block_ = nullptr; current_flow_block_ = nullptr;
// If both branches went somewhere, then they both returned, continued or broke. So, // If both branches went somewhere, then they both returned, continued or broke. So,
// there is no need for the if merge-block and there is nothing to branch to the merge // there is no need for the if merge-block and there is nothing to branch to the merge
// block anyway. // block anyway.
if (IsConnected(if_inst->Merge(), 1)) { if (IsConnected(if_inst->Merge())) {
current_flow_block_ = if_inst->Merge(); current_flow_block_ = if_inst->Merge();
} }
} }
@ -580,7 +564,7 @@ class Impl {
SetBranch(builder_.Continue(loop_inst)); SetBranch(builder_.Continue(loop_inst));
} }
if (IsConnected(loop_inst->Continuing(), 0)) { if (IsConnected(loop_inst->Continuing())) {
// Note, even if there is no continuing block, we may have branched into the // Note, even if there is no continuing block, we may have branched into the
// continue so we have to set the current block and then emit the branch if needed // continue so we have to set the current block and then emit the branch if needed
// below otherwise empty continuing blocks will fail to branch back to the start // below otherwise empty continuing blocks will fail to branch back to the start
@ -600,7 +584,7 @@ class Impl {
// target branches, eventually, to the merge, but nothing branched to the // target branches, eventually, to the merge, but nothing branched to the
// Continuing() block. // Continuing() block.
current_flow_block_ = loop_inst->Merge(); current_flow_block_ = loop_inst->Merge();
if (!IsConnected(loop_inst->Merge(), 0)) { if (!IsConnected(loop_inst->Merge())) {
current_flow_block_ = nullptr; current_flow_block_ = nullptr;
} }
} }
@ -626,10 +610,14 @@ class Impl {
// Create an `if (cond) {} else {break;}` control flow // Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get()); auto* if_inst = builder_.CreateIf(reg.Get());
if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
current_flow_block_->Instructions().Push(if_inst); current_flow_block_->Instructions().Push(if_inst);
current_flow_block_ = if_inst->True();
SetBranch(builder_.ExitIf(if_inst));
current_flow_block_ = if_inst->False();
SetBranch(builder_.ExitLoop(loop_inst));
current_flow_block_ = if_inst->Merge(); current_flow_block_ = if_inst->Merge();
EmitBlock(stmt->body); EmitBlock(stmt->body);
@ -669,10 +657,14 @@ class Impl {
// Create an `if (cond) {} else {break;}` control flow // Create an `if (cond) {} else {break;}` control flow
auto* if_inst = builder_.CreateIf(reg.Get()); auto* if_inst = builder_.CreateIf(reg.Get());
if_inst->True()->Instructions().Push(builder_.Branch(if_inst->Merge()));
if_inst->False()->Instructions().Push(builder_.Branch(loop_inst->Merge()));
current_flow_block_->Instructions().Push(if_inst); current_flow_block_->Instructions().Push(if_inst);
current_flow_block_ = if_inst->True();
SetBranch(builder_.ExitIf(if_inst));
current_flow_block_ = if_inst->False();
SetBranch(builder_.ExitLoop(loop_inst));
current_flow_block_ = if_inst->Merge(); current_flow_block_ = if_inst->Merge();
} }
@ -719,12 +711,14 @@ class Impl {
current_flow_block_ = builder_.CreateCase(switch_inst, selectors); current_flow_block_ = builder_.CreateCase(switch_inst, selectors);
EmitBlock(c->Body()->Declaration()); EmitBlock(c->Body()->Declaration());
BranchToIfNeeded(switch_inst->Merge()); if (NeedBranch()) {
SetBranch(builder_.ExitSwitch(switch_inst));
}
} }
} }
current_flow_block_ = nullptr; current_flow_block_ = nullptr;
if (IsConnected(switch_inst->Merge(), 1)) { if (IsConnected(switch_inst->Merge())) {
current_flow_block_ = switch_inst->Merge(); current_flow_block_ = switch_inst->Merge();
} }
} }
@ -746,9 +740,9 @@ class Impl {
TINT_ASSERT(IR, current_control); TINT_ASSERT(IR, current_control);
if (auto* c = current_control->As<Loop>()) { if (auto* c = current_control->As<Loop>()) {
BranchTo(c->Merge()); SetBranch(builder_.ExitLoop(c));
} else if (auto* s = current_control->As<Switch>()) { } else if (auto* s = current_control->As<Switch>()) {
BranchTo(s->Merge()); SetBranch(builder_.ExitSwitch(s));
} else { } else {
TINT_UNREACHABLE(IR, diagnostics_); TINT_UNREACHABLE(IR, diagnostics_);
} }
@ -964,14 +958,14 @@ class Impl {
// If the lhs is false, then that is the result we want to pass to the merge // If the lhs is false, then that is the result we want to pass to the merge
// block as our argument // block as our argument
current_flow_block_ = if_inst->False(); current_flow_block_ = if_inst->False();
BranchTo(if_inst->Merge(), std::move(alt_args)); SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
current_flow_block_ = if_inst->True(); current_flow_block_ = if_inst->True();
} else { } else {
// If the lhs is true, then that is the result we want to pass to the merge // If the lhs is true, then that is the result we want to pass to the merge
// block as our argument // block as our argument
current_flow_block_ = if_inst->True(); current_flow_block_ = if_inst->True();
BranchTo(if_inst->Merge(), std::move(alt_args)); SetBranch(builder_.ExitIf(if_inst, std::move(alt_args)));
current_flow_block_ = if_inst->False(); current_flow_block_ = if_inst->False();
} }
@ -983,7 +977,7 @@ class Impl {
utils::Vector<Value*, 1> args; utils::Vector<Value*, 1> args;
args.Push(rhs.Get()); args.Push(rhs.Get());
BranchTo(if_inst->Merge(), std::move(args)); SetBranch(builder_.ExitIf(if_inst, std::move(args)));
} }
current_flow_block_ = if_inst->Merge(); current_flow_block_ = if_inst->Merge();

View File

@ -60,13 +60,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = add %3, 1u %4:u32 = add %3, 1u
store %v1, %4 store %v1, %4
@ -87,13 +84,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = add %3, 1u %4:u32 = add %3, 1u
store %v1, %4 store %v1, %4
@ -137,13 +131,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, i32, read_write> = var %v1:ptr<private, i32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:i32 = load %v1 %3:i32 = load %v1
%4:i32 = sub %3, 1i %4:i32 = sub %3, 1i
store %v1, %4 store %v1, %4
@ -164,13 +155,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = sub %3, 1u %4:u32 = sub %3, 1u
store %v1, %4 store %v1, %4
@ -214,13 +202,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = mul %3, 1u %4:u32 = mul %3, 1u
store %v1, %4 store %v1, %4
@ -264,13 +249,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = div %3, 1u %4:u32 = div %3, 1u
store %v1, %4 store %v1, %4
@ -314,13 +296,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = mod %3, 1u %4:u32 = mod %3, 1u
store %v1, %4 store %v1, %4
@ -364,13 +343,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, bool, read_write> = var %v1:ptr<private, bool, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:bool = load %v1 %3:bool = load %v1
%4:bool = and %3, false %4:bool = and %3, false
store %v1, %4 store %v1, %4
@ -414,13 +390,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, bool, read_write> = var %v1:ptr<private, bool, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:bool = load %v1 %3:bool = load %v1
%4:bool = or %3, false %4:bool = or %3, false
store %v1, %4 store %v1, %4
@ -464,13 +437,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = xor %3, 1u %4:u32 = xor %3, 1u
store %v1, %4 store %v1, %4
@ -499,12 +469,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
if %3 [t: %b3, f: %b4, m: %b5] if %3 [t: %b3, f: %b4, m: %b5]
# True block # True block
%b3 = block { %b3 = block {
br %b5 false exit_if %b5 false
} }
# False block # False block
%b4 = block { %b4 = block {
br %b5 %3 exit_if %b5 %3
} }
# Merge block # Merge block
@ -512,12 +482,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
if %4:bool [t: %b6, f: %b7, m: %b8] if %4:bool [t: %b6, f: %b7, m: %b8]
# True block # True block
%b6 = block { %b6 = block {
br %b8 exit_if %b8
} }
# False block # False block
%b7 = block { %b7 = block {
br %b8 exit_if %b8
} }
# Merge block # Merge block
@ -527,9 +497,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
} }
} }
} }
)"); )");
} }
@ -553,12 +521,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
if %3 [t: %b3, f: %b4, m: %b5] if %3 [t: %b3, f: %b4, m: %b5]
# True block # True block
%b3 = block { %b3 = block {
br %b5 %3 exit_if %b5 %3
} }
# False block # False block
%b4 = block { %b4 = block {
br %b5 true exit_if %b5 true
} }
# Merge block # Merge block
@ -566,12 +534,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
if %4:bool [t: %b6, f: %b7, m: %b8] if %4:bool [t: %b6, f: %b7, m: %b8]
# True block # True block
%b6 = block { %b6 = block {
br %b8 exit_if %b8
} }
# False block # False block
%b7 = block { %b7 = block {
br %b8 exit_if %b8
} }
# Merge block # Merge block
@ -581,9 +549,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
} }
} }
} }
)"); )");
} }
@ -760,13 +726,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = shiftl %3, 1u %4:u32 = shiftl %3, 1u
store %v1, %4 store %v1, %4
@ -810,13 +773,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v1:ptr<private, u32, read_write> = var %v1:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = shiftr %3, 1u %4:u32 = shiftr %3, 1u
store %v1, %4 store %v1, %4
@ -853,12 +813,12 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
%7:f32 = mul 2.29999995231628417969f, %6 %7:f32 = mul 2.29999995231628417969f, %6
%8:f32 = div %5, %7 %8:f32 = div %5, %7
%9:bool = gt 2.5f, %8 %9:bool = gt 2.5f, %8
br %b5 %9 exit_if %b5 %9
} }
# False block # False block
%b4 = block { %b4 = block {
br %b5 %4 exit_if %b5 %4
} }
# Merge block # Merge block
@ -867,7 +827,6 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
} }
} }
} }
)"); )");
} }

View File

@ -37,13 +37,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%i:ptr<private, f32, read_write> = var, 1.0f %i:ptr<private, f32, read_write> = var, 1.0f
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:f32 = load %i %3:f32 = load %i
%tint_symbol:f32 = asin %3 %tint_symbol:f32 = asin %3
ret ret

View File

@ -102,13 +102,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%i:ptr<private, i32, read_write> = var, 1i %i:ptr<private, i32, read_write> = var, 1i
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:i32 = load %i %3:i32 = load %i
%tint_symbol:f32 = convert i32, %3 %tint_symbol:f32 = convert i32, %3
ret ret
@ -127,11 +124,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f %i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
br %b2 # root_end
} }
%b2 = root_terminator
)"); )");
} }
@ -146,13 +140,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%i:ptr<private, f32, read_write> = var, 1.0f %i:ptr<private, f32, read_write> = var, 1.0f
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
%3:f32 = load %i %3:f32 = load %i
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3 %tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %3
ret ret

View File

@ -38,13 +38,10 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%a:ptr<private, u32, read_write> = var %a:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
store %a, 4u store %a, 4u
ret ret
} }

View File

@ -141,7 +141,7 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -149,12 +149,12 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
if true [t: %b2, f: %b3, m: %b4] if true [t: %b2, f: %b3, m: %b4]
# True block # True block
%b2 = block { %b2 = block {
br %b4 exit_if %b4
} }
# False block # False block
%b3 = block { %b3 = block {
br %b4 exit_if %b4
} }
# Merge block # Merge block
@ -163,7 +163,6 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
} }
} }
} }
)"); )");
} }
@ -182,7 +181,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -192,9 +191,10 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
%b2 = block { %b2 = block {
ret ret
} }
# False block # False block
%b3 = block { %b3 = block {
br %b4 exit_if %b4
} }
# Merge block # Merge block
@ -203,7 +203,6 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
} }
} }
} }
)"); )");
} }
@ -222,7 +221,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -230,20 +229,20 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
if true [t: %b2, f: %b3, m: %b4] if true [t: %b2, f: %b3, m: %b4]
# True block # True block
%b2 = block { %b2 = block {
br %b4 exit_if %b4
} }
# False block # False block
%b3 = block { %b3 = block {
ret ret
} }
# Merge block # Merge block
%b4 = block { %b4 = block {
ret ret
} }
} }
} }
)"); )");
} }
@ -262,7 +261,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
EXPECT_EQ(1u, flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, flow->False()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -272,13 +271,13 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
%b2 = block { %b2 = block {
ret ret
} }
# False block # False block
%b3 = block { %b3 = block {
ret ret
} }
} }
} }
)"); )");
} }
@ -306,20 +305,19 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
%b2 = block { %b2 = block {
loop [s: %b5, m: %b6] loop [s: %b5, m: %b6]
%b5 = block { %b5 = block {
br %b6 exit_loop %b6
} }
# Merge block # Merge block
%b6 = block { %b6 = block {
br %b4 exit_if %b4
} }
} }
# False block # False block
%b3 = block { %b3 = block {
br %b4 exit_if %b4
} }
# Merge block # Merge block
@ -328,7 +326,6 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
} }
} }
} }
)"); )");
} }
@ -345,7 +342,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(1u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@ -354,7 +351,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
%b1 = block { %b1 = block {
loop [s: %b2, m: %b3] loop [s: %b2, m: %b3]
%b2 = block { %b2 = block {
br %b3 exit_loop %b3
} }
# Merge block # Merge block
@ -363,7 +360,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
} }
} }
} }
)"); )");
} }
@ -383,12 +379,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -398,12 +394,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
if true [t: %b5, f: %b6, m: %b7] if true [t: %b5, f: %b6, m: %b7]
# True block # True block
%b5 = block { %b5 = block {
br %b4 exit_loop %b4
} }
# False block # False block
%b6 = block { %b6 = block {
br %b7 exit_if %b7
} }
# Merge block # Merge block
@ -411,7 +407,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
continue %b3 continue %b3
} }
} }
# Continuing block # Continuing block
@ -425,7 +420,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
} }
} }
} }
)"); )");
} }
@ -443,7 +437,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
@ -466,7 +460,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
} }
} }
} }
)"); )");
} }
@ -500,7 +493,6 @@ TEST_F(IR_BuilderImplTest, Loop_Continuing_Body_Scope) {
} }
} }
} }
)"); )");
} }
@ -519,12 +511,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(2u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -536,9 +528,10 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
%b5 = block { %b5 = block {
ret ret
} }
# False block # False block
%b6 = block { %b6 = block {
br %b7 exit_if %b7
} }
# Merge block # Merge block
@ -546,7 +539,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
continue %b3 continue %b3
} }
} }
# Continuing block # Continuing block
@ -560,7 +552,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
} }
} }
} }
)"); )");
} }
@ -577,7 +568,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
@ -590,7 +581,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
} }
} }
} }
)"); )");
} }
@ -616,7 +606,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Merge()->InboundBranches().Length());
@ -629,7 +619,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
} }
} }
} }
)"); )");
} }
@ -648,12 +637,12 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(1u, loop_flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(0u, loop_flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, loop_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(0u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -663,15 +652,14 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
if true [t: %b4, f: %b5] if true [t: %b4, f: %b5]
# True block # True block
%b4 = block { %b4 = block {
br %b3 exit_loop %b3
} }
# False block # False block
%b5 = block { %b5 = block {
br %b3 exit_loop %b3
} }
} }
# Merge block # Merge block
@ -680,7 +668,6 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
} }
} }
} }
)"); )");
} }
@ -712,12 +699,12 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
if true [t: %b8, f: %b9, m: %b10] if true [t: %b8, f: %b9, m: %b10]
# True block # True block
%b8 = block { %b8 = block {
br %b7 exit_loop %b7
} }
# False block # False block
%b9 = block { %b9 = block {
br %b10 exit_if %b10
} }
# Merge block # Merge block
@ -730,7 +717,7 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
# False block # False block
%b12 = block { %b12 = block {
br %b13 exit_if %b13
} }
# Merge block # Merge block
@ -738,17 +725,15 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
continue %b6 continue %b6
} }
} }
} }
# Continuing block # Continuing block
%b6 = block { %b6 = block {
loop [s: %b14, m: %b15] loop [s: %b14, m: %b15]
%b14 = block { %b14 = block {
br %b15 exit_loop %b15
} }
# Merge block # Merge block
@ -768,10 +753,8 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
break_if false %b5 break_if false %b5
} }
} }
} }
# Merge block # Merge block
@ -779,12 +762,12 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
if true [t: %b19, f: %b20, m: %b21] if true [t: %b19, f: %b20, m: %b21]
# True block # True block
%b19 = block { %b19 = block {
br %b4 exit_loop %b4
} }
# False block # False block
%b20 = block { %b20 = block {
br %b21 exit_if %b21
} }
# Merge block # Merge block
@ -792,10 +775,8 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
continue %b3 continue %b3
} }
} }
} }
# Continuing block # Continuing block
@ -809,7 +790,6 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
} }
} }
} }
)"); )");
} }
@ -830,12 +810,12 @@ TEST_F(IR_BuilderImplTest, While) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -845,12 +825,12 @@ TEST_F(IR_BuilderImplTest, While) {
if false [t: %b5, f: %b6, m: %b7] if false [t: %b5, f: %b6, m: %b7]
# True block # True block
%b5 = block { %b5 = block {
br %b7 exit_if %b7
} }
# False block # False block
%b6 = block { %b6 = block {
br %b4 exit_loop %b4
} }
# Merge block # Merge block
@ -858,7 +838,6 @@ TEST_F(IR_BuilderImplTest, While) {
continue %b3 continue %b3
} }
} }
# Continuing block # Continuing block
@ -872,7 +851,6 @@ TEST_F(IR_BuilderImplTest, While) {
} }
} }
} }
)"); )");
} }
@ -893,12 +871,12 @@ TEST_F(IR_BuilderImplTest, While_Return) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(2u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -908,12 +886,12 @@ TEST_F(IR_BuilderImplTest, While_Return) {
if true [t: %b5, f: %b6, m: %b7] if true [t: %b5, f: %b6, m: %b7]
# True block # True block
%b5 = block { %b5 = block {
br %b7 exit_if %b7
} }
# False block # False block
%b6 = block { %b6 = block {
br %b4 exit_loop %b4
} }
# Merge block # Merge block
@ -934,7 +912,6 @@ TEST_F(IR_BuilderImplTest, While_Return) {
} }
} }
} }
)"); )");
} }
@ -973,7 +950,7 @@ TEST_F(IR_BuilderImplTest, DISABLED_For) {
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->True()->InboundBranches().Length());
EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->False()->InboundBranches().Length());
EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), R"()"); EXPECT_EQ(Disassemble(m), R"()");
} }
@ -990,7 +967,7 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
ASSERT_EQ(1u, m.functions.Length()); ASSERT_EQ(1u, m.functions.Length());
EXPECT_EQ(1u, flow->Start()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Start()->InboundBranches().Length());
EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Continuing()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
@ -999,7 +976,7 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
%b1 = block { %b1 = block {
loop [s: %b2, m: %b3] loop [s: %b2, m: %b3]
%b2 = block { %b2 = block {
br %b3 exit_loop %b3
} }
# Merge block # Merge block
@ -1008,7 +985,6 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
} }
} }
} }
)"); )");
} }
@ -1047,7 +1023,7 @@ TEST_F(IR_BuilderImplTest, Switch) {
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[2].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[2].Start()->InboundBranches().Length());
EXPECT_EQ(4u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -1055,17 +1031,17 @@ TEST_F(IR_BuilderImplTest, Switch) {
switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4), m: %b5] switch 1i [c: (0i, %b2), c: (1i, %b3), c: (default, %b4), m: %b5]
# Case block # Case block
%b2 = block { %b2 = block {
br %b5 exit_switch %b5
} }
# Case block # Case block
%b3 = block { %b3 = block {
br %b5 exit_switch %b5
} }
# Case block # Case block
%b4 = block { %b4 = block {
br %b5 exit_switch %b5
} }
# Merge block # Merge block
@ -1074,7 +1050,6 @@ TEST_F(IR_BuilderImplTest, Switch) {
} }
} }
} }
)"); )");
} }
@ -1109,7 +1084,7 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
EXPECT_TRUE(cases[0].selectors[2].IsDefault()); EXPECT_TRUE(cases[0].selectors[2].IsDefault());
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -1117,7 +1092,7 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
switch 1i [c: (0i 1i default, %b2), m: %b3] switch 1i [c: (0i 1i default, %b2), m: %b3]
# Case block # Case block
%b2 = block { %b2 = block {
br %b3 exit_switch %b3
} }
# Merge block # Merge block
@ -1126,7 +1101,6 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
} }
} }
} }
)"); )");
} }
@ -1149,7 +1123,7 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
EXPECT_TRUE(cases[0].selectors[0].IsDefault()); EXPECT_TRUE(cases[0].selectors[0].IsDefault());
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -1157,7 +1131,7 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
switch 1i [c: (default, %b2), m: %b3] switch 1i [c: (default, %b2), m: %b3]
# Case block # Case block
%b2 = block { %b2 = block {
br %b3 exit_switch %b3
} }
# Merge block # Merge block
@ -1166,7 +1140,6 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
} }
} }
} }
)"); )");
} }
@ -1197,7 +1170,7 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
EXPECT_EQ(3u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, flow->Merge()->InboundBranches().Length());
// This is 1 because the if is dead-code eliminated and the return doesn't happen. // This is 1 because the if is dead-code eliminated and the return doesn't happen.
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
@ -1206,12 +1179,12 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
switch 1i [c: (0i, %b2), c: (default, %b3), m: %b4] switch 1i [c: (0i, %b2), c: (default, %b3), m: %b4]
# Case block # Case block
%b2 = block { %b2 = block {
br %b4 exit_switch %b4
} }
# Case block # Case block
%b3 = block { %b3 = block {
br %b4 exit_switch %b4
} }
# Merge block # Merge block
@ -1220,7 +1193,6 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
} }
} }
} }
)"); )");
} }
@ -1254,7 +1226,7 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[0].Start()->InboundBranches().Length());
EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length()); EXPECT_EQ(1u, cases[1].Start()->InboundBranches().Length());
EXPECT_EQ(1u, flow->Merge()->InboundBranches().Length()); EXPECT_EQ(0u, flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b1 {
@ -1264,13 +1236,13 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
%b2 = block { %b2 = block {
ret ret
} }
# Case block # Case block
%b3 = block { %b3 = block {
ret ret
} }
} }
} }
)"); )");
} }

View File

@ -107,13 +107,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v2:ptr<private, i32, read_write> = var %v2:ptr<private, i32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
ret ret
} }
} }
@ -134,13 +131,10 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%v3:ptr<private, i32, read_write> = var %v3:ptr<private, i32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b2 {
%b2 = block {
%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %b3 {
%b3 = block {
store %v3, 42i store %v3, 42i
ret ret
} }

View File

@ -35,11 +35,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%a:ptr<private, u32, read_write> = var %a:ptr<private, u32, read_write> = var
br %b2 # root_end
} }
%b2 = root_terminator
)"); )");
} }
@ -53,11 +50,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%b1 = block { %b1 = block {
%a:ptr<private, u32, read_write> = var, 2u %a:ptr<private, u32, read_write> = var, 2u
br %b2 # root_end
} }
%b2 = root_terminator
)"); )");
} }

View File

@ -19,7 +19,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::If);
namespace tint::ir { namespace tint::ir {
If::If(Value* cond, Block* t, Block* f, Block* m) If::If(Value* cond, Block* t, Block* f, Block* m)
: Base(m), condition_(cond), true_(t), false_(f), merge_(m) { : Base(utils::Empty), condition_(cond), true_(t), false_(f), merge_(m) {
TINT_ASSERT(IR, true_); TINT_ASSERT(IR, true_);
TINT_ASSERT(IR, false_); TINT_ASSERT(IR, false_);
TINT_ASSERT(IR, merge_); TINT_ASSERT(IR, merge_);

View File

@ -18,7 +18,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Loop);
namespace tint::ir { namespace tint::ir {
Loop::Loop(Block* s, Block* c, Block* m) : Base(s), start_(s), continuing_(c), merge_(m) { Loop::Loop(Block* s, Block* c, Block* m)
: Base(utils::Empty), start_(s), continuing_(c), merge_(m) {
TINT_ASSERT(IR, start_); TINT_ASSERT(IR, start_);
TINT_ASSERT(IR, continuing_); TINT_ASSERT(IR, continuing_);
TINT_ASSERT(IR, merge_); TINT_ASSERT(IR, merge_);

View File

@ -18,7 +18,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Switch);
namespace tint::ir { namespace tint::ir {
Switch::Switch(Value* cond, Block* m) : Base(m), condition_(cond), merge_(m) { Switch::Switch(Value* cond, Block* m) : Base(utils::Empty), condition_(cond), merge_(m) {
TINT_ASSERT(IR, condition_); TINT_ASSERT(IR, condition_);
TINT_ASSERT(IR, merge_); TINT_ASSERT(IR, merge_);
condition_->AddUsage(this); condition_->AddUsage(this);

View File

@ -20,6 +20,7 @@
#include "src/tint/ir/block.h" #include "src/tint/ir/block.h"
#include "src/tint/ir/call.h" #include "src/tint/ir/call.h"
#include "src/tint/ir/constant.h" #include "src/tint/ir/constant.h"
#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/if.h" #include "src/tint/ir/if.h"
#include "src/tint/ir/instruction.h" #include "src/tint/ir/instruction.h"
#include "src/tint/ir/load.h" #include "src/tint/ir/load.h"
@ -121,46 +122,27 @@ class State {
while (block) { while (block) {
TINT_ASSERT(IR, block->HasBranchTarget()); TINT_ASSERT(IR, block->HasBranchTarget());
enum Status { kContinue, kStop, kError }; for (auto* inst : block->Instructions()) {
auto stmt = Stmt(inst);
Status status = tint::Switch( if (TINT_UNLIKELY(!stmt)) {
block, return nullptr;
}
[&](const ir::Block* blk) { if (auto* s = stmt.Get()) {
for (auto* inst : blk->Instructions()) { stmts.Push(s);
auto stmt = Stmt(inst); }
if (TINT_UNLIKELY(!stmt)) {
return kError;
}
if (auto* s = stmt.Get()) {
stmts.Push(s);
}
}
if (auto* if_ = blk->Branch()->As<ir::If>()) {
if (if_->Merge()->HasBranchTarget()) {
block = if_->Merge();
return kContinue;
}
} else if (auto* switch_ = blk->Branch()->As<ir::Switch>()) {
if (switch_->Merge()->HasBranchTarget()) {
block = switch_->Merge();
return kContinue;
}
}
return kStop;
},
[&](Default) {
UNHANDLED_CASE(block);
return kError;
});
if (TINT_UNLIKELY(status == kError)) {
return nullptr;
} }
if (status == kStop) { if (auto* if_ = block->Branch()->As<ir::If>()) {
break; if (if_->Merge()->HasBranchTarget()) {
block = if_->Merge();
continue;
}
} else if (auto* switch_ = block->Branch()->As<ir::Switch>()) {
if (switch_->Merge()->HasBranchTarget()) {
block = switch_->Merge();
continue;
}
} }
break;
} }
return b.Block(std::move(stmts)); return b.Block(std::move(stmts));
@ -174,16 +156,20 @@ class State {
return nullptr; return nullptr;
} }
if (!IsEmpty(i->False(), i->Merge())) { auto* false_blk = i->False();
if (false_blk->Instructions().Length() > 1 ||
(false_blk->Instructions().Length() == 1 && false_blk->HasBranchTarget() &&
!false_blk->Branch()->Is<ir::ExitIf>())) {
// If the else target is an `if` which has a merge target that just bounces to the outer // If the else target is an `if` which has a merge target that just bounces to the outer
// if merge target then emit an 'else if' instead of a block statement for the else. // if merge target then emit an 'else if' instead of a block statement for the else.
if (auto* inst = i->False()->Instructions().Front()->As<ir::If>(); if (auto* inst = i->False()->Instructions().Front()->As<ir::If>()) {
inst && inst->Merge()->IsTrampoline(i->Merge())) { if (auto* br = inst->Merge()->Branch()->As<ir::ExitIf>(); br && br->If() == i) {
auto* f = If(inst); auto* f = If(inst);
if (!f) { if (!f) {
return nullptr; return nullptr;
}
return b.If(cond, t, b.Else(f));
} }
return b.If(cond, t, b.Else(f));
} else { } else {
auto* f = BlockGraph(i->False()); auto* f = BlockGraph(i->False());
if (!f) { if (!f) {
@ -192,7 +178,6 @@ class State {
return b.If(cond, t, b.Else(f)); return b.If(cond, t, b.Else(f));
} }
} }
return b.If(cond, t); return b.If(cond, t);
} }
@ -265,17 +250,6 @@ class State {
return b.Return(val); return b.Return(val);
} }
/// @return true if there are no instructions between @p node and and @p stop_at
bool IsEmpty(const ir::Block* node, const ir::Block* stop_at) {
if (node->Instructions().IsEmpty()) {
return true;
}
if (auto* br = node->Instructions().Front()->As<Branch>()) {
return !br->Is<ir::Return>() && br->To() == stop_at;
}
return false;
}
utils::Result<const ast::Statement*> Stmt(const ir::Instruction* inst) { utils::Result<const ast::Statement*> Stmt(const ir::Instruction* inst) {
return tint::Switch<utils::Result<const ast::Statement*>>( return tint::Switch<utils::Result<const ast::Statement*>>(
inst, // inst, //

View File

@ -19,6 +19,7 @@
#include "spirv/unified1/spirv.h" #include "spirv/unified1/spirv.h"
#include "src/tint/ir/binary.h" #include "src/tint/ir/binary.h"
#include "src/tint/ir/block.h" #include "src/tint/ir/block.h"
#include "src/tint/ir/exit_if.h"
#include "src/tint/ir/if.h" #include "src/tint/ir/if.h"
#include "src/tint/ir/load.h" #include "src/tint/ir/load.h"
#include "src/tint/ir/module.h" #include "src/tint/ir/module.h"
@ -354,25 +355,24 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
} }
void GeneratorImplIr::EmitBranch(const ir::Branch* b) { void GeneratorImplIr::EmitBranch(const ir::Branch* b) {
if (b->Is<ir::Return>()) { tint::Switch( //
if (!b->Args().IsEmpty()) { b, //
TINT_ASSERT(Writer, b->Args().Length() == 1u); [&](const ir::Return*) {
OperandList operands; if (!b->Args().IsEmpty()) {
operands.push_back(Value(b->Args()[0])); TINT_ASSERT(Writer, b->Args().Length() == 1u);
current_function_.push_inst(spv::Op::OpReturnValue, operands); OperandList operands;
} else { operands.push_back(Value(b->Args()[0]));
current_function_.push_inst(spv::Op::OpReturn, {}); current_function_.push_inst(spv::Op::OpReturnValue, operands);
} } else {
return; current_function_.push_inst(spv::Op::OpReturn, {});
} }
return;
Switch( },
b->To(), [&](const ir::ExitIf* if_) {
[&](const ir::Block* blk) { current_function_.push_inst(spv::Op::OpBranch, {Label(blk)}); }, current_function_.push_inst(spv::Op::OpBranch, {Label(if_->If()->Merge())});
},
[&](Default) { [&](Default) {
// A block may not have an outward branch (e.g. an unreachable merge TINT_ICE(Writer, diagnostics_) << "unimplemented branch: " << b->TypeInfo().name;
// block).
current_function_.push_inst(spv::Op::OpUnreachable, {});
}); });
} }
@ -388,10 +388,12 @@ void GeneratorImplIr::EmitIf(const ir::If* i) {
uint32_t merge_label = Label(merge_block); uint32_t merge_label = Label(merge_block);
uint32_t true_label = merge_label; uint32_t true_label = merge_label;
uint32_t false_label = merge_label; uint32_t false_label = merge_label;
if (true_block->Instructions().Length() > 1 || true_block->Branch()->To() != merge_block) { if (true_block->Instructions().Length() > 1 ||
(true_block->HasBranchTarget() && !true_block->Branch()->Is<ir::ExitIf>())) {
true_label = Label(true_block); true_label = Label(true_block);
} }
if (false_block->Instructions().Length() > 1 || false_block->Branch()->To() != merge_block) { if (false_block->Instructions().Length() > 1 ||
(false_block->HasBranchTarget() && !false_block->Branch()->Is<ir::ExitIf>())) {
false_label = Label(false_block); false_label = Label(false_block);
} }

View File

@ -23,8 +23,8 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) {
auto* func = b.CreateFunction("foo", mod.Types().void_()); auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true)); auto* i = b.CreateIf(b.Constant(true));
i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); i->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
func->StartTarget()->SetInstructions(utils::Vector{i}); func->StartTarget()->SetInstructions(utils::Vector{i});
@ -49,12 +49,12 @@ TEST_F(SpvGeneratorImplTest, If_FalseEmpty) {
auto* func = b.CreateFunction("foo", mod.Types().void_()); auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true)); auto* i = b.CreateIf(b.Constant(true));
i->False()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); i->False()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
auto* true_block = i->True(); auto* true_block = i->True();
true_block->SetInstructions(utils::Vector{ true_block->SetInstructions(
b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())}); utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
func->StartTarget()->SetInstructions(utils::Vector{i}); func->StartTarget()->SetInstructions(utils::Vector{i});
@ -83,12 +83,12 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty) {
auto* func = b.CreateFunction("foo", mod.Types().void_()); auto* func = b.CreateFunction("foo", mod.Types().void_());
auto* i = b.CreateIf(b.Constant(true)); auto* i = b.CreateIf(b.Constant(true));
i->True()->SetInstructions(utils::Vector{b.Branch(i->Merge())}); i->True()->SetInstructions(utils::Vector{b.ExitIf(i)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); i->Merge()->SetInstructions(utils::Vector{b.Return(func)});
auto* false_block = i->False(); auto* false_block = i->False();
false_block->SetInstructions(utils::Vector{ false_block->SetInstructions(
b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.Branch(i->Merge())}); utils::Vector{b.Add(mod.Types().i32(), b.Constant(1_i), b.Constant(1_i)), b.ExitIf(i)});
func->StartTarget()->SetInstructions(utils::Vector{i}); func->StartTarget()->SetInstructions(utils::Vector{i});

View File

@ -100,7 +100,7 @@ TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) {
v->SetInitializer(b.Constant(42_i)); v->SetInitializer(b.Constant(42_i));
auto* i = b.CreateIf(b.Constant(true)); auto* i = b.CreateIf(b.Constant(true));
i->True()->SetInstructions(utils::Vector{v, b.Branch(i->Merge())}); i->True()->SetInstructions(utils::Vector{v, b.ExitIf(i)});
i->False()->SetInstructions(utils::Vector{b.Return(func)}); i->False()->SetInstructions(utils::Vector{b.Return(func)});
i->Merge()->SetInstructions(utils::Vector{b.Return(func)}); i->Merge()->SetInstructions(utils::Vector{b.Return(func)});