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 type display in disassembly, remove string methods. 

This CL consolidates the various ToInstruction and ToValue methods into
the disassembler. The type output is updated to prefix `:` instead of
surrounding by `()`.

Bug: tint:1718
Change-Id: I69e2d96ffbfe2113932740ce69d0967d29d41541
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131460
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
This commit is contained in:
dan sinclair 2023-05-04 13:34:06 +00:00 committed by Dawn LUCI CQ
parent b298b6a222
commit e964f5163c
36 changed files with 379 additions and 543 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

59
src/tint/ir/binary.cc
View File

@ -29,63 +29,4 @@ Binary::Binary(uint32_t id, Kind kind, const type::Type* ty, Value* lhs, Value*
Binary::~Binary() = default;
utils::StringStream& Binary::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = ";
switch (GetKind()) {
case Binary::Kind::kAdd:
out << "add";
break;
case Binary::Kind::kSubtract:
out << "sub";
break;
case Binary::Kind::kMultiply:
out << "mul";
break;
case Binary::Kind::kDivide:
out << "div";
break;
case Binary::Kind::kModulo:
out << "mod";
break;
case Binary::Kind::kAnd:
out << "and";
break;
case Binary::Kind::kOr:
out << "or";
break;
case Binary::Kind::kXor:
out << "xor";
break;
case Binary::Kind::kEqual:
out << "eq";
break;
case Binary::Kind::kNotEqual:
out << "neq";
break;
case Binary::Kind::kLessThan:
out << "lt";
break;
case Binary::Kind::kGreaterThan:
out << "gt";
break;
case Binary::Kind::kLessThanEqual:
out << "lte";
break;
case Binary::Kind::kGreaterThanEqual:
out << "gte";
break;
case Binary::Kind::kShiftLeft:
out << "shiftl";
break;
case Binary::Kind::kShiftRight:
out << "shiftr";
break;
}
out << " ";
lhs_->ToValue(out) << ", ";
rhs_->ToValue(out);
return out;
}
} // namespace tint::ir

6
src/tint/ir/binary.h
View File

@ -17,7 +17,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -70,11 +69,6 @@ class Binary : public utils::Castable<Binary, Instruction> {
/// @returns the right-hand-side value for the instruction
const Value* RHS() const { return rhs_; }
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
private:
Kind kind_;
Value* lhs_ = nullptr;

65
src/tint/ir/binary_test.cc
View File

@ -14,7 +14,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/test_helper.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
namespace {
@ -42,10 +41,6 @@ TEST_F(IR_InstructionTest, CreateAnd) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = and 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateOr) {
@ -66,10 +61,6 @@ TEST_F(IR_InstructionTest, CreateOr) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = or 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateXor) {
@ -90,10 +81,6 @@ TEST_F(IR_InstructionTest, CreateXor) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = xor 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateEqual) {
@ -114,10 +101,6 @@ TEST_F(IR_InstructionTest, CreateEqual) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(bool) = eq 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateNotEqual) {
@ -138,10 +121,6 @@ TEST_F(IR_InstructionTest, CreateNotEqual) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(bool) = neq 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateLessThan) {
@ -162,10 +141,6 @@ TEST_F(IR_InstructionTest, CreateLessThan) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(bool) = lt 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateGreaterThan) {
@ -186,10 +161,6 @@ TEST_F(IR_InstructionTest, CreateGreaterThan) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(bool) = gt 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateLessThanEqual) {
@ -210,10 +181,6 @@ TEST_F(IR_InstructionTest, CreateLessThanEqual) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(bool) = lte 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateGreaterThanEqual) {
@ -234,10 +201,6 @@ TEST_F(IR_InstructionTest, CreateGreaterThanEqual) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(bool) = gte 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateShiftLeft) {
@ -258,10 +221,6 @@ TEST_F(IR_InstructionTest, CreateShiftLeft) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = shiftl 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateShiftRight) {
@ -282,10 +241,6 @@ TEST_F(IR_InstructionTest, CreateShiftRight) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = shiftr 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateAdd) {
@ -306,10 +261,6 @@ TEST_F(IR_InstructionTest, CreateAdd) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = add 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateSubtract) {
@ -330,10 +281,6 @@ TEST_F(IR_InstructionTest, CreateSubtract) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = sub 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateMultiply) {
@ -354,10 +301,6 @@ TEST_F(IR_InstructionTest, CreateMultiply) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = mul 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateDivide) {
@ -378,10 +321,6 @@ TEST_F(IR_InstructionTest, CreateDivide) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = div 4i, 2i");
}
TEST_F(IR_InstructionTest, CreateModulo) {
@ -402,10 +341,6 @@ TEST_F(IR_InstructionTest, CreateModulo) {
auto rhs = inst->RHS()->As<Constant>()->value;
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = mod 4i, 2i");
}
TEST_F(IR_InstructionTest, Binary_Usage) {

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

@ -24,10 +24,4 @@ Bitcast::Bitcast(uint32_t id, const type::Type* type, Value* val)
Bitcast::~Bitcast() = default;
utils::StringStream& Bitcast::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = bitcast ";
EmitArgs(out);
return out;
}
} // namespace tint::ir

6
src/tint/ir/bitcast.h
View File

@ -17,7 +17,6 @@
#include "src/tint/ir/call.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -35,11 +34,6 @@ class Bitcast : public utils::Castable<Bitcast, Call> {
Bitcast& operator=(const Bitcast& inst) = delete;
Bitcast& operator=(Bitcast&& inst) = delete;
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
};
} // namespace tint::ir

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

