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 disassembly output. 

This CL switches the disassembler back to being recursive and fixes up
the indenting of blocks to better highlight if branches and switch
cases.

Bug: tint:1718
Change-Id: I14e8d7c68a083bf3afd9ff7241d49b2aa76be4ba
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/134300
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
This commit is contained in:
dan sinclair 2023-05-24 22:43:18 +00:00 committed by Dawn LUCI CQ
parent b3e680bb88
commit a518707db4
10 changed files with 1324 additions and 1140 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

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

@ -89,111 +89,106 @@ std::string_view Disassembler::IdOf(const Value* value) {
std::string Disassembler::Disassemble() {
if (mod_.root_block) {
walk_list_.push_back(mod_.root_block);
Walk();
TINT_ASSERT(IR, walk_list_.empty());
Indent() << "# Root block" << std::endl;
Walk(mod_.root_block);
Walk(mod_.root_block->Branch()->To());
}
for (auto* func : mod_.functions) {
walk_list_.push_back(func);
Walk();
TINT_ASSERT(IR, walk_list_.empty());
EmitFunction(func);
}
return out_.str();
}
void Disassembler::Walk() {
utils::Hashset<const FlowNode*, 32> visited_;
void Disassembler::Walk(const Block* blk) {
if (visited_.Contains(blk)) {
return;
}
visited_.Add(blk);
while (!walk_list_.empty()) {
const FlowNode* node = walk_list_.front();
walk_list_.pop_front();
tint::Switch(
blk,
[&](const ir::FunctionTerminator* t) {
TINT_ASSERT(IR, in_function_);
Indent() << "%fn" << IdOf(t) << " = func_terminator" << std::endl;
in_function_ = false;
},
[&](const ir::RootTerminator* t) {
TINT_ASSERT(IR, !in_function_);
Indent() << "%fn" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
},
[&](const ir::Block* b) {
// If this block is dead, nothing to do
if (!b->HasBranchTarget()) {
return;
}
if (visited_.Contains(node)) {
continue;
}
visited_.Add(node);
tint::Switch(
node,
[&](const ir::Function* f) {
in_function_ = true;
Indent() << "%fn" << IdOf(f) << " = func " << f->Name().Name() << "(";
for (auto* p : f->Params()) {
if (p != f->Params().Front()) {
Indent() << "%fn" << IdOf(b) << " = block";
if (!b->Params().IsEmpty()) {
out_ << " (";
for (auto* p : b->Params()) {
if (p != b->Params().Front()) {
out_ << ", ";
}
out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
EmitValue(p);
}
out_ << "):" << f->ReturnType()->FriendlyName();
out_ << ")";
}
if (f->Stage() != Function::PipelineStage::kUndefined) {
out_ << " [@" << f->Stage();
out_ << " {" << std::endl;
{
ScopedIndent si(indent_size_);
EmitBlockInstructions(b);
}
Indent() << "}" << std::endl;
if (f->WorkgroupSize()) {
auto arr = f->WorkgroupSize().value();
out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2]
<< ")";
}
if (!b->Branch()->To()->Is<FunctionTerminator>()) {
out_ << std::endl;
}
});
}
if (!f->ReturnAttributes().IsEmpty()) {
out_ << " ra:";
void Disassembler::EmitFunction(const Function* func) {
in_function_ = true;
for (auto attr : f->ReturnAttributes()) {
out_ << " @" << attr;
if (attr == Function::ReturnAttribute::kLocation) {
out_ << "(" << f->ReturnLocation().value() << ")";
}
}
}
out_ << "]";
}
out_ << " -> %fn" << IdOf(f->StartTarget()) << std::endl;
walk_list_.push_back(f->StartTarget());
},
[&](const ir::FunctionTerminator* t) {
TINT_ASSERT(IR, in_function_);
Indent() << "%fn" << IdOf(t) << " = func_terminator" << std::endl << std::endl;
in_function_ = false;
},
[&](const ir::RootTerminator* t) {
TINT_ASSERT(IR, !in_function_);
Indent() << "%fn" << IdOf(t) << " = root_terminator" << std::endl << std::endl;
},
[&](const ir::Block* b) {
// If this block is dead, nothing to do
if (!b->HasBranchTarget()) {
return;
}
Indent() << "%fn" << IdOf(b) << " = block";
if (!b->Params().IsEmpty()) {
out_ << " (";
for (auto* p : b->Params()) {
if (p != b->Params().Front()) {
out_ << ", ";
}
EmitValue(p);
}
out_ << ")";
}
out_ << " {" << std::endl;
{
ScopedIndent si(indent_size_);
EmitBlockInstructions(b);
}
Indent() << "}" << std::endl;
if (!b->Branch()->To()->Is<FunctionTerminator>()) {
out_ << std::endl;
}
walk_list_.push_back(b->Branch()->To());
});
Indent() << "%fn" << IdOf(func) << " = func " << func->Name().Name() << "(";
for (auto* p : func->Params()) {
if (p != func->Params().Front()) {
out_ << ", ";
}
out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
}
out_ << "):" << func->ReturnType()->FriendlyName();
if (func->Stage() != Function::PipelineStage::kUndefined) {
out_ << " [@" << func->Stage();
if (func->WorkgroupSize()) {
auto arr = func->WorkgroupSize().value();
out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2] << ")";
}
if (!func->ReturnAttributes().IsEmpty()) {
out_ << " ra:";
for (auto attr : func->ReturnAttributes()) {
out_ << " @" << attr;
if (attr == Function::ReturnAttribute::kLocation) {
out_ << "(" << func->ReturnLocation().value() << ")";
}
}
}
out_ << "]";
}
out_ << " -> %fn" << IdOf(func->StartTarget()) << " {" << std::endl;
{
ScopedIndent si(indent_size_);
Walk(func->StartTarget());
Walk(func->EndTarget());
}
Indent() << "}" << std::endl;
}
void Disassembler::EmitValueWithType(const Value* val) {
@ -335,16 +330,21 @@ void Disassembler::EmitIf(const If* i) {
if (i->Merge()->IsConnected()) {
out_ << ", m: %fn" << IdOf(i->Merge());
}
out_ << "]";
out_ << "]" << std::endl;
if (has_true) {
walk_list_.push_back(i->True());
ScopedIndent si(indent_size_);
Indent() << "# True block" << std::endl;
Walk(i->True());
}
if (has_false) {
walk_list_.push_back(i->False());
ScopedIndent si(indent_size_);
Indent() << "# False block" << std::endl;
Walk(i->False());
}
if (i->Merge()->IsConnected()) {
walk_list_.push_back(i->Merge());
Indent() << "# Merge block" << std::endl;
Walk(i->Merge());
}
}
@ -357,15 +357,21 @@ void Disassembler::EmitLoop(const Loop* l) {
if (l->Merge()->IsConnected()) {
out_ << ", m: %fn" << IdOf(l->Merge());
}
out_ << "]";
out_ << "]" << std::endl;
{ walk_list_.push_back(l->Start()); }
{
ScopedIndent si(indent_size_);
Walk(l->Start());
}
if (l->Continuing()->IsConnected()) {
walk_list_.push_back(l->Continuing());
ScopedIndent si(indent_size_);
Indent() << "# Continuing block" << std::endl;
Walk(l->Continuing());
}
if (l->Merge()->IsConnected()) {
walk_list_.push_back(l->Merge());
Indent() << "# Merge block" << std::endl;
Walk(l->Merge());
}
}
@ -394,13 +400,16 @@ void Disassembler::EmitSwitch(const Switch* s) {
if (s->Merge()->IsConnected()) {
out_ << ", m: %fn" << IdOf(s->Merge());
}
out_ << "]";
out_ << "]" << std::endl;
for (auto& c : s->Cases()) {
walk_list_.push_back(c.Start());
ScopedIndent si(indent_size_);
Indent() << "# Case block" << std::endl;
Walk(c.Start());
}
if (s->Merge()->IsConnected()) {
walk_list_.push_back(s->Merge());
Indent() << "# Merge block" << std::endl;
Walk(s->Merge());
}
}

7
src/tint/ir/disassembler.h
View File

@ -15,7 +15,6 @@
#ifndef SRC_TINT_IR_DISASSEMBLER_H_
#define SRC_TINT_IR_DISASSEMBLER_H_
#include <deque>
#include <string>
#include "src/tint/ir/binary.h"
@ -27,6 +26,7 @@
#include "src/tint/ir/switch.h"
#include "src/tint/ir/unary.h"
#include "src/tint/utils/hashmap.h"
#include "src/tint/utils/hashset.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -56,7 +56,8 @@ class Disassembler {
size_t IdOf(const FlowNode* node);
std::string_view IdOf(const Value* node);
void Walk();
void Walk(const Block* blk);
void EmitFunction(const Function* func);
void EmitInstruction(const Instruction* inst);
void EmitValueWithType(const Value* val);
void EmitValue(const Value* val);
@ -70,7 +71,7 @@ class Disassembler {
const Module& mod_;
utils::StringStream out_;
std::deque<const FlowNode*> walk_list_;
utils::Hashset<const Block*, 32> visited_;
utils::Hashmap<const FlowNode*, size_t, 32> flow_node_ids_;
utils::Hashmap<const Value*, std::string, 32> value_ids_;
uint32_t indent_size_ = 0;

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

File diff suppressed because it is too large Load Diff

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

@ -35,20 +35,20 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 -> %fn2
%fn2 = block {
br %fn3 0.0f # return
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():f32 -> %fn2 {
%fn2 = block {
br %fn3 0.0f # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
%fn5 = block {
%1:f32 = call my_func
%tint_symbol:f32 = bitcast %1
br %fn6 # return
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 {
%fn5 = block {
%1:f32 = call my_func
%tint_symbol:f32 = bitcast %1
br %fn6 # return
}
%fn6 = func_terminator
}
%fn6 = func_terminator
)");
}
@ -62,13 +62,13 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Discard) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():void [@fragment] -> %fn2
%fn2 = block {
discard
br %fn3 # return
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():void [@fragment] -> %fn2 {
%fn2 = block {
discard
br %fn3 # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
)");
}
@ -80,19 +80,19 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void -> %fn2
%fn2 = block {
br %fn3 # return
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void -> %fn2 {
%fn2 = block {
br %fn3 # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
%fn5 = block {
%2:void = call my_func, 6.0f
br %fn6 # return
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 {
%fn5 = block {
%2:void = call my_func, 6.0f
br %fn6 # return
}
%fn6 = func_terminator
}
%fn6 = func_terminator
)");
}
@ -104,21 +104,22 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%i:ptr<private, i32, read_write> = var, 1i
br %fn2 # root_end
}
%fn2 = root_terminator
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
%fn4 = block {
%2:i32 = load %i
%tint_symbol:f32 = convert i32, %2
br %fn5 # return
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 {
%fn4 = block {
%2:i32 = load %i
%tint_symbol:f32 = convert i32, %2
br %fn5 # return
}
%fn5 = func_terminator
}
%fn5 = func_terminator
)");
}
@ -129,7 +130,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%i:ptr<private, vec3<f32>, read_write> = var, vec3<f32> 0.0f
br %fn2 # root_end
}
@ -147,21 +149,22 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%i:ptr<private, f32, read_write> = var, 1.0f
br %fn2 # root_end
}
%fn2 = root_terminator
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
%fn4 = block {
%2:f32 = load %i
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %2
br %fn5 # return
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 {
%fn4 = block {
%2:f32 = load %i
%tint_symbol:vec3<f32> = construct 2.0f, 3.0f, %2
br %fn5 # return
}
%fn5 = func_terminator
}
%fn5 = func_terminator
)");
}

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

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

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

@ -35,20 +35,21 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%a:ptr<private, u32, read_write> = var
br %fn2 # root_end
}
%fn2 = root_terminator
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
%fn4 = block {
store %a, 4u
br %fn5 # return
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 {
%fn4 = block {
store %a, 4u
br %fn5 # return
}
%fn5 = func_terminator
}
%fn5 = func_terminator
)");
}

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

File diff suppressed because it is too large Load Diff

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

@ -34,20 +34,20 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Not) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool -> %fn2
%fn2 = block {
br %fn3 false # return
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool -> %fn2 {
%fn2 = block {
br %fn3 false # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
%fn5 = block {
%1:bool = call my_func
%tint_symbol:bool = eq %1, false
br %fn6 # return
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 {
%fn5 = block {
%1:bool = call my_func
%tint_symbol:bool = eq %1, false
br %fn6 # return
}
%fn6 = func_terminator
}
%fn6 = func_terminator
)");
}
@ -59,20 +59,20 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Complement) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2
%fn2 = block {
br %fn3 1u # return
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():u32 -> %fn2 {
%fn2 = block {
br %fn3 1u # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
%fn5 = block {
%1:u32 = call my_func
%tint_symbol:u32 = complement %1
br %fn6 # return
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 {
%fn5 = block {
%1:u32 = call my_func
%tint_symbol:u32 = complement %1
br %fn6 # return
}
%fn6 = func_terminator
}
%fn6 = func_terminator
)");
}
@ -84,20 +84,20 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Negation) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():i32 -> %fn2
%fn2 = block {
br %fn3 1i # return
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():i32 -> %fn2 {
%fn2 = block {
br %fn3 1i # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5
%fn5 = block {
%1:i32 = call my_func
%tint_symbol:i32 = negation %1
br %fn6 # return
%fn4 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn5 {
%fn5 = block {
%1:i32 = call my_func
%tint_symbol:i32 = negation %1
br %fn6 # return
}
%fn6 = func_terminator
}
%fn6 = func_terminator
)");
}
@ -110,19 +110,20 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%v2:ptr<private, i32, read_write> = var
br %fn2 # root_end
}
%fn2 = root_terminator
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
%fn4 = block {
br %fn5 # return
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 {
%fn4 = block {
br %fn5 # return
}
%fn5 = func_terminator
}
%fn5 = func_terminator
)");
}
@ -137,20 +138,21 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%v3:ptr<private, i32, read_write> = var
br %fn2 # root_end
}
%fn2 = root_terminator
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4
%fn4 = block {
store %v3, 42i
br %fn5 # return
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn4 {
%fn4 = block {
store %v3, 42i
br %fn5 # return
}
%fn5 = func_terminator
}
%fn5 = func_terminator
)");
}

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

@ -32,7 +32,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%a:ptr<private, u32, read_write> = var
br %fn2 # root_end
}
@ -49,7 +50,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = block {
EXPECT_EQ(Disassemble(m.Get()), R"(# Root block
%fn1 = block {
%a:ptr<private, u32, read_write> = var, 2u
br %fn2 # root_end
}
@ -67,13 +69,13 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2
%fn2 = block {
%a:ptr<function, u32, read_write> = var
br %fn3 # return
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 {
%fn2 = block {
%a:ptr<function, u32, read_write> = var
br %fn3 # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
)");
}
@ -86,13 +88,13 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2
%fn2 = block {
%a:ptr<function, u32, read_write> = var, 2u
br %fn3 # return
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] -> %fn2 {
%fn2 = block {
%a:ptr<function, u32, read_write> = var, 2u
br %fn3 # return
}
%fn3 = func_terminator
}
%fn3 = func_terminator
)");
}
} // namespace

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

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