tint/ir: Remove value id field.
There's a few reasons for this change: * Not all values have identifiers, and carrying redundant fields is inefficent. * Not all IDs will be integers - much like LLVM IR and SPIR-V, we will likely want to disassemble with textual identifiers, so a uint32_t is not ideal, and a std::string is even more bloat for each value. * Transforms don't use identifiers, but instead raw pointers. We don't want to encourage using IDs as they're simply a less-efficient way to refer to values. * This makes values consistent with types and flow-control blocks, as they will both have their disassembly ID generated by the disassembler. The next step will be to add a hashmap to the module so that pre-declared value names can be stored out-of-band. Bug: tint:1718 Change-Id: I80beafc165f2bde54cc44a91015776926ca952b2 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131740 Kokoro: Kokoro <noreply+kokoro@google.com> Reviewed-by: James Price <jrprice@google.com> Commit-Queue: Ben Clayton <bclayton@google.com>
This commit is contained in:
parent
1c56265ee1
commit
a8236a5529
|
@ -19,8 +19,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Binary);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Binary::Binary(uint32_t identifier, Kind kind, const type::Type* ty, Value* lhs, Value* rhs)
|
||||
: Base(identifier, ty), kind_(kind), lhs_(lhs), rhs_(rhs) {
|
||||
Binary::Binary(Kind kind, const type::Type* ty, Value* lhs, Value* rhs)
|
||||
: Base(ty), kind_(kind), lhs_(lhs), rhs_(rhs) {
|
||||
TINT_ASSERT(IR, lhs_);
|
||||
TINT_ASSERT(IR, rhs_);
|
||||
lhs_->AddUsage(this);
|
||||
|
|
|
@ -47,12 +47,11 @@ class Binary : public utils::Castable<Binary, Instruction> {
|
|||
};
|
||||
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param kind the kind of binary instruction
|
||||
/// @param type the result type
|
||||
/// @param lhs the lhs of the instruction
|
||||
/// @param rhs the rhs of the instruction
|
||||
Binary(uint32_t id, Kind kind, const type::Type* type, Value* lhs, Value* rhs);
|
||||
Binary(Kind kind, const type::Type* type, Value* lhs, Value* rhs);
|
||||
Binary(const Binary& inst) = delete;
|
||||
Binary(Binary&& inst) = delete;
|
||||
~Binary() override;
|
||||
|
|
|
@ -19,8 +19,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Bitcast);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Bitcast::Bitcast(uint32_t identifier, const type::Type* ty, Value* val)
|
||||
: Base(identifier, ty, utils::Vector{val}) {}
|
||||
Bitcast::Bitcast(const type::Type* ty, Value* val) : Base(ty, utils::Vector{val}) {}
|
||||
|
||||
Bitcast::~Bitcast() = default;
|
||||
|
||||
|
|
|
@ -24,10 +24,9 @@ namespace tint::ir {
|
|||
class Bitcast : public utils::Castable<Bitcast, Call> {
|
||||
public:
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param type the result type
|
||||
/// @param val the value being bitcast
|
||||
Bitcast(uint32_t id, const type::Type* type, Value* val);
|
||||
Bitcast(const type::Type* type, Value* val);
|
||||
Bitcast(const Bitcast& inst) = delete;
|
||||
Bitcast(Bitcast&& inst) = delete;
|
||||
~Bitcast() override;
|
||||
|
|
|
@ -109,7 +109,7 @@ void Builder::Branch(Block* from, FlowNode* to, utils::VectorRef<Value*> args) {
|
|||
}
|
||||
|
||||
Binary* Builder::CreateBinary(Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs) {
|
||||
return ir.instructions.Create<ir::Binary>(next_inst_id(), kind, type, lhs, rhs);
|
||||
return ir.instructions.Create<ir::Binary>(kind, type, lhs, rhs);
|
||||
}
|
||||
|
||||
Binary* Builder::And(const type::Type* type, Value* lhs, Value* rhs) {
|
||||
|
@ -177,7 +177,7 @@ Binary* Builder::Modulo(const type::Type* type, Value* lhs, Value* rhs) {
|
|||
}
|
||||
|
||||
Unary* Builder::CreateUnary(Unary::Kind kind, const type::Type* type, Value* val) {
|
||||
return ir.instructions.Create<ir::Unary>(next_inst_id(), kind, type, val);
|
||||
return ir.instructions.Create<ir::Unary>(kind, type, val);
|
||||
}
|
||||
|
||||
Unary* Builder::AddressOf(const type::Type* type, Value* val) {
|
||||
|
@ -201,7 +201,7 @@ Binary* Builder::Not(const type::Type* type, Value* val) {
|
|||
}
|
||||
|
||||
ir::Bitcast* Builder::Bitcast(const type::Type* type, Value* val) {
|
||||
return ir.instructions.Create<ir::Bitcast>(next_inst_id(), type, val);
|
||||
return ir.instructions.Create<ir::Bitcast>(type, val);
|
||||
}
|
||||
|
||||
ir::Discard* Builder::Discard() {
|
||||
|
@ -211,23 +211,23 @@ ir::Discard* Builder::Discard() {
|
|||
ir::UserCall* Builder::UserCall(const type::Type* type,
|
||||
Symbol name,
|
||||
utils::VectorRef<Value*> args) {
|
||||
return ir.instructions.Create<ir::UserCall>(next_inst_id(), type, name, std::move(args));
|
||||
return ir.instructions.Create<ir::UserCall>(type, name, std::move(args));
|
||||
}
|
||||
|
||||
ir::Convert* Builder::Convert(const type::Type* to,
|
||||
const type::Type* from,
|
||||
utils::VectorRef<Value*> args) {
|
||||
return ir.instructions.Create<ir::Convert>(next_inst_id(), to, from, std::move(args));
|
||||
return ir.instructions.Create<ir::Convert>(to, from, std::move(args));
|
||||
}
|
||||
|
||||
ir::Construct* Builder::Construct(const type::Type* to, utils::VectorRef<Value*> args) {
|
||||
return ir.instructions.Create<ir::Construct>(next_inst_id(), to, std::move(args));
|
||||
return ir.instructions.Create<ir::Construct>(to, std::move(args));
|
||||
}
|
||||
|
||||
ir::Builtin* Builder::Builtin(const type::Type* type,
|
||||
builtin::Function func,
|
||||
utils::VectorRef<Value*> args) {
|
||||
return ir.instructions.Create<ir::Builtin>(next_inst_id(), type, func, args);
|
||||
return ir.instructions.Create<ir::Builtin>(type, func, args);
|
||||
}
|
||||
|
||||
ir::Store* Builder::Store(Value* to, Value* from) {
|
||||
|
@ -237,7 +237,7 @@ ir::Store* Builder::Store(Value* to, Value* from) {
|
|||
ir::Var* Builder::Declare(const type::Type* type,
|
||||
builtin::AddressSpace address_space,
|
||||
builtin::Access access) {
|
||||
return ir.instructions.Create<ir::Var>(next_inst_id(), type, address_space, access);
|
||||
return ir.instructions.Create<ir::Var>(type, address_space, access);
|
||||
}
|
||||
|
||||
} // namespace tint::ir
|
||||
|
|
|
@ -364,11 +364,6 @@ class Builder {
|
|||
|
||||
/// The IR module.
|
||||
Module ir;
|
||||
|
||||
private:
|
||||
uint32_t next_inst_id() { return next_instruction_id_++; }
|
||||
|
||||
uint32_t next_instruction_id_ = 1;
|
||||
};
|
||||
|
||||
} // namespace tint::ir
|
||||
|
|
|
@ -53,13 +53,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAdd) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = add %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -95,13 +95,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundSubtract) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = sub %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -137,13 +137,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundMultiply) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = mul %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -179,13 +179,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundDiv) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = div %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -221,13 +221,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundModulo) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = mod %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -263,13 +263,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundAnd) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, bool, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, bool, read_write> = and %1:ref<private, bool, read_write>, false
|
||||
store %1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
|
||||
ret
|
||||
|
@ -305,13 +305,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundOr) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, bool, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, bool, read_write> = or %1:ref<private, bool, read_write>, false
|
||||
store %1:ref<private, bool, read_write>, %2:ref<private, bool, read_write>
|
||||
ret
|
||||
|
@ -347,13 +347,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundXor) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = xor %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -371,26 +371,26 @@ 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
|
||||
%fn1 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = func my_func():bool
|
||||
%fn2 = block
|
||||
ret true
|
||||
func_end
|
||||
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn4 = block
|
||||
%1:bool = call my_func
|
||||
%2:bool = var function read_write
|
||||
store %2:bool, %1:bool
|
||||
branch %fn4
|
||||
branch %fn5
|
||||
|
||||
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
|
||||
%fn5 = if %1:bool [t: %fn6, f: %fn7, m: %fn8]
|
||||
# true branch
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
store %2:bool, false
|
||||
branch %fn7
|
||||
branch %fn8
|
||||
|
||||
# if merge
|
||||
%fn7 = block
|
||||
%fn8 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -406,27 +406,27 @@ 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
|
||||
%fn1 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = func my_func():bool
|
||||
%fn2 = block
|
||||
ret true
|
||||
func_end
|
||||
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn4 = block
|
||||
%1:bool = call my_func
|
||||
%2:bool = var function read_write
|
||||
store %2:bool, %1:bool
|
||||
branch %fn4
|
||||
branch %fn5
|
||||
|
||||
%fn4 = if %1:bool [t: %fn5, f: %fn6, m: %fn7]
|
||||
%fn5 = if %1:bool [t: %fn6, f: %fn7, m: %fn8]
|
||||
# true branch
|
||||
# false branch
|
||||
%fn6 = block
|
||||
%fn7 = block
|
||||
store %2:bool, true
|
||||
branch %fn7
|
||||
branch %fn8
|
||||
|
||||
# if merge
|
||||
%fn7 = block
|
||||
%fn8 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -568,13 +568,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftLeft) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = shiftl %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -610,13 +610,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Binary_CompoundShiftRight) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ref<private, u32, read_write> = shiftr %1:ref<private, u32, read_write>, 1u
|
||||
store %1:ref<private, u32, read_write>, %2:ref<private, u32, read_write>
|
||||
ret
|
||||
|
@ -636,32 +636,32 @@ 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
|
||||
%fn1 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = func my_func():f32
|
||||
%fn2 = block
|
||||
ret 0.0f
|
||||
func_end
|
||||
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn4 = block
|
||||
%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
|
||||
branch %fn5
|
||||
|
||||
%fn4 = if %2:bool [t: %fn5, f: %fn6, m: %fn7]
|
||||
%fn5 = if %2:bool [t: %fn6, f: %fn7, m: %fn8]
|
||||
# true branch
|
||||
%fn5 = block
|
||||
%fn6 = 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
|
||||
branch %fn7
|
||||
branch %fn8
|
||||
|
||||
# if merge
|
||||
%fn7 = block
|
||||
%fn8 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -678,13 +678,13 @@ 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
|
||||
%fn1 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = func my_func():bool
|
||||
%fn2 = block
|
||||
ret true
|
||||
func_end
|
||||
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn4 = block
|
||||
%1:bool = call my_func, false
|
||||
ret
|
||||
func_end
|
||||
|
|
|
@ -91,14 +91,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Convert) {
|
|||
auto m = r.Move();
|
||||
ASSERT_TRUE(r);
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%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)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:f32 = convert i32, %1:ref<private, i32, read_write>
|
||||
ret
|
||||
func_end
|
||||
|
@ -115,7 +115,7 @@ TEST_F(IR_BuilderImplTest, EmitExpression_ConstructEmpty) {
|
|||
auto m = r.Move();
|
||||
ASSERT_TRUE(r);
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, vec3<f32>, read_write> = var private read_write
|
||||
store %1:ref<private, vec3<f32>, read_write>, vec3<f32> 0.0f
|
||||
|
||||
|
@ -134,14 +134,14 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Construct) {
|
|||
auto m = r.Move();
|
||||
ASSERT_TRUE(r);
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%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)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:vec3<f32> = construct 2.0f, 3.0f, %1:ref<private, f32, read_write>
|
||||
ret
|
||||
func_end
|
||||
|
|
|
@ -35,8 +35,8 @@ TEST_F(IR_BuilderImplTest, EmitExpression_MaterializedCall) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = func test_function():f32
|
||||
%fn1 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = func test_function():f32
|
||||
%fn2 = block
|
||||
ret 2.0f
|
||||
func_end
|
||||
|
||||
|
|
|
@ -36,13 +36,13 @@ TEST_F(IR_BuilderImplTest, EmitStatement_Assign) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
store %1:ref<private, u32, read_write>, 4u
|
||||
ret
|
||||
func_end
|
||||
|
|
|
@ -42,8 +42,8 @@ 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
|
||||
%fn1 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = func f():void
|
||||
%fn2 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -89,21 +89,21 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = if true [t: %fn3, f: %fn4, m: %fn5]
|
||||
%fn3 = if true [t: %fn4, f: %fn5, m: %fn6]
|
||||
# true branch
|
||||
%fn3 = block
|
||||
branch %fn5
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
|
||||
# false branch
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
%fn5 = block
|
||||
branch %fn6
|
||||
|
||||
# if merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -138,20 +138,20 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = if true [t: %fn3, f: %fn4, m: %fn5]
|
||||
%fn3 = if true [t: %fn4, f: %fn5, m: %fn6]
|
||||
# true branch
|
||||
%fn3 = block
|
||||
%fn4 = block
|
||||
ret
|
||||
# false branch
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
%fn5 = block
|
||||
branch %fn6
|
||||
|
||||
# if merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -186,20 +186,20 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = if true [t: %fn3, f: %fn4, m: %fn5]
|
||||
%fn3 = if true [t: %fn4, f: %fn5, m: %fn6]
|
||||
# true branch
|
||||
%fn3 = block
|
||||
branch %fn5
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
|
||||
# false branch
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
# if merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -234,16 +234,16 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = if true [t: %fn3, f: %fn4]
|
||||
%fn3 = if true [t: %fn4, f: %fn5]
|
||||
# true branch
|
||||
%fn3 = block
|
||||
%fn4 = block
|
||||
ret
|
||||
# false branch
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -278,30 +278,30 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
|
|||
ASSERT_NE(loop_flow->merge.target, nullptr);
|
||||
|
||||
EXPECT_EQ(Disassemble(m),
|
||||
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = if true [t: %fn3, f: %fn4, m: %fn5]
|
||||
%fn3 = if true [t: %fn4, f: %fn5, m: %fn6]
|
||||
# true branch
|
||||
%fn3 = block
|
||||
branch %fn6
|
||||
%fn4 = block
|
||||
branch %fn7
|
||||
|
||||
%fn6 = loop [s: %fn7, m: %fn8]
|
||||
%fn7 = loop [s: %fn8, m: %fn9]
|
||||
# loop start
|
||||
%fn7 = block
|
||||
branch %fn8
|
||||
%fn8 = block
|
||||
branch %fn9
|
||||
|
||||
# loop merge
|
||||
%fn8 = block
|
||||
branch %fn5
|
||||
%fn9 = block
|
||||
branch %fn6
|
||||
|
||||
# false branch
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
%fn5 = block
|
||||
branch %fn6
|
||||
|
||||
# if merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -336,17 +336,17 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, m: %fn4]
|
||||
%fn3 = loop [s: %fn4, m: %fn5]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn4
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
|
||||
# loop merge
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -395,34 +395,34 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, c: %fn4, m: %fn5]
|
||||
%fn3 = loop [s: %fn4, c: %fn5, m: %fn6]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn6
|
||||
%fn4 = block
|
||||
branch %fn7
|
||||
|
||||
%fn6 = if true [t: %fn7, f: %fn8, m: %fn9]
|
||||
%fn7 = if true [t: %fn8, f: %fn9, m: %fn10]
|
||||
# true branch
|
||||
%fn7 = block
|
||||
branch %fn5
|
||||
%fn8 = block
|
||||
branch %fn6
|
||||
|
||||
# false branch
|
||||
%fn8 = block
|
||||
branch %fn9
|
||||
%fn9 = block
|
||||
branch %fn10
|
||||
|
||||
# if merge
|
||||
%fn9 = block
|
||||
branch %fn4
|
||||
%fn10 = block
|
||||
branch %fn5
|
||||
|
||||
# loop continuing
|
||||
%fn4 = block
|
||||
branch %fn3
|
||||
%fn5 = block
|
||||
branch %fn4
|
||||
|
||||
# loop merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -471,34 +471,34 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, c: %fn4, m: %fn5]
|
||||
%fn3 = loop [s: %fn4, c: %fn5, m: %fn6]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn4
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
|
||||
# loop continuing
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
%fn5 = block
|
||||
branch %fn7
|
||||
|
||||
%fn6 = if true [t: %fn7, f: %fn8, m: %fn9]
|
||||
%fn7 = if true [t: %fn8, f: %fn9, m: %fn10]
|
||||
# true branch
|
||||
%fn7 = block
|
||||
branch %fn5
|
||||
%fn8 = block
|
||||
branch %fn6
|
||||
|
||||
# false branch
|
||||
%fn8 = block
|
||||
branch %fn9
|
||||
%fn9 = block
|
||||
branch %fn10
|
||||
|
||||
# if merge
|
||||
%fn9 = block
|
||||
branch %fn3
|
||||
%fn10 = block
|
||||
branch %fn4
|
||||
|
||||
# loop merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -547,30 +547,30 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, c: %fn4]
|
||||
%fn3 = loop [s: %fn4, c: %fn5]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn5
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
|
||||
%fn5 = if true [t: %fn6, f: %fn7, m: %fn8]
|
||||
%fn6 = if true [t: %fn7, f: %fn8, m: %fn9]
|
||||
# true branch
|
||||
%fn6 = block
|
||||
%fn7 = block
|
||||
ret
|
||||
# false branch
|
||||
%fn7 = block
|
||||
branch %fn8
|
||||
%fn8 = block
|
||||
branch %fn9
|
||||
|
||||
# if merge
|
||||
%fn8 = block
|
||||
branch %fn4
|
||||
%fn9 = block
|
||||
branch %fn5
|
||||
|
||||
# loop continuing
|
||||
%fn4 = block
|
||||
branch %fn3
|
||||
%fn5 = block
|
||||
branch %fn4
|
||||
|
||||
func_end
|
||||
|
||||
|
@ -605,13 +605,13 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3]
|
||||
%fn3 = loop [s: %fn4]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
%fn4 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -668,13 +668,13 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3]
|
||||
%fn3 = loop [s: %fn4]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
%fn4 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -723,26 +723,26 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, m: %fn4]
|
||||
%fn3 = loop [s: %fn4, m: %fn5]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn5
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
|
||||
%fn5 = if true [t: %fn6, f: %fn7]
|
||||
%fn6 = if true [t: %fn7, f: %fn8]
|
||||
# true branch
|
||||
%fn6 = block
|
||||
branch %fn4
|
||||
%fn7 = block
|
||||
branch %fn5
|
||||
|
||||
# false branch
|
||||
%fn7 = block
|
||||
branch %fn4
|
||||
%fn8 = block
|
||||
branch %fn5
|
||||
|
||||
# loop merge
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -870,108 +870,108 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, c: %fn4, m: %fn5]
|
||||
%fn3 = loop [s: %fn4, c: %fn5, m: %fn6]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn6
|
||||
%fn4 = block
|
||||
branch %fn7
|
||||
|
||||
%fn6 = loop [s: %fn7, c: %fn8, m: %fn9]
|
||||
%fn7 = loop [s: %fn8, c: %fn9, m: %fn10]
|
||||
# loop start
|
||||
%fn7 = block
|
||||
branch %fn10
|
||||
%fn8 = block
|
||||
branch %fn11
|
||||
|
||||
%fn10 = if true [t: %fn11, f: %fn12, m: %fn13]
|
||||
%fn11 = if true [t: %fn12, f: %fn13, m: %fn14]
|
||||
# true branch
|
||||
%fn11 = block
|
||||
%fn12 = block
|
||||
branch %fn10
|
||||
|
||||
# false branch
|
||||
%fn13 = block
|
||||
branch %fn14
|
||||
|
||||
# if merge
|
||||
%fn14 = block
|
||||
branch %fn15
|
||||
|
||||
%fn15 = if true [t: %fn16, f: %fn17, m: %fn18]
|
||||
# true branch
|
||||
%fn16 = block
|
||||
branch %fn9
|
||||
|
||||
# false branch
|
||||
%fn12 = block
|
||||
branch %fn13
|
||||
%fn17 = block
|
||||
branch %fn18
|
||||
|
||||
# if merge
|
||||
%fn13 = block
|
||||
branch %fn14
|
||||
|
||||
%fn14 = if true [t: %fn15, f: %fn16, m: %fn17]
|
||||
# true branch
|
||||
%fn15 = block
|
||||
branch %fn8
|
||||
|
||||
# false branch
|
||||
%fn16 = block
|
||||
branch %fn17
|
||||
|
||||
# if merge
|
||||
%fn17 = block
|
||||
branch %fn8
|
||||
%fn18 = block
|
||||
branch %fn9
|
||||
|
||||
# loop continuing
|
||||
%fn8 = block
|
||||
branch %fn18
|
||||
%fn9 = block
|
||||
branch %fn19
|
||||
|
||||
%fn18 = loop [s: %fn19, m: %fn20]
|
||||
%fn19 = loop [s: %fn20, m: %fn21]
|
||||
# loop start
|
||||
%fn19 = block
|
||||
branch %fn20
|
||||
%fn20 = block
|
||||
branch %fn21
|
||||
|
||||
# loop merge
|
||||
%fn20 = block
|
||||
branch %fn21
|
||||
%fn21 = block
|
||||
branch %fn22
|
||||
|
||||
%fn21 = loop [s: %fn22, c: %fn23, m: %fn24]
|
||||
%fn22 = loop [s: %fn23, c: %fn24, m: %fn25]
|
||||
# loop start
|
||||
%fn22 = block
|
||||
branch %fn23
|
||||
%fn23 = block
|
||||
branch %fn24
|
||||
|
||||
# loop continuing
|
||||
%fn23 = block
|
||||
branch %fn25
|
||||
%fn24 = block
|
||||
branch %fn26
|
||||
|
||||
%fn25 = if true [t: %fn26, f: %fn27, m: %fn28]
|
||||
%fn26 = if true [t: %fn27, f: %fn28, m: %fn29]
|
||||
# true branch
|
||||
%fn26 = block
|
||||
branch %fn24
|
||||
%fn27 = block
|
||||
branch %fn25
|
||||
|
||||
# false branch
|
||||
%fn27 = block
|
||||
branch %fn28
|
||||
%fn28 = block
|
||||
branch %fn29
|
||||
|
||||
# if merge
|
||||
%fn28 = block
|
||||
branch %fn22
|
||||
%fn29 = block
|
||||
branch %fn23
|
||||
|
||||
# loop merge
|
||||
%fn24 = block
|
||||
branch %fn7
|
||||
%fn25 = block
|
||||
branch %fn8
|
||||
|
||||
# loop merge
|
||||
%fn9 = block
|
||||
branch %fn29
|
||||
%fn10 = block
|
||||
branch %fn30
|
||||
|
||||
%fn29 = if true [t: %fn30, f: %fn31, m: %fn32]
|
||||
%fn30 = if true [t: %fn31, f: %fn32, m: %fn33]
|
||||
# true branch
|
||||
%fn30 = block
|
||||
branch %fn5
|
||||
%fn31 = block
|
||||
branch %fn6
|
||||
|
||||
# false branch
|
||||
%fn31 = block
|
||||
branch %fn32
|
||||
%fn32 = block
|
||||
branch %fn33
|
||||
|
||||
# if merge
|
||||
%fn32 = block
|
||||
branch %fn4
|
||||
%fn33 = block
|
||||
branch %fn5
|
||||
|
||||
# loop continuing
|
||||
%fn4 = block
|
||||
branch %fn3
|
||||
%fn5 = block
|
||||
branch %fn4
|
||||
|
||||
# loop merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1015,34 +1015,34 @@ TEST_F(IR_BuilderImplTest, While) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, c: %fn4, m: %fn5]
|
||||
%fn3 = loop [s: %fn4, c: %fn5, m: %fn6]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn6
|
||||
%fn4 = block
|
||||
branch %fn7
|
||||
|
||||
%fn6 = if false [t: %fn7, f: %fn8, m: %fn9]
|
||||
%fn7 = if false [t: %fn8, f: %fn9, m: %fn10]
|
||||
# true branch
|
||||
%fn7 = block
|
||||
branch %fn9
|
||||
%fn8 = block
|
||||
branch %fn10
|
||||
|
||||
# false branch
|
||||
%fn8 = block
|
||||
branch %fn5
|
||||
%fn9 = block
|
||||
branch %fn6
|
||||
|
||||
# if merge
|
||||
%fn9 = block
|
||||
branch %fn4
|
||||
%fn10 = block
|
||||
branch %fn5
|
||||
|
||||
# loop continuing
|
||||
%fn4 = block
|
||||
branch %fn3
|
||||
%fn5 = block
|
||||
branch %fn4
|
||||
|
||||
# loop merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1086,29 +1086,29 @@ TEST_F(IR_BuilderImplTest, While_Return) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, m: %fn4]
|
||||
%fn3 = loop [s: %fn4, m: %fn5]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn5
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
|
||||
%fn5 = if true [t: %fn6, f: %fn7, m: %fn8]
|
||||
%fn6 = if true [t: %fn7, f: %fn8, m: %fn9]
|
||||
# true branch
|
||||
%fn6 = block
|
||||
branch %fn8
|
||||
%fn7 = block
|
||||
branch %fn9
|
||||
|
||||
# false branch
|
||||
%fn7 = block
|
||||
branch %fn4
|
||||
%fn8 = block
|
||||
branch %fn5
|
||||
|
||||
# if merge
|
||||
%fn8 = block
|
||||
%fn9 = block
|
||||
ret
|
||||
# loop merge
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1194,17 +1194,17 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = loop [s: %fn3, m: %fn4]
|
||||
%fn3 = loop [s: %fn4, m: %fn5]
|
||||
# loop start
|
||||
%fn3 = block
|
||||
branch %fn4
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
|
||||
# loop merge
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1254,25 +1254,25 @@ TEST_F(IR_BuilderImplTest, Switch) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = switch 1i [c: (0i, %fn3), c: (1i, %fn4), c: (default, %fn5), m: %fn6]
|
||||
%fn3 = switch 1i [c: (0i, %fn4), c: (1i, %fn5), c: (default, %fn6), m: %fn7]
|
||||
# case 0i
|
||||
%fn3 = block
|
||||
branch %fn6
|
||||
%fn4 = block
|
||||
branch %fn7
|
||||
|
||||
# case 1i
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
%fn5 = block
|
||||
branch %fn7
|
||||
|
||||
# case default
|
||||
%fn5 = block
|
||||
branch %fn6
|
||||
%fn6 = block
|
||||
branch %fn7
|
||||
|
||||
# switch merge
|
||||
%fn6 = block
|
||||
%fn7 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1319,17 +1319,17 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = switch 1i [c: (0i 1i default, %fn3), m: %fn4]
|
||||
%fn3 = switch 1i [c: (0i 1i default, %fn4), m: %fn5]
|
||||
# case 0i 1i default
|
||||
%fn3 = block
|
||||
branch %fn4
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
|
||||
# switch merge
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1364,17 +1364,17 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = switch 1i [c: (default, %fn3), m: %fn4]
|
||||
%fn3 = switch 1i [c: (default, %fn4), m: %fn5]
|
||||
# case default
|
||||
%fn3 = block
|
||||
branch %fn4
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
|
||||
# switch merge
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1418,21 +1418,21 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = switch 1i [c: (0i, %fn3), c: (default, %fn4), m: %fn5]
|
||||
%fn3 = switch 1i [c: (0i, %fn4), c: (default, %fn5), m: %fn6]
|
||||
# case 0i
|
||||
%fn3 = block
|
||||
branch %fn5
|
||||
%fn4 = block
|
||||
branch %fn6
|
||||
|
||||
# case default
|
||||
%fn4 = block
|
||||
branch %fn5
|
||||
%fn5 = block
|
||||
branch %fn6
|
||||
|
||||
# switch merge
|
||||
%fn5 = block
|
||||
%fn6 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
@ -1478,16 +1478,16 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
|
|||
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)]
|
||||
%fn1 = block
|
||||
branch %fn2
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
branch %fn3
|
||||
|
||||
%fn2 = switch 1i [c: (0i, %fn3), c: (default, %fn4)]
|
||||
%fn3 = switch 1i [c: (0i, %fn4), c: (default, %fn5)]
|
||||
# case 0i
|
||||
%fn3 = block
|
||||
%fn4 = block
|
||||
ret
|
||||
# case default
|
||||
%fn4 = block
|
||||
%fn5 = block
|
||||
ret
|
||||
func_end
|
||||
|
||||
|
|
|
@ -90,13 +90,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_AddressOf) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, i32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ptr<private, i32, read_write> = addr_of %1:ref<private, i32, read_write>
|
||||
ret
|
||||
func_end
|
||||
|
@ -116,13 +116,13 @@ TEST_F(IR_BuilderImplTest, EmitExpression_Unary_Indirection) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, i32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
||||
%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%fn2 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn3 = block
|
||||
%2:ptr<private, i32, read_write> = addr_of %1:ref<private, i32, read_write>
|
||||
%3:i32 = indirection %2:ptr<private, i32, read_write>
|
||||
ret
|
||||
|
|
|
@ -33,7 +33,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_NoInit) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
|
||||
|
||||
|
@ -49,7 +49,7 @@ TEST_F(IR_BuilderImplTest, Emit_GlobalVar_Init) {
|
|||
ASSERT_TRUE(r) << Error();
|
||||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn0 = block
|
||||
EXPECT_EQ(Disassemble(m), R"(%fn1 = block
|
||||
%1:ref<private, u32, read_write> = var private read_write
|
||||
store %1:ref<private, u32, read_write>, 2u
|
||||
|
||||
|
@ -67,8 +67,8 @@ TEST_F(IR_BuilderImplTest, Emit_Var_NoInit) {
|
|||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m),
|
||||
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn1 = block
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%1:ref<function, u32, read_write> = var function read_write
|
||||
ret
|
||||
func_end
|
||||
|
@ -86,8 +86,8 @@ TEST_F(IR_BuilderImplTest, Emit_Var_Init) {
|
|||
auto m = r.Move();
|
||||
|
||||
EXPECT_EQ(Disassemble(m),
|
||||
R"(%fn0 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn1 = block
|
||||
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)]
|
||||
%fn2 = block
|
||||
%1:ref<function, u32, read_write> = var function read_write
|
||||
store %1:ref<function, u32, read_write>, 2u
|
||||
ret
|
||||
|
|
|
@ -23,11 +23,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Builtin);
|
|||
// \cond DO_NOT_DOCUMENT
|
||||
namespace tint::ir {
|
||||
|
||||
Builtin::Builtin(uint32_t identifier,
|
||||
const type::Type* ty,
|
||||
builtin::Function func,
|
||||
utils::VectorRef<Value*> arguments)
|
||||
: Base(identifier, ty, std::move(arguments)), func_(func) {}
|
||||
Builtin::Builtin(const type::Type* ty, builtin::Function func, utils::VectorRef<Value*> arguments)
|
||||
: Base(ty, std::move(arguments)), func_(func) {}
|
||||
|
||||
Builtin::~Builtin() = default;
|
||||
|
||||
|
|
|
@ -25,14 +25,10 @@ namespace tint::ir {
|
|||
class Builtin : public utils::Castable<Builtin, Call> {
|
||||
public:
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param type the result type
|
||||
/// @param func the builtin function
|
||||
/// @param args the conversion arguments
|
||||
Builtin(uint32_t id,
|
||||
const type::Type* type,
|
||||
builtin::Function func,
|
||||
utils::VectorRef<Value*> args);
|
||||
Builtin(const type::Type* type, builtin::Function func, utils::VectorRef<Value*> args);
|
||||
Builtin(const Builtin& inst) = delete;
|
||||
Builtin(Builtin&& inst) = delete;
|
||||
~Builtin() override;
|
||||
|
|
|
@ -20,8 +20,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Call);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Call::Call(uint32_t identifier, const type::Type* ty, utils::VectorRef<Value*> arguments)
|
||||
: Base(identifier, ty), args(std::move(arguments)) {
|
||||
Call::Call(const type::Type* ty, utils::VectorRef<Value*> arguments)
|
||||
: Base(ty), args(std::move(arguments)) {
|
||||
for (auto* arg : args) {
|
||||
arg->AddUsage(this);
|
||||
}
|
||||
|
|
|
@ -37,10 +37,9 @@ class Call : public utils::Castable<Call, Instruction> {
|
|||
/// Constructor
|
||||
Call() = delete;
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param type the result type
|
||||
/// @param args the constructor arguments
|
||||
Call(uint32_t id, const type::Type* type, utils::VectorRef<Value*> args);
|
||||
Call(const type::Type* type, utils::VectorRef<Value*> args);
|
||||
};
|
||||
|
||||
} // namespace tint::ir
|
||||
|
|
|
@ -22,8 +22,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Construct);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Construct::Construct(uint32_t identifier, const type::Type* ty, utils::VectorRef<Value*> arguments)
|
||||
: Base(identifier, ty, std::move(arguments)) {}
|
||||
Construct::Construct(const type::Type* ty, utils::VectorRef<Value*> arguments)
|
||||
: Base(ty, std::move(arguments)) {}
|
||||
|
||||
Construct::~Construct() = default;
|
||||
|
||||
|
|
|
@ -24,10 +24,9 @@ namespace tint::ir {
|
|||
class Construct : public utils::Castable<Construct, Call> {
|
||||
public:
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param type the result type
|
||||
/// @param args the constructor arguments
|
||||
Construct(uint32_t id, const type::Type* type, utils::VectorRef<Value*> args);
|
||||
Construct(const type::Type* type, utils::VectorRef<Value*> args);
|
||||
Construct(const Construct& inst) = delete;
|
||||
Construct(Construct&& inst) = delete;
|
||||
~Construct() override;
|
||||
|
|
|
@ -19,11 +19,10 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Convert);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Convert::Convert(uint32_t identifier,
|
||||
const type::Type* to_type,
|
||||
Convert::Convert(const type::Type* to_type,
|
||||
const type::Type* from_type,
|
||||
utils::VectorRef<Value*> arguments)
|
||||
: Base(identifier, to_type, arguments), from_type_(from_type) {}
|
||||
: Base(to_type, arguments), from_type_(from_type) {}
|
||||
|
||||
Convert::~Convert() = default;
|
||||
|
||||
|
|
|
@ -25,12 +25,10 @@ namespace tint::ir {
|
|||
class Convert : public utils::Castable<Convert, Call> {
|
||||
public:
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param result_type the result type
|
||||
/// @param from_type the type being converted from
|
||||
/// @param args the conversion arguments
|
||||
Convert(uint32_t id,
|
||||
const type::Type* result_type,
|
||||
Convert(const type::Type* result_type,
|
||||
const type::Type* from_type,
|
||||
utils::VectorRef<Value*> args);
|
||||
Convert(const Convert& inst) = delete;
|
||||
|
|
|
@ -41,27 +41,29 @@ namespace tint::ir {
|
|||
namespace {
|
||||
|
||||
class ScopedStopNode {
|
||||
static constexpr size_t N = 32;
|
||||
|
||||
public:
|
||||
ScopedStopNode(std::unordered_set<const FlowNode*>* stop_nodes, const FlowNode* node)
|
||||
ScopedStopNode(utils::Hashset<const FlowNode*, N>& stop_nodes, const FlowNode* node)
|
||||
: stop_nodes_(stop_nodes), node_(node) {
|
||||
stop_nodes_->insert(node_);
|
||||
stop_nodes_.Add(node_);
|
||||
}
|
||||
|
||||
~ScopedStopNode() { stop_nodes_->erase(node_); }
|
||||
~ScopedStopNode() { stop_nodes_.Remove(node_); }
|
||||
|
||||
private:
|
||||
std::unordered_set<const FlowNode*>* stop_nodes_;
|
||||
utils::Hashset<const FlowNode*, N>& stop_nodes_;
|
||||
const FlowNode* node_;
|
||||
};
|
||||
|
||||
class ScopedIndent {
|
||||
public:
|
||||
explicit ScopedIndent(uint32_t* indent) : indent_(indent) { (*indent_) += 2; }
|
||||
explicit ScopedIndent(uint32_t& indent) : indent_(indent) { indent_ += 2; }
|
||||
|
||||
~ScopedIndent() { (*indent_) -= 2; }
|
||||
~ScopedIndent() { indent_ -= 2; }
|
||||
|
||||
private:
|
||||
uint32_t* indent_;
|
||||
uint32_t& indent_;
|
||||
};
|
||||
|
||||
} // namespace
|
||||
|
@ -85,30 +87,28 @@ void Disassembler::EmitBlockInstructions(const Block* b) {
|
|||
}
|
||||
}
|
||||
|
||||
size_t Disassembler::GetIdForNode(const FlowNode* node) {
|
||||
size_t Disassembler::IdOf(const FlowNode* node) {
|
||||
TINT_ASSERT(IR, node);
|
||||
return flow_node_ids_.GetOrCreate(node, [&] { return flow_node_ids_.Count(); });
|
||||
}
|
||||
|
||||
auto it = flow_node_to_id_.find(node);
|
||||
if (it != flow_node_to_id_.end()) {
|
||||
return it->second;
|
||||
}
|
||||
size_t id = next_node_id_++;
|
||||
flow_node_to_id_[node] = id;
|
||||
return id;
|
||||
std::string_view Disassembler::IdOf(const Value* value) {
|
||||
TINT_ASSERT(IR, value);
|
||||
return value_ids_.GetOrCreate(value, [&] { return std::to_string(value_ids_.Count()); });
|
||||
}
|
||||
|
||||
void Disassembler::Walk(const FlowNode* node) {
|
||||
if ((visited_.count(node) > 0) || (stop_nodes_.count(node) > 0)) {
|
||||
if (visited_.Contains(node) || stop_nodes_.Contains(node)) {
|
||||
return;
|
||||
}
|
||||
visited_.insert(node);
|
||||
visited_.Add(node);
|
||||
|
||||
tint::Switch(
|
||||
node,
|
||||
[&](const ir::Function* f) {
|
||||
TINT_SCOPED_ASSIGNMENT(in_function_, true);
|
||||
|
||||
Indent() << "%fn" << GetIdForNode(f) << " = func " << f->name.Name()
|
||||
Indent() << "%fn" << IdOf(f) << " = func " << f->name.Name()
|
||||
<< "():" << f->return_type->FriendlyName();
|
||||
|
||||
if (f->pipeline_stage != Function::PipelineStage::kUndefined) {
|
||||
|
@ -136,8 +136,8 @@ void Disassembler::Walk(const FlowNode* node) {
|
|||
out_ << std::endl;
|
||||
|
||||
{
|
||||
ScopedIndent func_indent(&indent_size_);
|
||||
ScopedStopNode scope(&stop_nodes_, f->end_target);
|
||||
ScopedIndent func_indent(indent_size_);
|
||||
ScopedStopNode scope(stop_nodes_, f->end_target);
|
||||
Walk(f->start_target);
|
||||
}
|
||||
Walk(f->end_target);
|
||||
|
@ -148,7 +148,7 @@ void Disassembler::Walk(const FlowNode* node) {
|
|||
return;
|
||||
}
|
||||
|
||||
Indent() << "%fn" << GetIdForNode(b) << " = block" << std::endl;
|
||||
Indent() << "%fn" << IdOf(b) << " = block" << std::endl;
|
||||
EmitBlockInstructions(b);
|
||||
|
||||
if (b->branch.target->Is<FunctionTerminator>()) {
|
||||
|
@ -157,7 +157,7 @@ void Disassembler::Walk(const FlowNode* node) {
|
|||
// Nothing to do
|
||||
} else {
|
||||
Indent() << "branch "
|
||||
<< "%fn" << GetIdForNode(b->branch.target);
|
||||
<< "%fn" << IdOf(b->branch.target);
|
||||
}
|
||||
if (!b->branch.args.IsEmpty()) {
|
||||
out_ << " ";
|
||||
|
@ -177,7 +177,7 @@ void Disassembler::Walk(const FlowNode* node) {
|
|||
Walk(b->branch.target);
|
||||
},
|
||||
[&](const ir::Switch* s) {
|
||||
Indent() << "%fn" << GetIdForNode(s) << " = switch ";
|
||||
Indent() << "%fn" << IdOf(s) << " = switch ";
|
||||
EmitValue(s->condition);
|
||||
out_ << " [";
|
||||
for (const auto& c : s->cases) {
|
||||
|
@ -196,16 +196,16 @@ void Disassembler::Walk(const FlowNode* node) {
|
|||
EmitValue(selector.val);
|
||||
}
|
||||
}
|
||||
out_ << ", %fn" << GetIdForNode(c.start.target) << ")";
|
||||
out_ << ", %fn" << IdOf(c.start.target) << ")";
|
||||
}
|
||||
if (s->merge.target->IsConnected()) {
|
||||
out_ << ", m: %fn" << GetIdForNode(s->merge.target);
|
||||
out_ << ", m: %fn" << IdOf(s->merge.target);
|
||||
}
|
||||
out_ << "]" << std::endl;
|
||||
|
||||
{
|
||||
ScopedIndent switch_indent(&indent_size_);
|
||||
ScopedStopNode scope(&stop_nodes_, s->merge.target);
|
||||
ScopedIndent switch_indent(indent_size_);
|
||||
ScopedStopNode scope(stop_nodes_, s->merge.target);
|
||||
for (const auto& c : s->cases) {
|
||||
Indent() << "# case ";
|
||||
for (const auto& selector : c.selectors) {
|
||||
|
@ -230,18 +230,17 @@ void Disassembler::Walk(const FlowNode* node) {
|
|||
}
|
||||
},
|
||||
[&](const ir::If* i) {
|
||||
Indent() << "%fn" << GetIdForNode(i) << " = if ";
|
||||
Indent() << "%fn" << IdOf(i) << " = if ";
|
||||
EmitValue(i->condition);
|
||||
out_ << " [t: %fn" << GetIdForNode(i->true_.target) << ", f: %fn"
|
||||
<< GetIdForNode(i->false_.target);
|
||||
out_ << " [t: %fn" << IdOf(i->true_.target) << ", f: %fn" << IdOf(i->false_.target);
|
||||
if (i->merge.target->IsConnected()) {
|
||||
out_ << ", m: %fn" << GetIdForNode(i->merge.target);
|
||||
out_ << ", m: %fn" << IdOf(i->merge.target);
|
||||
}
|
||||
out_ << "]" << std::endl;
|
||||
|
||||
{
|
||||
ScopedIndent if_indent(&indent_size_);
|
||||
ScopedStopNode scope(&stop_nodes_, i->merge.target);
|
||||
ScopedIndent if_indent(indent_size_);
|
||||
ScopedStopNode scope(stop_nodes_, i->merge.target);
|
||||
|
||||
Indent() << "# true branch" << std::endl;
|
||||
Walk(i->true_.target);
|
||||
|
@ -258,22 +257,21 @@ void Disassembler::Walk(const FlowNode* node) {
|
|||
}
|
||||
},
|
||||
[&](const ir::Loop* l) {
|
||||
Indent() << "%fn" << GetIdForNode(l) << " = loop [s: %fn"
|
||||
<< GetIdForNode(l->start.target);
|
||||
Indent() << "%fn" << IdOf(l) << " = loop [s: %fn" << IdOf(l->start.target);
|
||||
|
||||
if (l->continuing.target->IsConnected()) {
|
||||
out_ << ", c: %fn" << GetIdForNode(l->continuing.target);
|
||||
out_ << ", c: %fn" << IdOf(l->continuing.target);
|
||||
}
|
||||
if (l->merge.target->IsConnected()) {
|
||||
out_ << ", m: %fn" << GetIdForNode(l->merge.target);
|
||||
out_ << ", m: %fn" << IdOf(l->merge.target);
|
||||
}
|
||||
out_ << "]" << std::endl;
|
||||
|
||||
{
|
||||
ScopedStopNode loop_scope(&stop_nodes_, l->merge.target);
|
||||
ScopedIndent loop_indent(&indent_size_);
|
||||
ScopedStopNode loop_scope(stop_nodes_, l->merge.target);
|
||||
ScopedIndent loop_indent(indent_size_);
|
||||
{
|
||||
ScopedStopNode inner_scope(&stop_nodes_, l->continuing.target);
|
||||
ScopedStopNode inner_scope(stop_nodes_, l->continuing.target);
|
||||
Indent() << "# loop start" << std::endl;
|
||||
Walk(l->start.target);
|
||||
}
|
||||
|
@ -355,11 +353,7 @@ void Disassembler::EmitValue(const Value* val) {
|
|||
emit(constant->value);
|
||||
},
|
||||
[&](const ir::Instruction* i) {
|
||||
if (i->id == ir::Instruction::kNoID) {
|
||||
out_ << "<no-id>";
|
||||
} else {
|
||||
out_ << "%" << i->id;
|
||||
}
|
||||
out_ << "%" << IdOf(i);
|
||||
if (i->Type() != nullptr) {
|
||||
out_ << ":" << i->Type()->FriendlyName();
|
||||
}
|
||||
|
|
|
@ -16,14 +16,14 @@
|
|||
#define SRC_TINT_IR_DISASSEMBLER_H_
|
||||
|
||||
#include <string>
|
||||
#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/hashmap.h"
|
||||
#include "src/tint/utils/hashset.h"
|
||||
#include "src/tint/utils/string_stream.h"
|
||||
|
||||
namespace tint::ir {
|
||||
|
@ -49,7 +49,9 @@ class Disassembler {
|
|||
|
||||
private:
|
||||
utils::StringStream& Indent();
|
||||
size_t GetIdForNode(const FlowNode* node);
|
||||
|
||||
size_t IdOf(const FlowNode* node);
|
||||
std::string_view IdOf(const Value* node);
|
||||
|
||||
void Walk(const FlowNode* node);
|
||||
void EmitInstruction(const Instruction* inst);
|
||||
|
@ -60,10 +62,10 @@ class Disassembler {
|
|||
|
||||
const Module& mod_;
|
||||
utils::StringStream out_;
|
||||
std::unordered_set<const FlowNode*> visited_;
|
||||
std::unordered_set<const FlowNode*> stop_nodes_;
|
||||
std::unordered_map<const FlowNode*, size_t> flow_node_to_id_;
|
||||
size_t next_node_id_ = 0;
|
||||
utils::Hashset<const FlowNode*, 32> visited_;
|
||||
utils::Hashset<const FlowNode*, 32> stop_nodes_;
|
||||
utils::Hashmap<const FlowNode*, size_t, 32> flow_node_ids_;
|
||||
utils::Hashmap<const Value*, std::string, 32> value_ids_;
|
||||
uint32_t indent_size_ = 0;
|
||||
bool in_function_ = false;
|
||||
};
|
||||
|
|
|
@ -19,7 +19,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Discard);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Discard::Discard() : Base(kNoID, nullptr, utils::Empty) {}
|
||||
Discard::Discard() : Base(nullptr, utils::Empty) {}
|
||||
|
||||
Discard::~Discard() = default;
|
||||
|
||||
|
|
|
@ -20,7 +20,7 @@ namespace tint::ir {
|
|||
|
||||
Instruction::Instruction() = default;
|
||||
|
||||
Instruction::Instruction(uint32_t identifier, const type::Type* ty) : id(identifier), type(ty) {}
|
||||
Instruction::Instruction(const type::Type* ty) : type(ty) {}
|
||||
|
||||
Instruction::~Instruction() = default;
|
||||
|
||||
|
|
|
@ -23,9 +23,6 @@ namespace tint::ir {
|
|||
/// An instruction in the IR.
|
||||
class Instruction : public utils::Castable<Instruction, Value> {
|
||||
public:
|
||||
/// The identifier used by instructions that have no value.
|
||||
static constexpr uint32_t kNoID = 0;
|
||||
|
||||
Instruction(const Instruction& inst) = delete;
|
||||
Instruction(Instruction&& inst) = delete;
|
||||
/// Destructor
|
||||
|
@ -37,9 +34,6 @@ class Instruction : public utils::Castable<Instruction, Value> {
|
|||
/// @returns the type of the value
|
||||
const type::Type* Type() const override { return type; }
|
||||
|
||||
/// The instruction identifier
|
||||
const uint32_t id = kNoID;
|
||||
|
||||
/// The instruction type
|
||||
const type::Type* type = nullptr;
|
||||
|
||||
|
@ -47,9 +41,8 @@ class Instruction : public utils::Castable<Instruction, Value> {
|
|||
/// Constructor
|
||||
Instruction();
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param type the result type
|
||||
Instruction(uint32_t id, const type::Type* type);
|
||||
explicit Instruction(const type::Type* type);
|
||||
};
|
||||
|
||||
} // namespace tint::ir
|
||||
|
|
|
@ -19,8 +19,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Unary);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Unary::Unary(uint32_t identifier, Kind kind, const type::Type* ty, Value* val)
|
||||
: Base(identifier, ty), kind_(kind), val_(val) {
|
||||
Unary::Unary(Kind kind, const type::Type* ty, Value* val) : Base(ty), kind_(kind), val_(val) {
|
||||
TINT_ASSERT(IR, val_);
|
||||
val_->AddUsage(this);
|
||||
}
|
||||
|
|
|
@ -32,11 +32,10 @@ class Unary : public utils::Castable<Unary, Instruction> {
|
|||
};
|
||||
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param kind the kind of unary instruction
|
||||
/// @param type the result type
|
||||
/// @param val the lhs of the instruction
|
||||
Unary(uint32_t id, Kind kind, const type::Type* type, Value* val);
|
||||
Unary(Kind kind, const type::Type* type, Value* val);
|
||||
Unary(const Unary& inst) = delete;
|
||||
Unary(Unary&& inst) = delete;
|
||||
~Unary() override;
|
||||
|
|
|
@ -22,11 +22,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::UserCall);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
UserCall::UserCall(uint32_t identifier,
|
||||
const type::Type* ty,
|
||||
Symbol n,
|
||||
utils::VectorRef<Value*> arguments)
|
||||
: Base(identifier, ty, std::move(arguments)), name(n) {}
|
||||
UserCall::UserCall(const type::Type* ty, Symbol n, utils::VectorRef<Value*> arguments)
|
||||
: Base(ty, std::move(arguments)), name(n) {}
|
||||
|
||||
UserCall::~UserCall() = default;
|
||||
|
||||
|
|
|
@ -25,11 +25,10 @@ namespace tint::ir {
|
|||
class UserCall : public utils::Castable<UserCall, Call> {
|
||||
public:
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param type the result type
|
||||
/// @param name the function name
|
||||
/// @param args the function arguments
|
||||
UserCall(uint32_t id, const type::Type* type, Symbol name, utils::VectorRef<Value*> args);
|
||||
UserCall(const type::Type* type, Symbol name, utils::VectorRef<Value*> args);
|
||||
UserCall(const UserCall& inst) = delete;
|
||||
UserCall(UserCall&& inst) = delete;
|
||||
~UserCall() override;
|
||||
|
|
|
@ -19,11 +19,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Var);
|
|||
|
||||
namespace tint::ir {
|
||||
|
||||
Var::Var(uint32_t identifier,
|
||||
const type::Type* ty,
|
||||
builtin::AddressSpace addr_space,
|
||||
builtin::Access acc)
|
||||
: Base(identifier, ty), address_space(addr_space), access(acc) {}
|
||||
Var::Var(const type::Type* ty, builtin::AddressSpace addr_space, builtin::Access acc)
|
||||
: Base(ty), address_space(addr_space), access(acc) {}
|
||||
|
||||
Var::~Var() = default;
|
||||
|
||||
|
|
|
@ -26,14 +26,10 @@ namespace tint::ir {
|
|||
class Var : public utils::Castable<Var, Instruction> {
|
||||
public:
|
||||
/// Constructor
|
||||
/// @param id the instruction id
|
||||
/// @param type the type
|
||||
/// @param address_space the address space of the var
|
||||
/// @param access the access mode of the var
|
||||
Var(uint32_t id,
|
||||
const type::Type* type,
|
||||
builtin::AddressSpace address_space,
|
||||
builtin::Access access);
|
||||
Var(const type::Type* type, builtin::AddressSpace address_space, builtin::Access access);
|
||||
Var(const Var& inst) = delete;
|
||||
Var(Var&& inst) = delete;
|
||||
~Var() override;
|
||||
|
|
Loading…
Reference in New Issue