@ -14,7 +14,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/test_helper.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
namespace {
@ -36,10 +35,6 @@ TEST_F(IR_InstructionTest, Bitcast) {
auto val = inst->Args()[0]->As<Constant>()->value;
ASSERT_TRUE(val->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, val->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = bitcast 4i");
}
TEST_F(IR_InstructionTest, Bitcast_Usage) {

261
src/tint/ir/builder_impl_test.cc
View File

@ -42,7 +42,7 @@ TEST_F(IR_BuilderImplTest, Func) {
EXPECT_EQ(1u, f->start_target->inbound_branches.Length());
EXPECT_EQ(1u, f->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func f(void)
EXPECT_EQ(Disassemble(m), R"(%fn0 = func f():void
%fn1 = block
ret
func_end
@ -88,7 +88,8 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -136,7 +137,8 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -183,7 +185,8 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -230,7 +233,8 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -273,7 +277,8 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
ASSERT_NE(loop_flow->continuing.target, nullptr);
ASSERT_NE(loop_flow->merge.target, nullptr);
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -330,7 +335,8 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -388,7 +394,8 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -463,7 +470,8 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -538,7 +546,8 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -595,7 +604,8 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -657,7 +667,8 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
// This is 1 because only the loop branch happens. The subsequent if return is dead code.
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -711,7 +722,8 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -857,7 +869,8 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1001,7 +1014,8 @@ TEST_F(IR_BuilderImplTest, While) {
EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1071,7 +1085,8 @@ TEST_F(IR_BuilderImplTest, While_Return) {
EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1178,7 +1193,8 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1237,7 +1253,8 @@ TEST_F(IR_BuilderImplTest, Switch) {
EXPECT_EQ(3u, flow->merge.target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1301,7 +1318,8 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1345,7 +1363,8 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1398,7 +1417,8 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
// This is 1 because the if is dead-code eliminated and the return doesn't happen.
EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1457,7 +1477,8 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length());
EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
branch %fn2
@ -1566,7 +1587,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1(ref<private, u32, read_write>) = var private read_write
%1:ref<private, u32, read_write> = var private read_write
@ -1582,8 +1603,8 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1(ref<private, u32, read_write>) = var private read_write
store %1(ref<private, u32, read_write>), 2u
%1:ref<private, u32, read_write> = var private read_write
store %1:ref<private, u32, read_write>, 2u
@ -1598,9 +1619,10 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
%1(ref<function, u32, read_write>) = var function read_write
%1:ref<function, u32, read_write> = var function read_write
ret
func_end
@ -1616,10 +1638,11 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
EXPECT_EQ(Disassemble(m),
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn1 = block
%1(ref<function, u32, read_write>) = var function read_write
store %1(ref<function, u32, read_write>), 2u
%1:ref<function, u32, read_write> = var function read_write
store %1:ref<function, u32, read_write>, 2u
ret
func_end
@ -1639,8 +1662,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Add) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = add %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = add %1:u32, 4u
)");
}
@ -1657,8 +1680,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Subtract) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = sub %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = sub %1:u32, 4u
)");
}
@ -1675,8 +1698,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Multiply) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = mul %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = mul %1:u32, 4u
)");
}
@ -1693,8 +1716,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Div) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = div %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = div %1:u32, 4u
)");
}
@ -1711,8 +1734,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Modulo) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = mod %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = mod %1:u32, 4u
)");
}
@ -1729,8 +1752,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_And) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = and %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = and %1:u32, 4u
)");
}
@ -1747,8 +1770,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Or) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = or %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = or %1:u32, 4u
)");
}
@ -1765,8 +1788,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Xor) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = xor %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = xor %1:u32, 4u
)");
}
@ -1779,22 +1802,22 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalAnd) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool)
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1(bool) = call my_func
%2(bool) = var function read_write
store %2(bool), %1(bool)
%1:bool = call my_func
%2:bool = var function read_write
store %2:bool, %1:bool
branch %fn4
%fn4 = if %1(bool) [t: %fn5, f: %fn6, m: %fn7]
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
%fn5 = block
store %2(bool), false
store %2:bool, false
branch %fn7
# if merge
@ -1814,23 +1837,23 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LogicalOr) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool)
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1(bool) = call my_func
%2(bool) = var function read_write
store %2(bool), %1(bool)
%1:bool = call my_func
%2:bool = var function read_write
store %2:bool, %1:bool
branch %fn4
%fn4 = if %1(bool) [t: %fn5, f: %fn6, m: %fn7]
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
# false branch
%fn6 = block
store %2(bool), true
store %2:bool, true
branch %fn7
# if merge
@ -1854,8 +1877,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Equal) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(bool) = eq %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = eq %1:u32, 4u
)");
}
@ -1872,8 +1895,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_NotEqual) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(bool) = neq %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = neq %1:u32, 4u
)");
}
@ -1890,8 +1913,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThan) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(bool) = lt %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = lt %1:u32, 4u
)");
}
@ -1908,8 +1931,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThan) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(bool) = gt %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = gt %1:u32, 4u
)");
}
@ -1926,8 +1949,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_LessThanEqual) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(bool) = lte %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = lte %1:u32, 4u
)");
}
@ -1944,8 +1967,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_GreaterThanEqual) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(bool) = gte %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:bool = gte %1:u32, 4u
)");
}
@ -1962,8 +1985,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftLeft) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = shiftl %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = shiftl %1:u32, 4u
)");
}
@ -1980,8 +2003,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_ShiftRight) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(u32) = call my_func
%2(u32) = shiftr %1(u32), 4u
EXPECT_EQ(d.AsString(), R"(%1:u32 = call my_func
%2:u32 = shiftr %1:u32, 4u
)");
}
@ -1996,28 +2019,28 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(f32)
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():f32
%fn1 = block
ret 0.0f
func_end
%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1(f32) = call my_func
%2(bool) = lt %1(f32), 2.0f
%3(bool) = var function read_write
store %3(bool), %2(bool)
%1:f32 = call my_func
%2:bool = lt %1:f32, 2.0f
%3:bool = var function read_write
store %3:bool, %2:bool
branch %fn4
%fn4 = if %2(bool) [t: %fn5, f: %fn6, m: %fn7]
%fn4 = if %2:bool [t: %fn5, f: %fn6, m: %fn7]
# true branch
%fn5 = block
%4(f32) = call my_func
%5(f32) = call my_func
%6(f32) = mul 2.29999995231628417969f, %5(f32)
%7(f32) = div %4(f32), %6(f32)
%8(bool) = gt 2.5f, %7(f32)
store %3(bool), %8(bool)
%4:f32 = call my_func
%5:f32 = call my_func
%6:f32 = mul 2.29999995231628417969f, %5:f32
%7:f32 = div %4:f32, %6:f32
%8:bool = gt 2.5f, %7:f32
store %3:bool, %8:bool
branch %fn7
# if merge
@ -2038,14 +2061,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_Compound_WithConstEval) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func(bool)
EXPECT_EQ(Disassemble(m), R"(%fn0 = func my_func():bool
%fn1 = block
ret true
func_end
%fn2 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn3 = block
%1(bool) = call my_func, false
%1:bool = call my_func, false
ret
func_end
@ -2066,8 +2089,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Bitcast) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(f32) = call my_func
%2(f32) = bitcast %1(f32)
EXPECT_EQ(d.AsString(), R"(%1:f32 = call my_func
%2:f32 = bitcast %1:f32
)");
}
@ -2103,7 +2126,7 @@ TEST_F(IR_BuilderImplTest, EmitStatement_UserFunction) {
Disassembler d(b.builder.ir);
d.EmitBlockInstructions(b.current_flow_block->As<ir::Block>());
EXPECT_EQ(d.AsString(), R"(%1(void) = call my_func, 6.0f
EXPECT_EQ(d.AsString(), R"(%1:void = call my_func, 6.0f
)");
}
@ -2117,8 +2140,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1(ref<private, vec3<f32>, read_write>) = var private read_write
store %1(ref<private, vec3<f32>, read_write>), vec3<f32> 0.0f
%1:ref<private, vec3<f32>, read_write> = var private read_write
store %1:ref<private, vec3<f32>, read_write>, vec3<f32> 0.0f
@ -2136,14 +2159,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1(ref<private, f32, read_write>) = var private read_write
store %1(ref<private, f32, read_write>), 1.0f
%1:ref<private, f32, read_write> = var private read_write
store %1:ref<private, f32, read_write>, 1.0f
%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2(vec3<f32>) = construct 2.0f, 3.0f, %1(ref<private, f32, read_write>)
%2:vec3<f32> = construct 2.0f, 3.0f, %1:ref<private, f32, read_write>
ret
func_end
@ -2161,14 +2184,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
ASSERT_TRUE(r);
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1(ref<private, i32, read_write>) = var private read_write
store %1(ref<private, i32, read_write>), 1i
%1:ref<private, i32, read_write> = var private read_write
store %1:ref<private, i32, read_write>, 1i
%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2(f32) = convert i32, %1(ref<private, i32, read_write>)
%2:f32 = convert i32, %1:ref<private, i32, read_write>
ret
func_end
@ -2184,7 +2207,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function(f32)
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function():f32
%fn1 = block
ret 2.0f
func_end
@ -2202,14 +2225,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Builtin) {
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
%1(ref<private, f32, read_write>) = var private read_write
store %1(ref<private, f32, read_write>), 1.0f
%1:ref<private, f32, read_write> = var private read_write
store %1:ref<private, f32, read_write>, 1.0f
%fn1 = func test_function(void) [@compute @workgroup_size(1, 1, 1)]
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
%fn2 = block
%2(f32) = asin %1(ref<private, f32, read_write>)
%2:f32 = asin %1:ref<private, f32, read_write>
ret
func_end
@ -2225,7 +2248,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Vertex) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@vertex ra: @position]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position]
%fn1 = block
ret vec4<f32> 0.0f
func_end
@ -2241,7 +2264,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Fragment) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(void) [@fragment]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():void [@fragment]
%fn1 = block
ret
func_end
@ -2257,7 +2280,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Compute) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(void) [@compute @workgroup_size(8, 4, 2)]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():void [@compute @workgroup_size(8, 4, 2)]
%fn1 = block
ret
func_end
@ -2273,7 +2296,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_Return) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec3<f32>)
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec3<f32>
%fn1 = block
ret vec3<f32> 0.0f
func_end
@ -2290,7 +2313,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPosition) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@vertex ra: @position]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
@ -2307,7 +2330,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnPositionInvariant) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@vertex ra: @position @invariant]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@vertex ra: @position @invariant]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
@ -2323,7 +2346,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnLocation) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(vec4<f32>) [@fragment ra: @location(1)]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():vec4<f32> [@fragment ra: @location(1)]
%fn1 = block
ret vec4<f32> 1.0f, 2.0f, 3.0f, 4.0f
func_end
@ -2340,7 +2363,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnFragDepth) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(f32) [@fragment ra: @frag_depth]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():f32 [@fragment ra: @frag_depth]
%fn1 = block
ret 1.0f
func_end
@ -2357,7 +2380,7 @@ TEST_F(IR_BuilderImplTest, EmitFunction_ReturnSampleMask) {
ASSERT_TRUE(r) << Error();
auto m = r.Move();
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test(u32) [@fragment ra: @sample_mask]
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test():u32 [@fragment ra: @sample_mask]
%fn1 = block
ret 1u
func_end

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

@ -28,11 +28,5 @@ Builtin::Builtin(uint32_t id,
Builtin::~Builtin() = default;
utils::StringStream& Builtin::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = " << builtin::str(func_) << " ";
EmitArgs(out);
return out;
}
} // namespace tint::ir
// \endcond

