encounter
/
dawn-cmake
mirror of https://github.com/encounter/dawn-cmake.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[ir] Update UserCall to hold Function Values. 

This CL updates the `UserCall` to hold a `Function` value instead of the
functions name symbol. The name symbol is also removed from the function
itself and stored into the module like all other values.

Bug: tint:1718
Change-Id: I6c94ce435a6a260f9fe953a04278129b35b3bd39
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134303
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
dan sinclair 2023-05-25 13:39:47 +00:00 committed by Dawn LUCI CQ
parent 79fe4ab42f
commit a2b489b900
23 changed files with 179 additions and 213 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

15
src/tint/ir/builder.cc
View File

@ -49,19 +49,12 @@ Function* Builder::CreateFunction(std::string_view name,
const type::Type* return_type, const type::Type* return_type,
Function::PipelineStage stage, Function::PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size) { std::optional<std::array<uint32_t, 3>> wg_size) {
return CreateFunction(ir.symbols.Register(name), return_type, stage, wg_size);
}
Function* Builder::CreateFunction(Symbol name,
const type::Type* return_type,
Function::PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size) {
TINT_ASSERT(IR, return_type); TINT_ASSERT(IR, return_type);
auto* ir_func = ir.values.Create<Function>(name, return_type, stage, wg_size); auto* ir_func = ir.values.Create<Function>(return_type, stage, wg_size);
ir_func->SetStartTarget(CreateBlock()); ir_func->SetStartTarget(CreateBlock());
ir_func->SetEndTarget(CreateFunctionTerminator()); ir_func->SetEndTarget(CreateFunctionTerminator());
ir.SetName(ir_func, name);
return ir_func; return ir_func;
} }
@ -182,9 +175,9 @@ ir::Discard* Builder::Discard() {
} }
ir::UserCall* Builder::UserCall(const type::Type* type, ir::UserCall* Builder::UserCall(const type::Type* type,
Symbol name, Function* func,
utils::VectorRef<Value*> args) { utils::VectorRef<Value*> args) {
return ir.values.Create<ir::UserCall>(type, name, std::move(args)); return ir.values.Create<ir::UserCall>(type, func, std::move(args));
} }
ir::Convert* Builder::Convert(const type::Type* to, ir::Convert* Builder::Convert(const type::Type* to,

15
src/tint/ir/builder.h
View File

@ -79,17 +79,6 @@ class Builder {
Function::PipelineStage stage = Function::PipelineStage::kUndefined, Function::PipelineStage stage = Function::PipelineStage::kUndefined,
std::optional<std::array<uint32_t, 3>> wg_size = {}); std::optional<std::array<uint32_t, 3>> wg_size = {});
/// Creates a function flow node
/// @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
Function* CreateFunction(Symbol 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 flow node
/// @param condition the if condition /// @param condition the if condition
/// @returns the flow node /// @returns the flow node
@ -299,10 +288,10 @@ class Builder {
/// Creates a user function call instruction /// Creates a user function call instruction
/// @param type the return type of the call /// @param type the return type of the call
/// @param name the name of the function being called /// @param func the function being called
/// @param args the call arguments /// @param args the call arguments
/// @returns the instruction /// @returns the instruction
ir::UserCall* UserCall(const type::Type* type, Symbol name, utils::VectorRef<Value*> args); ir::UserCall* UserCall(const type::Type* type, Function* func, utils::VectorRef<Value*> args);
/// Creates a value conversion instruction /// Creates a value conversion instruction
/// @param to the type converted to /// @param to the type converted to

4
src/tint/ir/debug.cc
View File

@ -78,8 +78,8 @@ std::string Debug::AsDotGraph(const Module* mod) {
out << "digraph G {" << std::endl; out << "digraph G {" << std::endl;
for (const auto* func : mod->functions) { for (const auto* func : mod->functions) {
// Cluster each function to label and draw a box around it. // Cluster each function to label and draw a box around it.
out << "subgraph cluster_" << func->Name().Name() << " {" << std::endl; out << "subgraph cluster_" << mod->NameOf(func).Name() << " {" << std::endl;
out << R"(label=")" << func->Name().Name() << R"(")" << std::endl; out << R"(label=")" << mod->NameOf(func).Name() << R"(")" << std::endl;
out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl; out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl;
out << name_for(func->EndTarget()) << R"( [label="end"])" << std::endl; out << name_for(func->EndTarget()) << R"( [label="end"])" << std::endl;
Graph(func->StartTarget()); Graph(func->StartTarget());

4
src/tint/ir/disassembler.cc
View File

@ -151,7 +151,7 @@ void Disassembler::Walk(const Block* blk) {
void Disassembler::EmitFunction(const Function* func) { void Disassembler::EmitFunction(const Function* func) {
in_function_ = true; in_function_ = true;
Indent() << "%" << IdOf(func) << " = func " << func->Name().Name() << "("; Indent() << "%" << IdOf(func) << " = func(";
for (auto* p : func->Params()) { for (auto* p : func->Params()) {
if (p != func->Params().Front()) { if (p != func->Params().Front()) {
out_ << ", "; out_ << ", ";
@ -292,7 +292,7 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
}, },
[&](const ir::UserCall* uc) { [&](const ir::UserCall* uc) {
EmitValueWithType(uc); EmitValueWithType(uc);
out_ << " = call " << uc->Name().Name(); out_ << " = call %" << IdOf(uc->Func());
if (!uc->Args().IsEmpty()) { if (!uc->Args().IsEmpty()) {
out_ << ", "; out_ << ", ";
} }

13
src/tint/ir/from_program.cc
View File

@ -196,10 +196,6 @@ class Impl {
return nullptr; return nullptr;
} }
Symbol CloneSymbol(Symbol sym) const {
return clone_ctx_.type_ctx.dst.st->Register(sym.Name());
}
ResultType EmitModule() { ResultType EmitModule() {
auto* sem = program_->Sem().Module(); auto* sem = program_->Sem().Module();
@ -250,11 +246,13 @@ class Impl {
const auto* sem = program_->Sem().Get(ast_func); const auto* sem = program_->Sem().Get(ast_func);
auto* ir_func = builder_.CreateFunction(CloneSymbol(ast_func->name->symbol), auto* ir_func = builder_.CreateFunction(ast_func->name->symbol.NameView(),
sem->ReturnType()->Clone(clone_ctx_.type_ctx)); sem->ReturnType()->Clone(clone_ctx_.type_ctx));
current_function_ = ir_func; current_function_ = ir_func;
builder_.ir.functions.Push(ir_func); builder_.ir.functions.Push(ir_func);
scopes_.Set(ast_func->name->symbol, ir_func);
if (ast_func->IsEntryPoint()) { if (ast_func->IsEntryPoint()) {
switch (ast_func->PipelineStage()) { switch (ast_func->PipelineStage()) {
case ast::PipelineStage::kVertex: case ast::PipelineStage::kVertex:
@ -1136,8 +1134,9 @@ class Impl {
return utils::Failure; return utils::Failure;
} else { } else {
// Not a builtin and not a templated call, so this is a user function. // Not a builtin and not a templated call, so this is a user function.
auto name = CloneSymbol(expr->target->identifier->symbol); inst = builder_.UserCall(
inst = builder_.UserCall(ty, name, std::move(args)); ty, scopes_.Get(expr->target->identifier->symbol)->As<ir::Function>(),
std::move(args));
} }
if (inst == nullptr) { if (inst == nullptr) {
return utils::Failure; return utils::Failure;

148
src/tint/ir/from_program_binary_test.cc
View File

@ -34,15 +34,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Add) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = add %3, 4u %tint_symbol:u32 = add %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -67,7 +67,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Increment) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = add %3, 1u %4:u32 = add %3, 1u
@ -95,7 +95,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = add %3, 1u %4:u32 = add %3, 1u
@ -115,15 +115,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Subtract) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = sub %3, 4u %tint_symbol:u32 = sub %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -148,7 +148,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Decrement) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:i32 = load %v1 %3:i32 = load %v1
%4:i32 = sub %3, 1i %4:i32 = sub %3, 1i
@ -176,7 +176,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = sub %3, 1u %4:u32 = sub %3, 1u
@ -196,15 +196,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Multiply) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = mul %3, 4u %tint_symbol:u32 = mul %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -229,7 +229,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = mul %3, 1u %4:u32 = mul %3, 1u
@ -249,15 +249,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Div) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = div %3, 4u %tint_symbol:u32 = div %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -282,7 +282,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = div %3, 1u %4:u32 = div %3, 1u
@ -302,15 +302,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Modulo) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = mod %3, 4u %tint_symbol:u32 = mod %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -335,7 +335,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = mod %3, 1u %4:u32 = mod %3, 1u
@ -355,15 +355,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_And) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = and %3, 4u %tint_symbol:u32 = and %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -388,7 +388,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:bool = load %v1 %3:bool = load %v1
%4:bool = and %3, false %4:bool = and %3, false
@ -408,15 +408,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Or) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = or %3, 4u %tint_symbol:u32 = or %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -441,7 +441,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:bool = load %v1 %3:bool = load %v1
%4:bool = or %3, false %4:bool = or %3, false
@ -461,15 +461,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Xor) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = xor %3, 4u %tint_symbol:u32 = xor %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -494,7 +494,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = xor %3, 1u %4:u32 = xor %3, 1u
@ -514,15 +514,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():bool -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 true # return br %fn2 true # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:bool = call my_func %3:bool = call %my_func
if %3 [t: %fn4, f: %fn5, m: %fn6] if %3 [t: %fn4, f: %fn5, m: %fn6]
# True block # True block
%fn4 = block { %fn4 = block {
@ -570,15 +570,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():bool -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 true # return br %fn2 true # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:bool = call my_func %3:bool = call %my_func
if %3 [t: %fn4, f: %fn5, m: %fn6] if %3 [t: %fn4, f: %fn5, m: %fn6]
# True block # True block
%fn4 = block { %fn4 = block {
@ -626,15 +626,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Equal) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:bool = eq %3, 4u %tint_symbol:bool = eq %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -651,15 +651,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_NotEqual) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:bool = neq %3, 4u %tint_symbol:bool = neq %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -676,15 +676,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThan) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:bool = lt %3, 4u %tint_symbol:bool = lt %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -701,15 +701,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThan) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:bool = gt %3, 4u %tint_symbol:bool = gt %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -726,15 +726,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThanEqual) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:bool = lte %3, 4u %tint_symbol:bool = lte %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -751,15 +751,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThanEqual) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:bool = gte %3, 4u %tint_symbol:bool = gte %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -776,15 +776,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftLeft) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = shiftl %3, 4u %tint_symbol:u32 = shiftl %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -809,7 +809,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = shiftl %3, 1u %4:u32 = shiftl %3, 1u
@ -829,15 +829,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftRight) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0u # return br %fn2 0u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = shiftr %3, 4u %tint_symbol:u32 = shiftr %3, 4u
br %fn4 # return br %fn4 # return
} }
@ -862,7 +862,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = load %v1 %3:u32 = load %v1
%4:u32 = shiftr %3, 1u %4:u32 = shiftr %3, 1u
@ -884,21 +884,21 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():f32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0.0f # return br %fn2 0.0f # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:f32 = call my_func %3:f32 = call %my_func
%4:bool = lt %3, 2.0f %4:bool = lt %3, 2.0f
if %4 [t: %fn4, f: %fn5, m: %fn6] if %4 [t: %fn4, f: %fn5, m: %fn6]
# True block # True block
%fn4 = block { %fn4 = block {
%5:f32 = call my_func %5:f32 = call %my_func
%6:f32 = call my_func %6:f32 = call %my_func
%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
@ -931,15 +931,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func(%p:bool):bool -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:bool):bool -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 true # return br %fn2 true # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%tint_symbol:bool = call my_func, false %tint_symbol:bool = call %my_func, false
br %fn4 # return br %fn4 # return
} }
%fn4 = func_terminator %fn4 = func_terminator

18
src/tint/ir/from_program_call_test.cc
View File

@ -35,15 +35,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():f32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():f32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 0.0f # return br %fn2 0.0f # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:f32 = call my_func %3:f32 = call %my_func
%tint_symbol:f32 = bitcast %3 %tint_symbol:f32 = bitcast %3
br %fn4 # return br %fn4 # return
} }
@ -62,7 +62,7 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Discard) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func test_function():void [@fragment] -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():void [@fragment] -> %fn1 {
%fn1 = block { %fn1 = block {
discard discard
br %fn2 # return br %fn2 # return
@ -80,15 +80,15 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func(%p:f32):void -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func(%p:f32):void -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 # return br %fn2 # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%4:void = call my_func, 6.0f %4:void = call %my_func, 6.0f
br %fn4 # return br %fn4 # return
} }
%fn4 = func_terminator %fn4 = func_terminator
@ -112,7 +112,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:i32 = load %i %3:i32 = load %i
%tint_symbol:f32 = convert i32, %3 %tint_symbol:f32 = convert i32, %3
@ -157,7 +157,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = 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

2
src/tint/ir/from_program_materialize_test.cc
View File

@ -34,7 +34,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func test_function():f32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%test_function = func():f32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 2.0f # return br %fn2 2.0f # return
} }

