diff --git a/src/tint/ir/binary.cc b/src/tint/ir/binary.cc
index a6a6aa796c..063924f1e1 100644
--- a/src/tint/ir/binary.cc
+++ b/src/tint/ir/binary.cc
@@ -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);
diff --git a/src/tint/ir/binary.h b/src/tint/ir/binary.h
index c686797994..e941b3f308 100644
--- a/src/tint/ir/binary.h
+++ b/src/tint/ir/binary.h
@@ -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;
diff --git a/src/tint/ir/bitcast.cc b/src/tint/ir/bitcast.cc
index 89eedcafb6..455bc69fb3 100644
--- a/src/tint/ir/bitcast.cc
+++ b/src/tint/ir/bitcast.cc
@@ -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;
 
diff --git a/src/tint/ir/bitcast.h b/src/tint/ir/bitcast.h
index 04437016c2..0ad1acbc86 100644
--- a/src/tint/ir/bitcast.h
+++ b/src/tint/ir/bitcast.h
@@ -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;
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index e826be4d93..fe350cd323 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -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
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index ff01e2deeb..87a256b97e 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -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
diff --git a/src/tint/ir/builder_impl_binary_test.cc b/src/tint/ir/builder_impl_binary_test.cc
index a9a8978939..ed131595d9 100644
--- a/src/tint/ir/builder_impl_binary_test.cc
+++ b/src/tint/ir/builder_impl_binary_test.cc
@@ -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
diff --git a/src/tint/ir/builder_impl_call_test.cc b/src/tint/ir/builder_impl_call_test.cc
index 456488054f..6339626109 100644
--- a/src/tint/ir/builder_impl_call_test.cc
+++ b/src/tint/ir/builder_impl_call_test.cc
@@ -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
diff --git a/src/tint/ir/builder_impl_materialize_test.cc b/src/tint/ir/builder_impl_materialize_test.cc
index 4f917f70b7..6b4ae843bf 100644
--- a/src/tint/ir/builder_impl_materialize_test.cc
+++ b/src/tint/ir/builder_impl_materialize_test.cc
@@ -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
 
diff --git a/src/tint/ir/builder_impl_store_test.cc b/src/tint/ir/builder_impl_store_test.cc
index f7cc6ab8bc..7cd074c96c 100644
--- a/src/tint/ir/builder_impl_store_test.cc
+++ b/src/tint/ir/builder_impl_store_test.cc
@@ -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
diff --git a/src/tint/ir/builder_impl_test.cc b/src/tint/ir/builder_impl_test.cc
index 4a6e713853..36eb3b6211 100644
--- a/src/tint/ir/builder_impl_test.cc
+++ b/src/tint/ir/builder_impl_test.cc
@@ -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
 
diff --git a/src/tint/ir/builder_impl_unary_test.cc b/src/tint/ir/builder_impl_unary_test.cc
index 9d35c020e3..091493d72f 100644
--- a/src/tint/ir/builder_impl_unary_test.cc
+++ b/src/tint/ir/builder_impl_unary_test.cc
@@ -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
diff --git a/src/tint/ir/builder_impl_var_test.cc b/src/tint/ir/builder_impl_var_test.cc
index 485684a4b3..d68ce8a60b 100644
--- a/src/tint/ir/builder_impl_var_test.cc
+++ b/src/tint/ir/builder_impl_var_test.cc
@@ -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
diff --git a/src/tint/ir/builtin.cc b/src/tint/ir/builtin.cc
index 5e53eecdc0..fac0793cc7 100644
--- a/src/tint/ir/builtin.cc
+++ b/src/tint/ir/builtin.cc
@@ -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;
 
diff --git a/src/tint/ir/builtin.h b/src/tint/ir/builtin.h
index 2106a68a05..9a82dbad95 100644
--- a/src/tint/ir/builtin.h
+++ b/src/tint/ir/builtin.h
@@ -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;
diff --git a/src/tint/ir/call.cc b/src/tint/ir/call.cc
index 9f5a8a5ea6..d8c1af696d 100644
--- a/src/tint/ir/call.cc
+++ b/src/tint/ir/call.cc
@@ -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);
     }
diff --git a/src/tint/ir/call.h b/src/tint/ir/call.h
index b0f2e7c06f..fde193e708 100644
--- a/src/tint/ir/call.h
+++ b/src/tint/ir/call.h
@@ -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
diff --git a/src/tint/ir/construct.cc b/src/tint/ir/construct.cc
index 44c8862034..18e830d80d 100644
--- a/src/tint/ir/construct.cc
+++ b/src/tint/ir/construct.cc
@@ -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;
 
diff --git a/src/tint/ir/construct.h b/src/tint/ir/construct.h
index b1711fb1bf..b3ed6b6bd1 100644
--- a/src/tint/ir/construct.h
+++ b/src/tint/ir/construct.h
@@ -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;
diff --git a/src/tint/ir/convert.cc b/src/tint/ir/convert.cc
index 969390c898..edc2fe7361 100644
--- a/src/tint/ir/convert.cc
+++ b/src/tint/ir/convert.cc
@@ -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;
 
diff --git a/src/tint/ir/convert.h b/src/tint/ir/convert.h
index 953a5fad80..9fdbc5ca9d 100644
--- a/src/tint/ir/convert.h
+++ b/src/tint/ir/convert.h
@@ -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;
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index 58fb1a060f..50d11cc522 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -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();
             }
diff --git a/src/tint/ir/disassembler.h b/src/tint/ir/disassembler.h
index d9182d3b43..2438d807ed 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 @@ 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;
 };
diff --git a/src/tint/ir/discard.cc b/src/tint/ir/discard.cc
index 6c7ded7468..feebb07742 100644
--- a/src/tint/ir/discard.cc
+++ b/src/tint/ir/discard.cc
@@ -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;
 
diff --git a/src/tint/ir/instruction.cc b/src/tint/ir/instruction.cc
index 0cf49662db..bbd4992cc0 100644
--- a/src/tint/ir/instruction.cc
+++ b/src/tint/ir/instruction.cc
@@ -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;
 
diff --git a/src/tint/ir/instruction.h b/src/tint/ir/instruction.h
index 1d4cafbb60..1bfde501f2 100644
--- a/src/tint/ir/instruction.h
+++ b/src/tint/ir/instruction.h
@@ -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
diff --git a/src/tint/ir/unary.cc b/src/tint/ir/unary.cc
index e3f20ebaba..2fd7e57784 100644
--- a/src/tint/ir/unary.cc
+++ b/src/tint/ir/unary.cc
@@ -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);
 }
diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h
index 64301d4d09..8eb0bb42d5 100644
--- a/src/tint/ir/unary.h
+++ b/src/tint/ir/unary.h
@@ -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;
diff --git a/src/tint/ir/user_call.cc b/src/tint/ir/user_call.cc
index e7e2e26855..f718284ee5 100644
--- a/src/tint/ir/user_call.cc
+++ b/src/tint/ir/user_call.cc
@@ -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;
 
diff --git a/src/tint/ir/user_call.h b/src/tint/ir/user_call.h
index 1edfa8b0e9..5ada8f8ddf 100644
--- a/src/tint/ir/user_call.h
+++ b/src/tint/ir/user_call.h
@@ -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;
diff --git a/src/tint/ir/var.cc b/src/tint/ir/var.cc
index b7ede1d795..3ab97e0045 100644
--- a/src/tint/ir/var.cc
+++ b/src/tint/ir/var.cc
@@ -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;
 
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index dd46f546f2..1387c609d4 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -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;