6
src/tint/ir/builtin.h
View File

@ -18,7 +18,6 @@
#include "src/tint/builtin/function.h"
#include "src/tint/ir/call.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -44,11 +43,6 @@ class Builtin : public utils::Castable<Builtin, Call> {
/// @returns the builtin function
builtin::Function Func() const { return func_; }
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
private:
const builtin::Function func_;
};

11
src/tint/ir/call.cc
View File

@ -29,15 +29,4 @@ Call::Call(uint32_t id, const type::Type* type, utils::VectorRef<Value*> args)
Call::~Call() = default;
void Call::EmitArgs(utils::StringStream& out) const {
bool first = true;
for (const auto* arg : args_) {
if (!first) {
out << ", ";
}
first = false;
arg->ToValue(out);
}
}
} // namespace tint::ir

5
src/tint/ir/call.h
View File

@ -17,7 +17,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -34,10 +33,6 @@ class Call : public utils::Castable<Call, Instruction> {
/// @returns the constructor arguments
utils::VectorRef<Value*> Args() const { return args_; }
/// Writes the call arguments to the given stream.
/// @param out the output stream
void EmitArgs(utils::StringStream& out) const;
protected:
/// Constructor
Call();

46
src/tint/ir/constant.cc
View File

@ -14,13 +14,6 @@
#include "src/tint/ir/constant.h"
#include <string>
#include "src/tint/constant/composite.h"
#include "src/tint/constant/scalar.h"
#include "src/tint/constant/splat.h"
#include "src/tint/switch.h"
TINT_INSTANTIATE_TYPEINFO(tint::ir::Constant);
namespace tint::ir {
@ -29,43 +22,4 @@ Constant::Constant(const constant::Value* val) : value(val) {}
Constant::~Constant() = default;
utils::StringStream& Constant::ToValue(utils::StringStream& out) const {
std::function<void(const constant::Value*)> emit = [&](const constant::Value* c) {
Switch(
c,
[&](const constant::Scalar<AFloat>* scalar) { out << scalar->ValueAs<AFloat>().value; },
[&](const constant::Scalar<AInt>* scalar) { out << scalar->ValueAs<AInt>().value; },
[&](const constant::Scalar<i32>* scalar) {
out << scalar->ValueAs<i32>().value << "i";
},
[&](const constant::Scalar<u32>* scalar) {
out << scalar->ValueAs<u32>().value << "u";
},
[&](const constant::Scalar<f32>* scalar) {
out << scalar->ValueAs<f32>().value << "f";
},
[&](const constant::Scalar<f16>* scalar) {
out << scalar->ValueAs<f16>().value << "h";
},
[&](const constant::Scalar<bool>* scalar) {
out << (scalar->ValueAs<bool>() ? "true" : "false");
},
[&](const constant::Splat* splat) {
out << splat->Type()->FriendlyName() << " ";
emit(splat->Index(0));
},
[&](const constant::Composite* composite) {
out << composite->Type()->FriendlyName() << " ";
for (const auto* elem : composite->elements) {
if (elem != composite->elements[0]) {
out << ", ";
}
emit(elem);
}
});
};
emit(value);
return out;
}
} // namespace tint::ir

6
src/tint/ir/constant.h
View File

@ -17,7 +17,6 @@
#include "src/tint/constant/value.h"
#include "src/tint/ir/value.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -32,11 +31,6 @@ class Constant : public utils::Castable<Constant, Value> {
/// @returns the type of the constant
const type::Type* Type() const override { return value->Type(); }
/// Write the constant to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToValue(utils::StringStream& out) const override;
/// The constants value
const constant::Value* const value;
};