2
src/tint/ir/from_program_store_test.cc
View File

@ -43,7 +43,7 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
store %a, 4u store %a, 4u
br %fn4 # return br %fn4 # return

56
src/tint/ir/from_program_test.cc
View File

@ -70,7 +70,7 @@ TEST_F(IR_BuilderImplTest, Func) {
EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func f():void -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%f = func():void -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 # return br %fn2 # return
} }
@ -95,7 +95,7 @@ TEST_F(IR_BuilderImplTest, Func_WithParam) {
EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func f(%a:u32):u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32):u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 %a # return br %fn2 %a # return
} }
@ -121,7 +121,7 @@ TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) {
EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined); EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func f(%a:u32, %b:i32, %c:bool):void -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%f = func(%a:u32, %b:i32, %c:bool):void -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 # return br %fn2 # return
} }
@ -159,7 +159,7 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
if true [t: %fn2, f: %fn3, m: %fn4] if true [t: %fn2, f: %fn3, m: %fn4]
# True block # True block
@ -203,7 +203,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
if true [t: %fn2, f: %fn3, m: %fn4] if true [t: %fn2, f: %fn3, m: %fn4]
# True block # True block
@ -246,7 +246,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
if true [t: %fn2, f: %fn3, m: %fn4] if true [t: %fn2, f: %fn3, m: %fn4]
# True block # True block
@ -289,7 +289,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
if true [t: %fn2, f: %fn3] if true [t: %fn2, f: %fn3]
# True block # True block
@ -324,7 +324,7 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
ASSERT_NE(loop_flow, nullptr); ASSERT_NE(loop_flow, nullptr);
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
if true [t: %fn2, f: %fn3, m: %fn4] if true [t: %fn2, f: %fn3, m: %fn4]
# True block # True block
@ -383,7 +383,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -432,7 +432,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -496,7 +496,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -547,7 +547,7 @@ TEST_F(IR_BuilderImplTest, Loop_Continuing_Body_Scope) {
auto m = res.Move(); auto m = res.Move();
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -611,7 +611,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3] loop [s: %fn2, c: %fn3]
%fn2 = block { %fn2 = block {
@ -665,7 +665,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3] loop [s: %fn2, c: %fn3]
%fn2 = block { %fn2 = block {
@ -712,7 +712,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
EXPECT_EQ(3u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(3u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -790,7 +790,7 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -843,7 +843,7 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), EXPECT_EQ(Disassemble(m.Get()),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -1002,7 +1002,7 @@ TEST_F(IR_BuilderImplTest, While) {
EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -1068,7 +1068,7 @@ TEST_F(IR_BuilderImplTest, While_Return) {
EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length()); EXPECT_EQ(2u, if_flow->Merge()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -1167,7 +1167,7 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
loop [s: %fn2, c: %fn3, m: %fn4] loop [s: %fn2, c: %fn3, m: %fn4]
%fn2 = block { %fn2 = block {
@ -1230,7 +1230,7 @@ TEST_F(IR_BuilderImplTest, Switch) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
switch 1i [c: (0i, %fn2), c: (1i, %fn3), c: (default, %fn4), m: %fn5] switch 1i [c: (0i, %fn2), c: (1i, %fn3), c: (default, %fn4), m: %fn5]
# Case block # Case block
@ -1295,7 +1295,7 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
switch 1i [c: (0i 1i default, %fn2), m: %fn3] switch 1i [c: (0i 1i default, %fn2), m: %fn3]
# Case block # Case block
@ -1338,7 +1338,7 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
switch 1i [c: (default, %fn2), m: %fn3] switch 1i [c: (default, %fn2), m: %fn3]
# Case block # Case block
@ -1390,7 +1390,7 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
switch 1i [c: (0i, %fn2), c: (default, %fn3), m: %fn4] switch 1i [c: (0i, %fn2), c: (default, %fn3), m: %fn4]
# Case block # Case block
@ -1449,7 +1449,7 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m), EXPECT_EQ(Disassemble(m),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
switch 1i [c: (0i, %fn2), c: (default, %fn3)] switch 1i [c: (0i, %fn2), c: (default, %fn3)]
# Case block # Case block
@ -1476,15 +1476,15 @@ TEST_F(IR_BuilderImplTest, Emit_Phony) {
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), EXPECT_EQ(Disassemble(m.Get()),
R"(%1 = func b():i32 -> %fn1 { R"(%b = func():i32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 1i # return br %fn2 1i # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:i32 = call b %3:i32 = call %b
br %fn4 # return br %fn4 # return
} }
%fn4 = func_terminator %fn4 = func_terminator

22
src/tint/ir/from_program_unary_test.cc
View File

@ -34,15 +34,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Not) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():bool -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():bool -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 false # return br %fn2 false # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:bool = call my_func %3:bool = call %my_func
%tint_symbol:bool = eq %3, false %tint_symbol:bool = eq %3, false
br %fn4 # return br %fn4 # return
} }
@ -59,15 +59,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Complement) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():u32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():u32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 1u # return br %fn2 1u # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:u32 = call my_func %3:u32 = call %my_func
%tint_symbol:u32 = complement %3 %tint_symbol:u32 = complement %3
br %fn4 # return br %fn4 # return
} }
@ -84,15 +84,15 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Negation) {
auto m = Build(); auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%1 = func my_func():i32 -> %fn1 { EXPECT_EQ(Disassemble(m.Get()), R"(%my_func = func():i32 -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 1i # return br %fn2 1i # return
} }
%fn2 = func_terminator %fn2 = func_terminator
} }
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
%3:i32 = call my_func %3:i32 = call %my_func
%tint_symbol:i32 = negation %3 %tint_symbol:i32 = negation %3
br %fn4 # return br %fn4 # return
} }
@ -118,7 +118,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
br %fn4 # return br %fn4 # return
} }
@ -146,7 +146,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
%fn2 = root_terminator %fn2 = root_terminator
%2 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 { %test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn3 {
%fn3 = block { %fn3 = block {
store %v3, 42i store %v3, 42i
br %fn4 # return br %fn4 # return

4
src/tint/ir/from_program_var_test.cc
View File

@ -69,7 +69,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), EXPECT_EQ(Disassemble(m.Get()),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
%a:ptr<function, u32, read_write> = var %a:ptr<function, u32, read_write> = var
br %fn2 # return br %fn2 # return
@ -88,7 +88,7 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
ASSERT_TRUE(m) << (!m ? m.Failure() : ""); ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), EXPECT_EQ(Disassemble(m.Get()),
R"(%1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { R"(%test_function = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
%a:ptr<function, u32, read_write> = var, 2u %a:ptr<function, u32, read_write> = var, 2u
br %fn2 # return br %fn2 # return

5
src/tint/ir/function.cc
View File

@ -18,11 +18,10 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Function);
namespace tint::ir { namespace tint::ir {
Function::Function(Symbol name, Function::Function(const type::Type* rt,
const type::Type* rt,
PipelineStage stage, PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size) std::optional<std::array<uint32_t, 3>> wg_size)
: Base(), name_(name), return_type_(rt), pipeline_stage_(stage), workgroup_size_(wg_size) {} : Base(), return_type_(rt), pipeline_stage_(stage), workgroup_size_(wg_size) {}
Function::~Function() = default; Function::~Function() = default;

9
src/tint/ir/function.h
View File

@ -21,7 +21,6 @@
#include "src/tint/ir/function_param.h" #include "src/tint/ir/function_param.h"
#include "src/tint/ir/value.h" #include "src/tint/ir/value.h"
#include "src/tint/symbol.h"
#include "src/tint/type/type.h" #include "src/tint/type/type.h"
// Forward declarations // Forward declarations
@ -64,19 +63,14 @@ class Function : public utils::Castable<Function, Value> {
}; };
/// Constructor /// Constructor
/// @param n the function name
/// @param rt the function return type /// @param rt the function return type
/// @param stage the function stage /// @param stage the function stage
/// @param wg_size the workgroup_size /// @param wg_size the workgroup_size
Function(Symbol n, Function(const type::Type* rt,
const type::Type* rt,
PipelineStage stage = PipelineStage::kUndefined, PipelineStage stage = PipelineStage::kUndefined,
std::optional<std::array<uint32_t, 3>> wg_size = {}); std::optional<std::array<uint32_t, 3>> wg_size = {});
~Function() override; ~Function() override;
/// @returns the function name
Symbol Name() const { return name_; }
/// Sets the function stage /// Sets the function stage
/// @param stage the stage to set /// @param stage the stage to set
void SetStage(PipelineStage stage) { pipeline_stage_ = stage; } void SetStage(PipelineStage stage) { pipeline_stage_ = stage; }
@ -130,7 +124,6 @@ class Function : public utils::Castable<Function, Value> {
FunctionTerminator* EndTarget() const { return end_target_; } FunctionTerminator* EndTarget() const { return end_target_; }
private: private:
Symbol name_;
const type::Type* return_type_; const type::Type* return_type_;
PipelineStage pipeline_stage_; PipelineStage pipeline_stage_;
std::optional<std::array<uint32_t, 3>> workgroup_size_; std::optional<std::array<uint32_t, 3>> workgroup_size_;

6
src/tint/ir/to_program.cc
View File

@ -91,7 +91,7 @@ class State {
const ast::Function* Fn(const Function* fn) { const ast::Function* Fn(const Function* fn) {
SCOPED_NESTING(); SCOPED_NESTING();
auto name = Sym(fn->Name()); auto name = NameOf(fn);
// TODO(crbug.com/tint/1915): Properly implement this when we've fleshed out Function // TODO(crbug.com/tint/1915): Properly implement this when we've fleshed out Function
utils::Vector<const ast::Parameter*, 1> params{}; utils::Vector<const ast::Parameter*, 1> params{};
auto ret_ty = Type(fn->ReturnType()); auto ret_ty = Type(fn->ReturnType());
@ -349,7 +349,7 @@ class State {
} }
return tint::Switch( return tint::Switch(
call, // call, //
[&](const ir::UserCall* c) { return b.Call(Sym(c->Name()), std::move(args)); }, [&](const ir::UserCall* c) { return b.Call(NameOf(c->Func()), std::move(args)); },
[&](Default) { [&](Default) {
UNHANDLED_CASE(call); UNHANDLED_CASE(call);
return nullptr; return nullptr;
@ -497,8 +497,6 @@ class State {
}); });
} }
Symbol Sym(const Symbol& s) { return b.Symbols().Register(s.NameView()); }
void Err(std::string str) { b.Diagnostics().add_error(diag::System::IR, std::move(str)); } void Err(std::string str) { b.Diagnostics().add_error(diag::System::IR, std::move(str)); }
}; };

2
src/tint/ir/transform/add_empty_entry_point.cc
View File

@ -35,7 +35,7 @@ void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
} }
ir::Builder builder(*ir); ir::Builder builder(*ir);
auto* ep = builder.CreateFunction(ir->symbols.New("unused_entry_point"), ir->Types().void_(), auto* ep = builder.CreateFunction("unused_entry_point", ir->Types().void_(),
Function::PipelineStage::kCompute, std::array{1u, 1u, 1u}); Function::PipelineStage::kCompute, std::array{1u, 1u, 1u});
ep->StartTarget()->SetInstructions(utils::Vector{builder.Branch(ep->EndTarget())}); ep->StartTarget()->SetInstructions(utils::Vector{builder.Branch(ep->EndTarget())});
ir->functions.Push(ep); ir->functions.Push(ep);

4
src/tint/ir/transform/add_empty_entry_point_test.cc
View File

@ -25,7 +25,7 @@ using IR_AddEmptyEntryPointTest = TransformTest;
TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) { TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
auto* expect = R"( auto* expect = R"(
%1 = func unused_entry_point():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 { %unused_entry_point = func():void [@compute @workgroup_size(1, 1, 1)] -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 # return br %fn2 # return
} }
@ -44,7 +44,7 @@ TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) {
mod.functions.Push(ep); mod.functions.Push(ep);
auto* expect = R"( auto* expect = R"(
%1 = func main():void [@fragment] -> %fn1 { %main = func():void [@fragment] -> %fn1 {
%fn1 = block { %fn1 = block {
br %fn2 # return br %fn2 # return
} }

6
src/tint/ir/user_call.cc
View File

@ -22,8 +22,10 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::UserCall);
namespace tint::ir { namespace tint::ir {
UserCall::UserCall(const type::Type* ty, Symbol n, utils::VectorRef<Value*> arguments) UserCall::UserCall(const type::Type* ty, Function* func, utils::VectorRef<Value*> arguments)
: Base(ty, std::move(arguments)), name_(n) {} : Base(ty, std::move(arguments)), func_(func) {
func->AddUsage(this);
}
UserCall::~UserCall() = default; UserCall::~UserCall() = default;

10
src/tint/ir/user_call.h
View File

@ -16,7 +16,7 @@
#define SRC_TINT_IR_USER_CALL_H_ #define SRC_TINT_IR_USER_CALL_H_
#include "src/tint/ir/call.h" #include "src/tint/ir/call.h"
#include "src/tint/symbol.h" #include "src/tint/ir/function.h"
#include "src/tint/utils/castable.h" #include "src/tint/utils/castable.h"
namespace tint::ir { namespace tint::ir {
@ -26,16 +26,16 @@ class UserCall : public utils::Castable<UserCall, Call> {
public: public:
/// Constructor /// Constructor
/// @param type the result type /// @param type the result type
/// @param name the function name /// @param func the function being called
/// @param args the function arguments /// @param args the function arguments
UserCall(const type::Type* type, Symbol name, utils::VectorRef<Value*> args); UserCall(const type::Type* type, Function* func, utils::VectorRef<Value*> args);
~UserCall() override; ~UserCall() override;
/// @returns the called function name /// @returns the called function name
Symbol Name() const { return name_; } const Function* Func() const { return func_; }
private: private:
Symbol name_; const Function* func_ = nullptr;
}; };
} // namespace tint::ir } // namespace tint::ir

22
src/tint/transform/manager_test.cc
View File

@ -50,8 +50,7 @@ class AST_AddFunction final : public ast::transform::Transform {
class IR_AddFunction final : public ir::transform::Transform { class IR_AddFunction final : public ir::transform::Transform {
void Run(ir::Module* mod, const DataMap&, DataMap&) const override { void Run(ir::Module* mod, const DataMap&, DataMap&) const override {
ir::Builder builder(*mod); ir::Builder builder(*mod);
auto* func = auto* func = builder.CreateFunction("ir_func", mod->Types().Get<type::Void>());
builder.CreateFunction(mod->symbols.New("ir_func"), mod->Types().Get<type::Void>());
func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())}); func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())});
mod->functions.Push(func); mod->functions.Push(func);
} }
@ -68,8 +67,7 @@ Program MakeAST() {
ir::Module MakeIR() { ir::Module MakeIR() {
ir::Module mod; ir::Module mod;
ir::Builder builder(mod); ir::Builder builder(mod);
auto* func = auto* func = builder.CreateFunction("main", mod.Types().Get<type::Void>());
builder.CreateFunction(builder.ir.symbols.New("main"), mod.Types().Get<type::Void>());
func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())}); func->StartTarget()->SetInstructions(utils::Vector{builder.Branch(func->EndTarget())});
builder.ir.functions.Push(func); builder.ir.functions.Push(func);
return mod; return mod;
@ -104,8 +102,8 @@ TEST_F(TransformManagerTest, IR_MutateInPlace) {
manager.Run(&ir, {}, outputs); manager.Run(&ir, {}, outputs);
ASSERT_EQ(ir.functions.Length(), 2u); ASSERT_EQ(ir.functions.Length(), 2u);
EXPECT_EQ(ir.functions[0]->Name().Name(), "main"); EXPECT_EQ(ir.NameOf(ir.functions[0]).Name(), "main");
EXPECT_EQ(ir.functions[1]->Name().Name(), "ir_func"); EXPECT_EQ(ir.NameOf(ir.functions[1]).Name(), "ir_func");
} }
TEST_F(TransformManagerTest, AST_MixedTransforms_AST_Before_IR) { TEST_F(TransformManagerTest, AST_MixedTransforms_AST_Before_IR) {
@ -150,9 +148,9 @@ TEST_F(TransformManagerTest, IR_MixedTransforms_AST_Before_IR) {
manager.Run(&ir, {}, outputs); manager.Run(&ir, {}, outputs);
ASSERT_EQ(ir.functions.Length(), 3u); ASSERT_EQ(ir.functions.Length(), 3u);
EXPECT_EQ(ir.functions[0]->Name().Name(), "ast_func"); EXPECT_EQ(ir.NameOf(ir.functions[0]).Name(), "ast_func");
EXPECT_EQ(ir.functions[1]->Name().Name(), "main"); EXPECT_EQ(ir.NameOf(ir.functions[1]).Name(), "main");
EXPECT_EQ(ir.functions[2]->Name().Name(), "ir_func"); EXPECT_EQ(ir.NameOf(ir.functions[2]).Name(), "ir_func");
} }
TEST_F(TransformManagerTest, IR_MixedTransforms_IR_Before_AST) { TEST_F(TransformManagerTest, IR_MixedTransforms_IR_Before_AST) {
@ -165,9 +163,9 @@ TEST_F(TransformManagerTest, IR_MixedTransforms_IR_Before_AST) {
manager.Run(&ir, {}, outputs); manager.Run(&ir, {}, outputs);
ASSERT_EQ(ir.functions.Length(), 3u); ASSERT_EQ(ir.functions.Length(), 3u);
EXPECT_EQ(ir.functions[0]->Name().Name(), "ast_func"); EXPECT_EQ(ir.NameOf(ir.functions[0]).Name(), "ast_func");
EXPECT_EQ(ir.functions[1]->Name().Name(), "main"); EXPECT_EQ(ir.NameOf(ir.functions[1]).Name(), "main");
EXPECT_EQ(ir.functions[2]->Name().Name(), "ir_func"); EXPECT_EQ(ir.NameOf(ir.functions[2]).Name(), "ir_func");
} }
#endif // TINT_BUILD_IR #endif // TINT_BUILD_IR

9
src/tint/writer/spirv/ir/generator_impl_ir.cc
View File

@ -223,10 +223,10 @@ uint32_t GeneratorImplIr::Label(const ir::Block* block) {
void GeneratorImplIr::EmitFunction(const ir::Function* func) { void GeneratorImplIr::EmitFunction(const ir::Function* func) {
// Make an ID for the function. // Make an ID for the function.
auto id = module_.NextId(); auto id = module_.NextId();
functions_.Add(func->Name(), id); values_.Add(func, id);
// Emit the function name. // Emit the function name.
module_.PushDebug(spv::Op::OpName, {id, Operand(func->Name().Name())}); module_.PushDebug(spv::Op::OpName, {id, Operand(ir_->NameOf(func).Name())});
// Emit OpEntryPoint and OpExecutionMode declarations if needed. // Emit OpEntryPoint and OpExecutionMode declarations if needed.
if (func->Stage() != ir::Function::PipelineStage::kUndefined) { if (func->Stage() != ir::Function::PipelineStage::kUndefined) {
@ -306,7 +306,8 @@ void GeneratorImplIr::EmitEntryPoint(const ir::Function* func, uint32_t id) {
} }
// TODO(jrprice): Add the interface list of all referenced global variables. // TODO(jrprice): Add the interface list of all referenced global variables.
module_.PushEntryPoint(spv::Op::OpEntryPoint, {U32Operand(stage), id, func->Name().Name()}); module_.PushEntryPoint(spv::Op::OpEntryPoint,
{U32Operand(stage), id, ir_->NameOf(func).Name()});
} }
void GeneratorImplIr::EmitBlock(const ir::Block* block) { void GeneratorImplIr::EmitBlock(const ir::Block* block) {
@ -525,7 +526,7 @@ void GeneratorImplIr::EmitStore(const ir::Store* store) {
uint32_t GeneratorImplIr::EmitUserCall(const ir::UserCall* call) { uint32_t GeneratorImplIr::EmitUserCall(const ir::UserCall* call) {
auto id = module_.NextId(); auto id = module_.NextId();
OperandList operands = {Type(call->Type()), id, functions_.Get(call->Name()).value()}; OperandList operands = {Type(call->Type()), id, values_.Get(call->Func()).value()};
for (auto* arg : call->Args()) { for (auto* arg : call->Args()) {
operands.push_back(Value(arg)); operands.push_back(Value(arg));
} }

4
src/tint/writer/spirv/ir/generator_impl_ir.h
View File

@ -178,10 +178,6 @@ class GeneratorImplIr {
/// The map of constants to their result IDs. /// The map of constants to their result IDs.
utils::Hashmap<const constant::Value*, uint32_t, 16> constants_; utils::Hashmap<const constant::Value*, uint32_t, 16> constants_;
/// The map of functions to their result IDs.
/// TODO(jrprice): Merge into `values_` map when `ir::Function` becomes an `ir::Value`.
utils::Hashmap<Symbol, uint32_t, 8> functions_;
/// The map of non-constant values to their result IDs. /// The map of non-constant values to their result IDs.
utils::Hashmap<const ir::Value*, uint32_t, 8> values_; utils::Hashmap<const ir::Value*, uint32_t, 8> values_;

12
src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
View File

@ -196,10 +196,9 @@ TEST_F(SpvGeneratorImplTest, Function_Call) {
utils::Vector{result, b.Branch(foo->EndTarget(), utils::Vector{result})}); utils::Vector{result, b.Branch(foo->EndTarget(), utils::Vector{result})});
auto* bar = b.CreateFunction("bar", mod.Types().void_()); auto* bar = b.CreateFunction("bar", mod.Types().void_());
bar->StartTarget()->SetInstructions( bar->StartTarget()->SetInstructions(utils::Vector{
utils::Vector{b.UserCall(i32_ty, mod.symbols.Get("foo"), b.UserCall(i32_ty, foo, utils::Vector{b.Constant(i32(2)), b.Constant(i32(3))}),
utils::Vector{b.Constant(i32(2)), b.Constant(i32(3))}), b.Branch(bar->EndTarget())});
b.Branch(bar->EndTarget())});
generator_.EmitFunction(foo); generator_.EmitFunction(foo);
generator_.EmitFunction(bar); generator_.EmitFunction(bar);
@ -231,9 +230,8 @@ TEST_F(SpvGeneratorImplTest, Function_Call_Void) {
foo->StartTarget()->SetInstructions(utils::Vector{b.Branch(foo->EndTarget())}); foo->StartTarget()->SetInstructions(utils::Vector{b.Branch(foo->EndTarget())});
auto* bar = b.CreateFunction("bar", mod.Types().void_()); auto* bar = b.CreateFunction("bar", mod.Types().void_());
bar->StartTarget()->SetInstructions( bar->StartTarget()->SetInstructions(utils::Vector{
utils::Vector{b.UserCall(mod.Types().void_(), mod.symbols.Get("foo"), utils::Empty), b.UserCall(mod.Types().void_(), foo, utils::Empty), b.Branch(bar->EndTarget())});
b.Branch(bar->EndTarget())});
generator_.EmitFunction(foo); generator_.EmitFunction(foo);
generator_.EmitFunction(bar); generator_.EmitFunction(bar);