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>
diff --git a/src/tint/ir/binary.cc b/src/tint/ir/binary.cc
index a6a6aa7..063924f 100644
--- a/src/tint/ir/binary.cc
+++ b/src/tint/ir/binary.cc
@@ -19,8 +19,8 @@
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);
diff --git a/src/tint/ir/binary.h b/src/tint/ir/binary.h
index c686797..e941b3f 100644
--- a/src/tint/ir/binary.h
+++ b/src/tint/ir/binary.h
@@ -47,12 +47,11 @@
};
/// 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;
diff --git a/src/tint/ir/bitcast.cc b/src/tint/ir/bitcast.cc
index 89eedca..455bc69 100644
--- a/src/tint/ir/bitcast.cc
+++ b/src/tint/ir/bitcast.cc
@@ -19,8 +19,7 @@
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;
diff --git a/src/tint/ir/bitcast.h b/src/tint/ir/bitcast.h
index 0443701..0ad1acb 100644
--- a/src/tint/ir/bitcast.h
+++ b/src/tint/ir/bitcast.h
@@ -24,10 +24,9 @@
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;
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index e826be4d..fe350cd 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -109,7 +109,7 @@
}
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 @@
}
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 @@
}
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::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::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
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index ff01e2d..87a256b 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -364,11 +364,6 @@
/// The IR module.
Module ir;
-
- private:
- uint32_t next_inst_id() { return next_instruction_id_++; }
-
- uint32_t next_instruction_id_ = 1;
};
} // namespace tint::ir
diff --git a/src/tint/ir/builder_impl_binary_test.cc b/src/tint/ir/builder_impl_binary_test.cc
index a9a8978..ed13159 100644
--- a/src/tint/ir/builder_impl_binary_test.cc
+++ b/src/tint/ir/builder_impl_binary_test.cc
@@ -53,13 +53,13 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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
diff --git a/src/tint/ir/builder_impl_call_test.cc b/src/tint/ir/builder_impl_call_test.cc
index 4564880..6339626 100644
--- a/src/tint/ir/builder_impl_call_test.cc
+++ b/src/tint/ir/builder_impl_call_test.cc
@@ -91,14 +91,14 @@
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 @@
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 @@
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
diff --git a/src/tint/ir/builder_impl_materialize_test.cc b/src/tint/ir/builder_impl_materialize_test.cc
index 4f917f7..6b4ae84 100644
--- a/src/tint/ir/builder_impl_materialize_test.cc
+++ b/src/tint/ir/builder_impl_materialize_test.cc
@@ -35,8 +35,8 @@
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
diff --git a/src/tint/ir/builder_impl_store_test.cc b/src/tint/ir/builder_impl_store_test.cc
index f7cc6ab..7cd074c 100644
--- a/src/tint/ir/builder_impl_store_test.cc
+++ b/src/tint/ir/builder_impl_store_test.cc
@@ -36,13 +36,13 @@
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
diff --git a/src/tint/ir/builder_impl_test.cc b/src/tint/ir/builder_impl_test.cc
index 4a6e713..36eb3b6 100644
--- a/src/tint/ir/builder_impl_test.cc
+++ b/src/tint/ir/builder_impl_test.cc
@@ -42,8 +42,8 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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 @@
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
diff --git a/src/tint/ir/builder_impl_unary_test.cc b/src/tint/ir/builder_impl_unary_test.cc
index 9d35c02..091493d7 100644
--- a/src/tint/ir/builder_impl_unary_test.cc
+++ b/src/tint/ir/builder_impl_unary_test.cc
@@ -90,13 +90,13 @@
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 @@
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
diff --git a/src/tint/ir/builder_impl_var_test.cc b/src/tint/ir/builder_impl_var_test.cc
index 485684a..d68ce8a 100644
--- a/src/tint/ir/builder_impl_var_test.cc
+++ b/src/tint/ir/builder_impl_var_test.cc
@@ -33,7 +33,7 @@
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 @@
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 @@
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 @@
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
diff --git a/src/tint/ir/builtin.cc b/src/tint/ir/builtin.cc
index 5e53eec..fac0793 100644
--- a/src/tint/ir/builtin.cc
+++ b/src/tint/ir/builtin.cc
@@ -23,11 +23,8 @@
// \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;
diff --git a/src/tint/ir/builtin.h b/src/tint/ir/builtin.h
index 2106a68..9a82dba 100644
--- a/src/tint/ir/builtin.h
+++ b/src/tint/ir/builtin.h
@@ -25,14 +25,10 @@
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;
diff --git a/src/tint/ir/call.cc b/src/tint/ir/call.cc
index 9f5a8a5..d8c1af6 100644
--- a/src/tint/ir/call.cc
+++ b/src/tint/ir/call.cc
@@ -20,8 +20,8 @@
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);
}
diff --git a/src/tint/ir/call.h b/src/tint/ir/call.h
index b0f2e7c..fde193e 100644
--- a/src/tint/ir/call.h
+++ b/src/tint/ir/call.h
@@ -37,10 +37,9 @@
/// 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
diff --git a/src/tint/ir/construct.cc b/src/tint/ir/construct.cc
index 44c8862..18e830d 100644
--- a/src/tint/ir/construct.cc
+++ b/src/tint/ir/construct.cc
@@ -22,8 +22,8 @@
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;
diff --git a/src/tint/ir/construct.h b/src/tint/ir/construct.h
index b1711fb..b3ed6b6 100644
--- a/src/tint/ir/construct.h
+++ b/src/tint/ir/construct.h
@@ -24,10 +24,9 @@
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;
diff --git a/src/tint/ir/convert.cc b/src/tint/ir/convert.cc
index 969390c..edc2fe7 100644
--- a/src/tint/ir/convert.cc
+++ b/src/tint/ir/convert.cc
@@ -19,11 +19,10 @@
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;
diff --git a/src/tint/ir/convert.h b/src/tint/ir/convert.h
index 953a5fa..9fdbc5c 100644
--- a/src/tint/ir/convert.h
+++ b/src/tint/ir/convert.h
@@ -25,12 +25,10 @@
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;
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 58fb1a0..50d11cc 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -41,27 +41,29 @@
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 @@
}
}
-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 @@
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 @@
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 @@
// 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 @@
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 @@
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 @@
}
},
[&](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 @@
}
},
[&](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 @@
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();
}
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index d9182d3..2438d80 100644
--- a/src/tint/ir/disassembler.h
+++ b/src/tint/ir/disassembler.h
@@ -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 @@
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 @@
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;
};
diff --git a/src/tint/ir/discard.cc b/src/tint/ir/discard.cc
index 6c7ded7..feebb07 100644
--- a/src/tint/ir/discard.cc
+++ b/src/tint/ir/discard.cc
@@ -19,7 +19,7 @@
namespace tint::ir {
-Discard::Discard() : Base(kNoID, nullptr, utils::Empty) {}
+Discard::Discard() : Base(nullptr, utils::Empty) {}
Discard::~Discard() = default;
diff --git a/src/tint/ir/instruction.cc b/src/tint/ir/instruction.cc
index 0cf4966..bbd4992 100644
--- a/src/tint/ir/instruction.cc
+++ b/src/tint/ir/instruction.cc
@@ -20,7 +20,7 @@
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;
diff --git a/src/tint/ir/instruction.h b/src/tint/ir/instruction.h
index 1d4cafb..1bfde50 100644
--- a/src/tint/ir/instruction.h
+++ b/src/tint/ir/instruction.h
@@ -23,9 +23,6 @@
/// 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 @@
/// @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 @@
/// 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
diff --git a/src/tint/ir/unary.cc b/src/tint/ir/unary.cc
index e3f20eb..2fd7e57 100644
--- a/src/tint/ir/unary.cc
+++ b/src/tint/ir/unary.cc
@@ -19,8 +19,7 @@
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);
}
diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h
index 64301d4..8eb0bb4 100644
--- a/src/tint/ir/unary.h
+++ b/src/tint/ir/unary.h
@@ -32,11 +32,10 @@
};
/// 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;
diff --git a/src/tint/ir/user_call.cc b/src/tint/ir/user_call.cc
index e7e2e26..f718284 100644
--- a/src/tint/ir/user_call.cc
+++ b/src/tint/ir/user_call.cc
@@ -22,11 +22,8 @@
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;
diff --git a/src/tint/ir/user_call.h b/src/tint/ir/user_call.h
index 1edfa8b..5ada8f8 100644
--- a/src/tint/ir/user_call.h
+++ b/src/tint/ir/user_call.h
@@ -25,11 +25,10 @@
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;
diff --git a/src/tint/ir/var.cc b/src/tint/ir/var.cc
index b7ede1d..3ab97e0 100644
--- a/src/tint/ir/var.cc
+++ b/src/tint/ir/var.cc
@@ -19,11 +19,8 @@
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;
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index dd46f54..1387c60 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -26,14 +26,10 @@
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;