19
src/tint/ir/constant_test.cc
View File

@ -14,7 +14,6 @@
#include "src/tint/ir/test_helper.h"
#include "src/tint/ir/value.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
namespace {
@ -31,9 +30,6 @@ TEST_F(IR_ConstantTest, f32) {
auto* c = b.builder.Constant(1.2_f);
EXPECT_EQ(1.2_f, c->value->As<constant::Scalar<f32>>()->ValueAs<f32>());
c->ToValue(str);
EXPECT_EQ("1.20000004768371582031f", str.str());
EXPECT_TRUE(c->value->Is<constant::Scalar<f32>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
@ -49,9 +45,6 @@ TEST_F(IR_ConstantTest, f16) {
auto* c = b.builder.Constant(1.1_h);
EXPECT_EQ(1.1_h, c->value->As<constant::Scalar<f16>>()->ValueAs<f16>());
c->ToValue(str);
EXPECT_EQ("1.099609375h", str.str());
EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
EXPECT_TRUE(c->value->Is<constant::Scalar<f16>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
@ -67,9 +60,6 @@ TEST_F(IR_ConstantTest, i32) {
auto* c = b.builder.Constant(1_i);
EXPECT_EQ(1_i, c->value->As<constant::Scalar<i32>>()->ValueAs<i32>());
c->ToValue(str);
EXPECT_EQ("1i", str.str());
EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
EXPECT_TRUE(c->value->Is<constant::Scalar<i32>>());
@ -85,9 +75,6 @@ TEST_F(IR_ConstantTest, u32) {
auto* c = b.builder.Constant(2_u);
EXPECT_EQ(2_u, c->value->As<constant::Scalar<u32>>()->ValueAs<u32>());
c->ToValue(str);
EXPECT_EQ("2u", str.str());
EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
@ -103,9 +90,6 @@ TEST_F(IR_ConstantTest, bool) {
auto* c = b.builder.Constant(false);
EXPECT_FALSE(c->value->As<constant::Scalar<bool>>()->ValueAs<bool>());
c->ToValue(str);
EXPECT_EQ("false", str.str());
}
{
@ -113,9 +97,6 @@ TEST_F(IR_ConstantTest, bool) {
auto c = b.builder.Constant(true);
EXPECT_TRUE(c->value->As<constant::Scalar<bool>>()->ValueAs<bool>());
c->ToValue(str);
EXPECT_EQ("true", str.str());
EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());

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

@ -24,10 +24,4 @@ Construct::Construct(uint32_t id, const type::Type* type, utils::VectorRef<Value
Construct::~Construct() = default;
utils::StringStream& Construct::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = construct ";
EmitArgs(out);
return out;
}
} // namespace tint::ir

6
src/tint/ir/construct.h
View File

@ -17,7 +17,6 @@
#include "src/tint/ir/call.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -35,11 +34,6 @@ class Construct : public utils::Castable<Construct, Call> {
Construct& operator=(const Construct& inst) = delete;
Construct& operator=(Construct&& inst) = delete;
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
};
} // namespace tint::ir

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

@ -27,10 +27,4 @@ Convert::Convert(uint32_t id,
Convert::~Convert() = default;
utils::StringStream& Convert::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = convert " << from_type_->FriendlyName() << ", ";
EmitArgs(out);
return out;
}
} // namespace tint::ir

6
src/tint/ir/convert.h
View File

@ -18,7 +18,6 @@
#include "src/tint/ir/call.h"
#include "src/tint/type/type.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -46,11 +45,6 @@ class Convert : public utils::Castable<Convert, Call> {
/// @returns the to type
const type::Type* ToType() const { return Type(); }
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
private:
const type::Type* from_type_ = nullptr;
};

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

@ -14,12 +14,25 @@
#include "src/tint/ir/disassembler.h"
#include "src//tint/ir/unary.h"
#include "src/tint/constant/composite.h"
#include "src/tint/constant/scalar.h"
#include "src/tint/constant/splat.h"
#include "src/tint/ir/binary.h"
#include "src/tint/ir/bitcast.h"
#include "src/tint/ir/block.h"
#include "src/tint/ir/builtin.h"
#include "src/tint/ir/construct.h"
#include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h"
#include "src/tint/ir/function_terminator.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/loop.h"
#include "src/tint/ir/root_terminator.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/switch.h"
#include "src/tint/ir/user_call.h"
#include "src/tint/ir/var.h"
#include "src/tint/switch.h"
#include "src/tint/type/type.h"
#include "src/tint/utils/scoped_assignment.h"
@ -67,7 +80,8 @@ utils::StringStream& Disassembler::Indent() {
void Disassembler::EmitBlockInstructions(const Block* b) {
for (const auto* inst : b->instructions) {
Indent();
inst->ToInstruction(out_) << std::endl;
EmitInstruction(inst);
out_ << std::endl;
}
}
@ -94,8 +108,8 @@ void Disassembler::Walk(const FlowNode* node) {
[&](const ir::Function* f) {
TINT_SCOPED_ASSIGNMENT(in_function_, true);
Indent() << "%fn" << GetIdForNode(f) << " = func " << f->name.Name() << "("
<< f->return_type->FriendlyName() << ")";
Indent() << "%fn" << GetIdForNode(f) << " = func " << f->name.Name()
<< "():" << f->return_type->FriendlyName();
if (f->pipeline_stage != Function::PipelineStage::kUndefined) {
out_ << " [@" << f->pipeline_stage;
@ -151,7 +165,7 @@ void Disassembler::Walk(const FlowNode* node) {
if (v != b->branch.args.Front()) {
out_ << ", ";
}
v->ToValue(out_);
EmitValue(v);
}
}
out_ << std::endl;
@ -164,7 +178,7 @@ void Disassembler::Walk(const FlowNode* node) {
},
[&](const ir::Switch* s) {
Indent() << "%fn" << GetIdForNode(s) << " = switch ";
s->condition->ToValue(out_);
EmitValue(s->condition);
out_ << " [";
for (const auto& c : s->cases) {
if (&c != &s->cases.Front()) {
@ -179,7 +193,7 @@ void Disassembler::Walk(const FlowNode* node) {
if (selector.IsDefault()) {
out_ << "default";
} else {
selector.val->ToValue(out_);
EmitValue(selector.val);
}
}
out_ << ", %fn" << GetIdForNode(c.start.target) << ")";
@ -202,7 +216,7 @@ void Disassembler::Walk(const FlowNode* node) {
if (selector.IsDefault()) {
out_ << "default";
} else {
selector.val->ToValue(out_);
EmitValue(selector.val);
}
}
out_ << std::endl;
@ -217,7 +231,7 @@ void Disassembler::Walk(const FlowNode* node) {
},
[&](const ir::If* i) {
Indent() << "%fn" << GetIdForNode(i) << " = if ";
i->condition->ToValue(out_);
EmitValue(i->condition);
out_ << " [t: %fn" << GetIdForNode(i->true_.target) << ", f: %fn"
<< GetIdForNode(i->false_.target);
if (i->merge.target->IsConnected()) {
@ -296,4 +310,195 @@ std::string Disassembler::Disassemble() {
return out_.str();
}
void Disassembler::EmitValue(const Value* val) {
tint::Switch(
val,
[&](const ir::Constant* constant) {
std::function<void(const constant::Value*)> emit = [&](const constant::Value* c) {
tint::Switch(
c,
[&](const constant::Scalar<AFloat>* scalar) {
out_ << scalar->ValueAs<AFloat>().value;
},
[&](const constant::Scalar<AInt>* scalar) {
out_ << scalar->ValueAs<AInt>().value;
},
[&](const constant::Scalar<i32>* scalar) {
out_ << scalar->ValueAs<i32>().value << "i";
},
[&](const constant::Scalar<u32>* scalar) {
out_ << scalar->ValueAs<u32>().value << "u";
},
[&](const constant::Scalar<f32>* scalar) {
out_ << scalar->ValueAs<f32>().value << "f";
},
[&](const constant::Scalar<f16>* scalar) {
out_ << scalar->ValueAs<f16>().value << "h";
},
[&](const constant::Scalar<bool>* scalar) {
out_ << (scalar->ValueAs<bool>() ? "true" : "false");
},
[&](const constant::Splat* splat) {
out_ << splat->Type()->FriendlyName() << " ";
emit(splat->Index(0));
},
[&](const constant::Composite* composite) {
out_ << composite->Type()->FriendlyName() << " ";
for (const auto* elem : composite->elements) {
if (elem != composite->elements[0]) {
out_ << ", ";
}
emit(elem);
}
});
};
emit(constant->value);
},
[&](const ir::Instruction* i) {
out_ << "%" << std::to_string(i->Id());
if (i->Type() != nullptr) {
out_ << ":" << i->Type()->FriendlyName();
}
});
}
void Disassembler::EmitInstruction(const Instruction* inst) {
tint::Switch(
inst, //
[&](const ir::Binary* b) { EmitBinary(b); }, [&](const ir::Unary* u) { EmitUnary(u); },
[&](const ir::Bitcast* b) {
EmitValue(b);
out_ << " = bitcast ";
EmitArgs(b);
},
[&](const ir::Discard*) { out_ << "discard"; },
[&](const ir::Builtin* b) {
EmitValue(b);
out_ << " = " << builtin::str(b->Func()) << " ";
EmitArgs(b);
},
[&](const ir::Construct* c) {
EmitValue(c);
out_ << " = construct ";
EmitArgs(c);
},
[&](const ir::Convert* c) {
EmitValue(c);
out_ << " = convert " << c->FromType()->FriendlyName() << ", ";
EmitArgs(c);
},
[&](const ir::Store* s) {
out_ << "store ";
EmitValue(s->To());
out_ << ", ";
EmitValue(s->From());
},
[&](const ir::UserCall* uc) {
EmitValue(uc);
out_ << " = call " << uc->Name().Name();
if (uc->Args().Length() > 0) {
out_ << ", ";
}
EmitArgs(uc);
},
[&](const ir::Var* v) {
EmitValue(v);
out_ << " = var " << v->AddressSpace() << " " << v->Access();
});
}
void Disassembler::EmitArgs(const Call* call) {
bool first = true;
for (const auto* arg : call->Args()) {
if (!first) {
out_ << ", ";
}
first = false;
EmitValue(arg);
}
}
void Disassembler::EmitBinary(const Binary* b) {
EmitValue(b);
out_ << " = ";
switch (b->GetKind()) {
case Binary::Kind::kAdd:
out_ << "add";
break;
case Binary::Kind::kSubtract:
out_ << "sub";
break;
case Binary::Kind::kMultiply:
out_ << "mul";
break;
case Binary::Kind::kDivide:
out_ << "div";
break;
case Binary::Kind::kModulo:
out_ << "mod";
break;
case Binary::Kind::kAnd:
out_ << "and";
break;
case Binary::Kind::kOr:
out_ << "or";
break;
case Binary::Kind::kXor:
out_ << "xor";
break;
case Binary::Kind::kEqual:
out_ << "eq";
break;
case Binary::Kind::kNotEqual:
out_ << "neq";
break;
case Binary::Kind::kLessThan:
out_ << "lt";
break;
case Binary::Kind::kGreaterThan:
out_ << "gt";
break;
case Binary::Kind::kLessThanEqual:
out_ << "lte";
break;
case Binary::Kind::kGreaterThanEqual:
out_ << "gte";
break;
case Binary::Kind::kShiftLeft:
out_ << "shiftl";
break;
case Binary::Kind::kShiftRight:
out_ << "shiftr";
break;
}
out_ << " ";
EmitValue(b->LHS());
out_ << ", ";
EmitValue(b->RHS());
}
void Disassembler::EmitUnary(const Unary* u) {
EmitValue(u);
out_ << " = ";
switch (u->GetKind()) {
case Unary::Kind::kAddressOf:
out_ << "addr_of";
break;
case Unary::Kind::kComplement:
out_ << "complement";
break;
case Unary::Kind::kIndirection:
out_ << "indirection";
break;
case Unary::Kind::kNegation:
out_ << "negation";
break;
case Unary::Kind::kNot:
out_ << "not";
break;
}
out_ << " ";
EmitValue(u->Val());
}
} // namespace tint::ir

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

@ -19,8 +19,11 @@
#include <unordered_map>
#include <unordered_set>
#include "src/tint/ir/binary.h"
#include "src/tint/ir/call.h"
#include "src/tint/ir/flow_node.h"
#include "src/tint/ir/module.h"
#include "src/tint/ir/unary.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -46,9 +49,15 @@ class Disassembler {
private:
utils::StringStream& Indent();
void Walk(const FlowNode* node);
size_t GetIdForNode(const FlowNode* node);
void Walk(const FlowNode* node);
void EmitInstruction(const Instruction* inst);
void EmitValue(const Value* val);
void EmitArgs(const Call* call);
void EmitBinary(const Binary* b);
void EmitUnary(const Unary* b);
const Module& mod_;
utils::StringStream out_;
std::unordered_set<const FlowNode*> visited_;

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

@ -23,9 +23,4 @@ Discard::Discard() : Base() {}
Discard::~Discard() = default;
utils::StringStream& Discard::ToInstruction(utils::StringStream& out) const {
out << "discard";
return out;
}
} // namespace tint::ir

6
src/tint/ir/discard.h
View File

@ -18,7 +18,6 @@
#include "src/tint/debug.h"
#include "src/tint/ir/call.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -33,11 +32,6 @@ class Discard : public utils::Castable<Discard, Call> {
Discard& operator=(const Discard& inst) = delete;
Discard& operator=(Discard&& inst) = delete;
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
};
} // namespace tint::ir

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

@ -14,7 +14,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/test_helper.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
namespace {
@ -26,10 +25,6 @@ TEST_F(IR_InstructionTest, Discard) {
const auto* inst = b.builder.Discard();
ASSERT_TRUE(inst->Is<ir::Discard>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "discard");
}
} // namespace

19
src/tint/ir/instruction.h
View File

@ -17,7 +17,6 @@
#include "src/tint/ir/value.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -32,25 +31,11 @@ class Instruction : public utils::Castable<Instruction, Value> {
Instruction& operator=(const Instruction& inst) = delete;
Instruction& operator=(Instruction&& inst) = delete;
/// @returns the id of the instruction
uint32_t Id() const { return id_; }
/// @returns the type of the value
const type::Type* Type() const override { return type_; }
/// Write the value to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToValue(utils::StringStream& out) const override {
out << "%" << std::to_string(id_);
if (type_ != nullptr) {
out << "(" << Type()->FriendlyName() << ")";
}
return out;
}
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
virtual utils::StringStream& ToInstruction(utils::StringStream& out) const = 0;
protected:
/// Constructor
Instruction();

7
src/tint/ir/store.cc
View File

@ -28,11 +28,4 @@ Store::Store(Value* to, Value* from) : Base(), to_(to), from_(from) {
Store::~Store() = default;
utils::StringStream& Store::ToInstruction(utils::StringStream& out) const {
out << "store ";
to_->ToValue(out) << ", ";
from_->ToValue(out);
return out;
}
} // namespace tint::ir

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

@ -17,7 +17,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -36,14 +35,9 @@ class Store : public utils::Castable<Store, Instruction> {
Store& operator=(Store&& inst) = delete;
/// @returns the value being stored too
const Value* to() const { return to_; }
const Value* To() const { return to_; }
/// @returns the value being stored
const Value* from() const { return from_; }
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
const Value* From() const { return from_; }
private:
Value* to_ = nullptr;

23
src/tint/ir/store_test.cc
View File

@ -14,7 +14,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/test_helper.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
namespace {
@ -32,16 +31,12 @@ TEST_F(IR_InstructionTest, CreateStore) {
const auto* inst = b.builder.Store(to, b.builder.Constant(4_i));
ASSERT_TRUE(inst->Is<Store>());
ASSERT_EQ(inst->to(), to);
ASSERT_EQ(inst->To(), to);
ASSERT_TRUE(inst->from()->Is<Constant>());
auto lhs = inst->from()->As<Constant>()->value;
ASSERT_TRUE(inst->From()->Is<Constant>());
auto lhs = inst->From()->As<Constant>()->value;
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "store %0, 4i");
}
TEST_F(IR_InstructionTest, Store_Usage) {
@ -50,13 +45,13 @@ TEST_F(IR_InstructionTest, Store_Usage) {
auto* to = b.builder.Discard();
const auto* inst = b.builder.Store(to, b.builder.Constant(4_i));
ASSERT_NE(inst->to(), nullptr);
ASSERT_EQ(inst->to()->Usage().Length(), 1u);
EXPECT_EQ(inst->to()->Usage()[0], inst);
ASSERT_NE(inst->To(), nullptr);
ASSERT_EQ(inst->To()->Usage().Length(), 1u);
EXPECT_EQ(inst->To()->Usage()[0], inst);
ASSERT_NE(inst->from(), nullptr);
ASSERT_EQ(inst->from()->Usage().Length(), 1u);
EXPECT_EQ(inst->from()->Usage()[0], inst);
ASSERT_NE(inst->From(), nullptr);
ASSERT_EQ(inst->From()->Usage().Length(), 1u);
EXPECT_EQ(inst->From()->Usage()[0], inst);
}
} // namespace

1
src/tint/ir/test_helper.h
View File

@ -24,6 +24,7 @@
#include "src/tint/ir/disassembler.h"
#include "src/tint/number.h"
#include "src/tint/program_builder.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {

24
src/tint/ir/unary.cc
View File

@ -27,28 +27,4 @@ Unary::Unary(uint32_t id, Kind kind, const type::Type* type, Value* val)
Unary::~Unary() = default;
utils::StringStream& Unary::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = ";
switch (GetKind()) {
case Unary::Kind::kAddressOf:
out << "addr_of";
break;
case Unary::Kind::kComplement:
out << "complement";
break;
case Unary::Kind::kIndirection:
out << "indirection";
break;
case Unary::Kind::kNegation:
out << "negation";
break;
case Unary::Kind::kNot:
out << "not";
break;
}
out << " ";
val_->ToValue(out);
return out;
}
} // namespace tint::ir

6
src/tint/ir/unary.h
View File

@ -17,7 +17,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -52,11 +51,6 @@ class Unary : public utils::Castable<Unary, Instruction> {
/// @returns the value for the instruction
const Value* Val() const { return val_; }
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
private:
Kind kind_;
Value* val_ = nullptr;

21
src/tint/ir/unary_test.cc
View File

@ -14,7 +14,6 @@
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/test_helper.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
namespace {
@ -42,10 +41,6 @@ TEST_F(IR_InstructionTest, CreateAddressOf) {
auto lhs = inst->Val()->As<Constant>()->value;
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(ptr<private, i32, read_write>) = addr_of 4i");
}
TEST_F(IR_InstructionTest, CreateComplement) {
@ -60,10 +55,6 @@ TEST_F(IR_InstructionTest, CreateComplement) {
auto lhs = inst->Val()->As<Constant>()->value;
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = complement 4i");
}
TEST_F(IR_InstructionTest, CreateIndirection) {
@ -80,10 +71,6 @@ TEST_F(IR_InstructionTest, CreateIndirection) {
auto lhs = inst->Val()->As<Constant>()->value;
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = indirection 4i");
}
TEST_F(IR_InstructionTest, CreateNegation) {
@ -98,10 +85,6 @@ TEST_F(IR_InstructionTest, CreateNegation) {
auto lhs = inst->Val()->As<Constant>()->value;
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(i32) = negation 4i");
}
TEST_F(IR_InstructionTest, CreateNot) {
@ -116,10 +99,6 @@ TEST_F(IR_InstructionTest, CreateNot) {
auto lhs = inst->Val()->As<Constant>()->value;
ASSERT_TRUE(lhs->Is<constant::Scalar<bool>>());
EXPECT_TRUE(lhs->As<constant::Scalar<bool>>()->ValueAs<bool>());
utils::StringStream str;
inst->ToInstruction(str);
EXPECT_EQ(str.str(), "%1(bool) = not true");
}
TEST_F(IR_InstructionTest, Unary_Usage) {

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

@ -24,13 +24,4 @@ UserCall::UserCall(uint32_t id, const type::Type* type, Symbol name, utils::Vect
UserCall::~UserCall() = default;
utils::StringStream& UserCall::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = call " << name_.Name();
if (Args().Length() > 0) {
out << ", ";
}
EmitArgs(out);
return out;
}
} // namespace tint::ir

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

@ -18,7 +18,6 @@
#include "src/tint/ir/call.h"
#include "src/tint/symbol.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -41,11 +40,6 @@ class UserCall : public utils::Castable<UserCall, Call> {
/// @returns the function name
Symbol Name() const { return name_; }
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
private:
Symbol name_{};
};

6
src/tint/ir/value.h
View File

@ -17,7 +17,6 @@
#include "src/tint/type/type.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
#include "src/tint/utils/unique_vector.h"
// Forward declarations
@ -50,11 +49,6 @@ class Value : public utils::Castable<Value> {
/// @returns the type of the value
virtual const type::Type* Type() const = 0;
/// Write the value to the given stream
/// @param out the stream to write to
/// @returns the stream
virtual utils::StringStream& ToValue(utils::StringStream& out) const = 0;
protected:
/// Constructor
Value();

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

@ -27,9 +27,4 @@ Var::Var(uint32_t id,
Var::~Var() = default;
utils::StringStream& Var::ToInstruction(utils::StringStream& out) const {
ToValue(out) << " = var " << address_space_ << " " << access_;
return out;
}
} // namespace tint::ir

6
src/tint/ir/var.h
View File

@ -19,7 +19,6 @@
#include "src/tint/builtin/address_space.h"
#include "src/tint/ir/instruction.h"
#include "src/tint/utils/castable.h"
#include "src/tint/utils/string_stream.h"
namespace tint::ir {
@ -48,11 +47,6 @@ class Var : public utils::Castable<Var, Instruction> {
/// @returns the access mode
builtin::Access Access() const { return access_; }
/// Write the instruction to the given stream
/// @param out the stream to write to
/// @returns the stream
utils::StringStream& ToInstruction(utils::StringStream& out) const override;
private:
builtin::AddressSpace address_space_;
builtin::Access access_;