diff --git a/src/tint/lang/core/ir/access.cc b/src/tint/lang/core/ir/access.cc
index b4434e6..b9f750e 100644
--- a/src/tint/lang/core/ir/access.cc
+++ b/src/tint/lang/core/ir/access.cc
@@ -51,7 +51,7 @@
     auto new_result = ctx.Clone(Result(0));
     auto obj = ctx.Remap(Object());
     auto indices = ctx.Remap<Access::kDefaultNumOperands>(Indices());
-    return ctx.ir.instructions.Create<Access>(new_result, obj, indices);
+    return ctx.ir.allocators.instructions.Create<Access>(new_result, obj, indices);
 }
 //! @endcond
 
diff --git a/src/tint/lang/core/ir/binary/decode.cc b/src/tint/lang/core/ir/binary/decode.cc
index a65c48f..825be15 100644
--- a/src/tint/lang/core/ir/binary/decode.cc
+++ b/src/tint/lang/core/ir/binary/decode.cc
@@ -153,7 +153,7 @@
     // Functions
     ////////////////////////////////////////////////////////////////////////////
     ir::Function* CreateFunction(const pb::Function&) {
-        return mod_out_.values.Create<ir::Function>();
+        return mod_out_.allocators.values.Create<ir::Function>();
     }
 
     void PopulateFunction(ir::Function* fn_out, const pb::Function& fn_in) {
@@ -344,69 +344,69 @@
     }
 
     ir::Access* CreateInstructionAccess(const pb::InstructionAccess&) {
-        return mod_out_.instructions.Create<ir::Access>();
+        return mod_out_.allocators.instructions.Create<ir::Access>();
     }
 
     ir::CoreBinary* CreateInstructionBinary(const pb::InstructionBinary& binary_in) {
-        auto* binary_out = mod_out_.instructions.Create<ir::CoreBinary>();
+        auto* binary_out = mod_out_.allocators.instructions.Create<ir::CoreBinary>();
         binary_out->SetOp(BinaryOp(binary_in.op()));
         return binary_out;
     }
 
     ir::Bitcast* CreateInstructionBitcast(const pb::InstructionBitcast&) {
-        return mod_out_.instructions.Create<ir::Bitcast>();
+        return mod_out_.allocators.instructions.Create<ir::Bitcast>();
     }
 
     ir::BreakIf* CreateInstructionBreakIf(const pb::InstructionBreakIf&) {
-        auto* break_if_out = mod_out_.instructions.Create<ir::BreakIf>();
+        auto* break_if_out = mod_out_.allocators.instructions.Create<ir::BreakIf>();
         break_ifs_.Push(break_if_out);
         return break_if_out;
     }
 
     ir::CoreBuiltinCall* CreateInstructionBuiltinCall(const pb::InstructionBuiltinCall& call_in) {
-        auto* call_out = mod_out_.instructions.Create<ir::CoreBuiltinCall>();
+        auto* call_out = mod_out_.allocators.instructions.Create<ir::CoreBuiltinCall>();
         call_out->SetFunc(BuiltinFn(call_in.builtin()));
         return call_out;
     }
 
     ir::Construct* CreateInstructionConstruct(const pb::InstructionConstruct&) {
-        return mod_out_.instructions.Create<ir::Construct>();
+        return mod_out_.allocators.instructions.Create<ir::Construct>();
     }
 
     ir::Continue* CreateInstructionContinue(const pb::InstructionContinue&) {
-        auto* continue_ = mod_out_.instructions.Create<ir::Continue>();
+        auto* continue_ = mod_out_.allocators.instructions.Create<ir::Continue>();
         continues_.Push(continue_);
         return continue_;
     }
 
     ir::Convert* CreateInstructionConvert(const pb::InstructionConvert&) {
-        return mod_out_.instructions.Create<ir::Convert>();
+        return mod_out_.allocators.instructions.Create<ir::Convert>();
     }
 
     ir::ExitIf* CreateInstructionExitIf(const pb::InstructionExitIf&) {
-        auto* exit_out = mod_out_.instructions.Create<ir::ExitIf>();
+        auto* exit_out = mod_out_.allocators.instructions.Create<ir::ExitIf>();
         exit_ifs_.Push(exit_out);
         return exit_out;
     }
 
     ir::ExitLoop* CreateInstructionExitLoop(const pb::InstructionExitLoop&) {
-        auto* exit_out = mod_out_.instructions.Create<ir::ExitLoop>();
+        auto* exit_out = mod_out_.allocators.instructions.Create<ir::ExitLoop>();
         exit_loops_.Push(exit_out);
         return exit_out;
     }
 
     ir::ExitSwitch* CreateInstructionExitSwitch(const pb::InstructionExitSwitch&) {
-        auto* exit_out = mod_out_.instructions.Create<ir::ExitSwitch>();
+        auto* exit_out = mod_out_.allocators.instructions.Create<ir::ExitSwitch>();
         exit_switches_.Push(exit_out);
         return exit_out;
     }
 
     ir::Discard* CreateInstructionDiscard(const pb::InstructionDiscard&) {
-        return mod_out_.instructions.Create<ir::Discard>();
+        return mod_out_.allocators.instructions.Create<ir::Discard>();
     }
 
     ir::If* CreateInstructionIf(const pb::InstructionIf& if_in) {
-        auto* if_out = mod_out_.instructions.Create<ir::If>();
+        auto* if_out = mod_out_.allocators.instructions.Create<ir::If>();
         if (if_in.has_true_()) {
             if_out->SetTrue(Block(if_in.true_()));
         }
@@ -417,20 +417,20 @@
     }
 
     ir::Let* CreateInstructionLet(const pb::InstructionLet&) {
-        return mod_out_.instructions.Create<ir::Let>();
+        return mod_out_.allocators.instructions.Create<ir::Let>();
     }
 
     ir::Load* CreateInstructionLoad(const pb::InstructionLoad&) {
-        return mod_out_.instructions.Create<ir::Load>();
+        return mod_out_.allocators.instructions.Create<ir::Load>();
     }
 
     ir::LoadVectorElement* CreateInstructionLoadVectorElement(
         const pb::InstructionLoadVectorElement&) {
-        return mod_out_.instructions.Create<ir::LoadVectorElement>();
+        return mod_out_.allocators.instructions.Create<ir::LoadVectorElement>();
     }
 
     ir::Loop* CreateInstructionLoop(const pb::InstructionLoop& loop_in) {
-        auto* loop_out = mod_out_.instructions.Create<ir::Loop>();
+        auto* loop_out = mod_out_.allocators.instructions.Create<ir::Loop>();
         if (loop_in.has_initalizer()) {
             loop_out->SetInitializer(Block(loop_in.initalizer()));
         } else {
@@ -446,26 +446,26 @@
     }
 
     ir::NextIteration* CreateInstructionNextIteration(const pb::InstructionNextIteration&) {
-        auto* next_it_out = mod_out_.instructions.Create<ir::NextIteration>();
+        auto* next_it_out = mod_out_.allocators.instructions.Create<ir::NextIteration>();
         next_iterations_.Push(next_it_out);
         return next_it_out;
     }
 
     ir::Return* CreateInstructionReturn(const pb::InstructionReturn&) {
-        return mod_out_.instructions.Create<ir::Return>();
+        return mod_out_.allocators.instructions.Create<ir::Return>();
     }
 
     ir::Store* CreateInstructionStore(const pb::InstructionStore&) {
-        return mod_out_.instructions.Create<ir::Store>();
+        return mod_out_.allocators.instructions.Create<ir::Store>();
     }
 
     ir::StoreVectorElement* CreateInstructionStoreVectorElement(
         const pb::InstructionStoreVectorElement&) {
-        return mod_out_.instructions.Create<ir::StoreVectorElement>();
+        return mod_out_.allocators.instructions.Create<ir::StoreVectorElement>();
     }
 
     ir::Swizzle* CreateInstructionSwizzle(const pb::InstructionSwizzle& swizzle_in) {
-        auto* swizzle_out = mod_out_.instructions.Create<ir::Swizzle>();
+        auto* swizzle_out = mod_out_.allocators.instructions.Create<ir::Swizzle>();
         Vector<uint32_t, 4> indices;
         for (auto idx : swizzle_in.indices()) {
             indices.Push(idx);
@@ -475,7 +475,7 @@
     }
 
     ir::Switch* CreateInstructionSwitch(const pb::InstructionSwitch& switch_in) {
-        auto* switch_out = mod_out_.instructions.Create<ir::Switch>();
+        auto* switch_out = mod_out_.allocators.instructions.Create<ir::Switch>();
         for (auto& case_in : switch_in.cases()) {
             ir::Switch::Case case_out{};
             case_out.block = Block(case_in.block());
@@ -495,17 +495,17 @@
     }
 
     ir::CoreUnary* CreateInstructionUnary(const pb::InstructionUnary& unary_in) {
-        auto* unary_out = mod_out_.instructions.Create<ir::CoreUnary>();
+        auto* unary_out = mod_out_.allocators.instructions.Create<ir::CoreUnary>();
         unary_out->SetOp(UnaryOp(unary_in.op()));
         return unary_out;
     }
 
     ir::UserCall* CreateInstructionUserCall(const pb::InstructionUserCall&) {
-        return mod_out_.instructions.Create<ir::UserCall>();
+        return mod_out_.allocators.instructions.Create<ir::UserCall>();
     }
 
     ir::Var* CreateInstructionVar(const pb::InstructionVar& var_in) {
-        auto* var_out = mod_out_.instructions.Create<ir::Var>();
+        auto* var_out = mod_out_.allocators.instructions.Create<ir::Var>();
         if (var_in.has_binding_point()) {
             auto& bp_in = var_in.binding_point();
             var_out->SetBindingPoint(bp_in.group(), bp_in.binding());
diff --git a/src/tint/lang/core/ir/bitcast.cc b/src/tint/lang/core/ir/bitcast.cc
index 8c6ece5..01c06b6 100644
--- a/src/tint/lang/core/ir/bitcast.cc
+++ b/src/tint/lang/core/ir/bitcast.cc
@@ -46,7 +46,7 @@
 Bitcast* Bitcast::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto* val = ctx.Remap(Val());
-    return ctx.ir.instructions.Create<Bitcast>(new_result, val);
+    return ctx.ir.allocators.instructions.Create<Bitcast>(new_result, val);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/block_param.cc b/src/tint/lang/core/ir/block_param.cc
index 8a663f8..0866efb 100644
--- a/src/tint/lang/core/ir/block_param.cc
+++ b/src/tint/lang/core/ir/block_param.cc
@@ -42,7 +42,7 @@
 BlockParam::~BlockParam() = default;
 
 BlockParam* BlockParam::Clone(CloneContext& ctx) {
-    auto* new_bp = ctx.ir.values.Create<BlockParam>(type_);
+    auto* new_bp = ctx.ir.allocators.values.Create<BlockParam>(type_);
 
     auto name = ctx.ir.NameOf(this);
     if (name.IsValid()) {
diff --git a/src/tint/lang/core/ir/break_if.cc b/src/tint/lang/core/ir/break_if.cc
index dc312e6..53a7145 100644
--- a/src/tint/lang/core/ir/break_if.cc
+++ b/src/tint/lang/core/ir/break_if.cc
@@ -59,7 +59,7 @@
     auto* loop = ctx.Remap(loop_);
     auto* cond = ctx.Remap(Condition());
     auto args = ctx.Remap<BreakIf::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<BreakIf>(cond, loop, args);
+    return ctx.ir.allocators.instructions.Create<BreakIf>(cond, loop, args);
 }
 
 void BreakIf::SetLoop(ir::Loop* loop) {
diff --git a/src/tint/lang/core/ir/builder.cc b/src/tint/lang/core/ir/builder.cc
index 48da157..16d0ed0 100644
--- a/src/tint/lang/core/ir/builder.cc
+++ b/src/tint/lang/core/ir/builder.cc
@@ -54,7 +54,7 @@
 Function* Builder::Function(const core::type::Type* return_type,
                             Function::PipelineStage stage,
                             std::optional<std::array<uint32_t, 3>> wg_size) {
-    auto* ir_func = ir.values.Create<ir::Function>(return_type, stage, wg_size);
+    auto* ir_func = ir.allocators.values.Create<ir::Function>(return_type, stage, wg_size);
     ir_func->SetBlock(Block());
     ir.functions.Push(ir_func);
     return ir_func;
@@ -70,7 +70,8 @@
 }
 
 ir::Loop* Builder::Loop() {
-    return Append(ir.instructions.Create<ir::Loop>(Block(), MultiInBlock(), MultiInBlock()));
+    return Append(
+        ir.allocators.instructions.Create<ir::Loop>(Block(), MultiInBlock(), MultiInBlock()));
 }
 
 Block* Builder::Case(ir::Switch* s, VectorRef<ir::Constant*> values) {
@@ -95,11 +96,11 @@
 }
 
 ir::Discard* Builder::Discard() {
-    return Append(ir.instructions.Create<ir::Discard>());
+    return Append(ir.allocators.instructions.Create<ir::Discard>());
 }
 
 ir::Var* Builder::Var(const core::type::MemoryView* type) {
-    return Append(ir.instructions.Create<ir::Var>(InstructionResult(type)));
+    return Append(ir.allocators.instructions.Create<ir::Var>(InstructionResult(type)));
 }
 
 ir::Var* Builder::Var(std::string_view name, const core::type::MemoryView* type) {
@@ -109,31 +110,31 @@
 }
 
 ir::BlockParam* Builder::BlockParam(const core::type::Type* type) {
-    return ir.values.Create<ir::BlockParam>(type);
+    return ir.allocators.values.Create<ir::BlockParam>(type);
 }
 
 ir::BlockParam* Builder::BlockParam(std::string_view name, const core::type::Type* type) {
-    auto* param = ir.values.Create<ir::BlockParam>(type);
+    auto* param = ir.allocators.values.Create<ir::BlockParam>(type);
     ir.SetName(param, name);
     return param;
 }
 
 ir::FunctionParam* Builder::FunctionParam(const core::type::Type* type) {
-    return ir.values.Create<ir::FunctionParam>(type);
+    return ir.allocators.values.Create<ir::FunctionParam>(type);
 }
 
 ir::FunctionParam* Builder::FunctionParam(std::string_view name, const core::type::Type* type) {
-    auto* param = ir.values.Create<ir::FunctionParam>(type);
+    auto* param = ir.allocators.values.Create<ir::FunctionParam>(type);
     ir.SetName(param, name);
     return param;
 }
 
 ir::TerminateInvocation* Builder::TerminateInvocation() {
-    return Append(ir.instructions.Create<ir::TerminateInvocation>());
+    return Append(ir.allocators.instructions.Create<ir::TerminateInvocation>());
 }
 
 ir::Unreachable* Builder::Unreachable() {
-    return Append(ir.instructions.Create<ir::Unreachable>());
+    return Append(ir.allocators.instructions.Create<ir::Unreachable>());
 }
 
 const core::type::Type* Builder::VectorPtrElementType(const core::type::Type* type) {
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index 4593c79..c705d0c 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -261,7 +261,7 @@
     template <typename T>
     ir::If* If(T&& condition) {
         auto* cond_val = Value(std::forward<T>(condition));
-        return Append(ir.instructions.Create<ir::If>(cond_val, Block(), Block()));
+        return Append(ir.allocators.instructions.Create<ir::If>(cond_val, Block(), Block()));
     }
 
     /// Creates a loop instruction
@@ -274,7 +274,7 @@
     template <typename T>
     ir::Switch* Switch(T&& condition) {
         auto* cond_val = Value(std::forward<T>(condition));
-        return Append(ir.instructions.Create<ir::Switch>(cond_val));
+        return Append(ir.allocators.instructions.Create<ir::Switch>(cond_val));
     }
 
     /// Creates a default case for the switch @p s
@@ -298,7 +298,8 @@
     /// @param val the constant value
     /// @returns the new constant
     ir::Constant* Constant(const core::constant::Value* val) {
-        return ir.constants.GetOrAdd(val, [&] { return ir.values.Create<ir::Constant>(val); });
+        return ir.constants.GetOrAdd(
+            val, [&] { return ir.allocators.values.Create<ir::Constant>(val); });
     }
 
     /// Creates a ir::Constant for an i32 Scalar
@@ -476,8 +477,8 @@
         CheckForNonDeterministicEvaluation<LHS, RHS>();
         auto* lhs_val = Value(std::forward<LHS>(lhs));
         auto* rhs_val = Value(std::forward<RHS>(rhs));
-        return Append(
-            ir.instructions.Create<ir::CoreBinary>(InstructionResult(type), op, lhs_val, rhs_val));
+        return Append(ir.allocators.instructions.Create<ir::CoreBinary>(InstructionResult(type), op,
+                                                                        lhs_val, rhs_val));
     }
 
     /// Creates an And operation
@@ -826,7 +827,8 @@
     template <typename VAL>
     ir::CoreUnary* Unary(UnaryOp op, const core::type::Type* type, VAL&& val) {
         auto* value = Value(std::forward<VAL>(val));
-        return Append(ir.instructions.Create<ir::CoreUnary>(InstructionResult(type), op, value));
+        return Append(
+            ir.allocators.instructions.Create<ir::CoreUnary>(InstructionResult(type), op, value));
     }
 
     /// Creates an op for `op val`
@@ -908,7 +910,8 @@
     template <typename VAL>
     ir::Bitcast* Bitcast(const core::type::Type* type, VAL&& val) {
         auto* value = Value(std::forward<VAL>(val));
-        return Append(ir.instructions.Create<ir::Bitcast>(InstructionResult(type), value));
+        return Append(
+            ir.allocators.instructions.Create<ir::Bitcast>(InstructionResult(type), value));
     }
 
     /// Creates a bitcast instruction
@@ -942,8 +945,8 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::UserCall* Call(const core::type::Type* type, ir::Function* func, ARGS&&... args) {
-        return Append(ir.instructions.Create<ir::UserCall>(InstructionResult(type), func,
-                                                           Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::UserCall>(
+            InstructionResult(type), func, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates a user function call instruction
@@ -954,8 +957,8 @@
     template <typename TYPE, typename... ARGS>
     ir::UserCall* Call(ir::Function* func, ARGS&&... args) {
         auto* type = ir.Types().Get<TYPE>();
-        return Append(ir.instructions.Create<ir::UserCall>(InstructionResult(type), func,
-                                                           Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::UserCall>(
+            InstructionResult(type), func, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates a core builtin call instruction
@@ -965,7 +968,7 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::CoreBuiltinCall* Call(const core::type::Type* type, core::BuiltinFn func, ARGS&&... args) {
-        return Append(ir.instructions.Create<ir::CoreBuiltinCall>(
+        return Append(ir.allocators.instructions.Create<ir::CoreBuiltinCall>(
             InstructionResult(type), func, Values(std::forward<ARGS>(args)...)));
     }
 
@@ -977,7 +980,7 @@
     template <typename TYPE, typename... ARGS>
     ir::CoreBuiltinCall* Call(core::BuiltinFn func, ARGS&&... args) {
         auto* type = ir.Types().Get<TYPE>();
-        return Append(ir.instructions.Create<ir::CoreBuiltinCall>(
+        return Append(ir.allocators.instructions.Create<ir::CoreBuiltinCall>(
             InstructionResult(type), func, Values(std::forward<ARGS>(args)...)));
     }
 
@@ -989,8 +992,8 @@
     template <typename KLASS, typename FUNC, typename... ARGS>
     tint::traits::EnableIf<tint::traits::IsTypeOrDerived<KLASS, ir::BuiltinCall>, KLASS*>
     Call(const core::type::Type* type, FUNC func, ARGS&&... args) {
-        return Append(ir.instructions.Create<KLASS>(InstructionResult(type), func,
-                                                    Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<KLASS>(
+            InstructionResult(type), func, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates a value conversion instruction to the template type T
@@ -1008,8 +1011,8 @@
     /// @returns the instruction
     template <typename VAL>
     ir::Convert* Convert(const core::type::Type* to, VAL&& val) {
-        return Append(ir.instructions.Create<ir::Convert>(InstructionResult(to),
-                                                          Value(std::forward<VAL>(val))));
+        return Append(ir.allocators.instructions.Create<ir::Convert>(
+            InstructionResult(to), Value(std::forward<VAL>(val))));
     }
 
     /// Creates a value constructor instruction to the template type T
@@ -1027,8 +1030,8 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::Construct* Construct(const core::type::Type* type, ARGS&&... args) {
-        return Append(ir.instructions.Create<ir::Construct>(InstructionResult(type),
-                                                            Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::Construct>(
+            InstructionResult(type), Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates a load instruction
@@ -1037,7 +1040,7 @@
     template <typename VAL>
     ir::Load* Load(VAL&& from) {
         auto* value = Value(std::forward<VAL>(from));
-        return Append(ir.instructions.Create<ir::Load>(
+        return Append(ir.allocators.instructions.Create<ir::Load>(
             InstructionResult(value->Type()->UnwrapPtrOrRef()), value));
     }
 
@@ -1050,7 +1053,7 @@
         CheckForNonDeterministicEvaluation<TO, FROM>();
         auto* to_val = Value(std::forward<TO>(to));
         auto* from_val = Value(std::forward<FROM>(from));
-        return Append(ir.instructions.Create<ir::Store>(to_val, from_val));
+        return Append(ir.allocators.instructions.Create<ir::Store>(to_val, from_val));
     }
 
     /// Creates a store vector element instruction
@@ -1064,7 +1067,8 @@
         auto* to_val = Value(std::forward<TO>(to));
         auto* index_val = Value(std::forward<INDEX>(index));
         auto* value_val = Value(std::forward<VALUE>(value));
-        return Append(ir.instructions.Create<ir::StoreVectorElement>(to_val, index_val, value_val));
+        return Append(ir.allocators.instructions.Create<ir::StoreVectorElement>(to_val, index_val,
+                                                                                value_val));
     }
 
     /// Creates a load vector element instruction
@@ -1077,7 +1081,8 @@
         auto* from_val = Value(std::forward<FROM>(from));
         auto* index_val = Value(std::forward<INDEX>(index));
         auto* res = InstructionResult(VectorPtrElementType(from_val->Type()));
-        return Append(ir.instructions.Create<ir::LoadVectorElement>(res, from_val, index_val));
+        return Append(
+            ir.allocators.instructions.Create<ir::LoadVectorElement>(res, from_val, index_val));
     }
 
     /// Creates a new `var` declaration
@@ -1151,7 +1156,8 @@
             TINT_ASSERT(val);
             return nullptr;
         }
-        auto* let = Append(ir.instructions.Create<ir::Let>(InstructionResult(val->Type()), val));
+        auto* let =
+            Append(ir.allocators.instructions.Create<ir::Let>(InstructionResult(val->Type()), val));
         ir.SetName(let->Result(0), name);
         return let;
     }
@@ -1160,7 +1166,7 @@
     /// @param type the let type
     /// @returns the instruction
     ir::Let* Let(const type::Type* type) {
-        auto* let = ir.instructions.Create<ir::Let>(InstructionResult(type), nullptr);
+        auto* let = ir.allocators.instructions.Create<ir::Let>(InstructionResult(type), nullptr);
         Append(let);
         return let;
     }
@@ -1169,7 +1175,7 @@
     /// @param func the function being returned
     /// @returns the instruction
     ir::Return* Return(ir::Function* func) {
-        return Append(ir.instructions.Create<ir::Return>(func));
+        return Append(ir.allocators.instructions.Create<ir::Return>(func));
     }
 
     /// Creates a return instruction
@@ -1180,11 +1186,11 @@
     ir::Return* Return(ir::Function* func, ARG&& value) {
         if constexpr (std::is_same_v<std::decay_t<ARG>, ir::Value*>) {
             if (value == nullptr) {
-                return Append(ir.instructions.Create<ir::Return>(func));
+                return Append(ir.allocators.instructions.Create<ir::Return>(func));
             }
         }
         auto* val = Value(std::forward<ARG>(value));
-        return Append(ir.instructions.Create<ir::Return>(func, val));
+        return Append(ir.allocators.instructions.Create<ir::Return>(func, val));
     }
 
     /// Creates a loop next iteration instruction
@@ -1193,8 +1199,8 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::NextIteration* NextIteration(ir::Loop* loop, ARGS&&... args) {
-        return Append(
-            ir.instructions.Create<ir::NextIteration>(loop, Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::NextIteration>(
+            loop, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates a loop break-if instruction
@@ -1206,8 +1212,8 @@
     ir::BreakIf* BreakIf(ir::Loop* loop, CONDITION&& condition, ARGS&&... args) {
         CheckForNonDeterministicEvaluation<CONDITION, ARGS...>();
         auto* cond_val = Value(std::forward<CONDITION>(condition));
-        return Append(ir.instructions.Create<ir::BreakIf>(cond_val, loop,
-                                                          Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::BreakIf>(
+            cond_val, loop, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates a continue instruction
@@ -1216,8 +1222,8 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::Continue* Continue(ir::Loop* loop, ARGS&&... args) {
-        return Append(
-            ir.instructions.Create<ir::Continue>(loop, Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::Continue>(
+            loop, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates an exit switch instruction
@@ -1226,8 +1232,8 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::ExitSwitch* ExitSwitch(ir::Switch* sw, ARGS&&... args) {
-        return Append(
-            ir.instructions.Create<ir::ExitSwitch>(sw, Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::ExitSwitch>(
+            sw, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates an exit loop instruction
@@ -1236,8 +1242,8 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::ExitLoop* ExitLoop(ir::Loop* loop, ARGS&&... args) {
-        return Append(
-            ir.instructions.Create<ir::ExitLoop>(loop, Values(std::forward<ARGS>(args)...)));
+        return Append(ir.allocators.instructions.Create<ir::ExitLoop>(
+            loop, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates an exit if instruction
@@ -1246,7 +1252,8 @@
     /// @returns the instruction
     template <typename... ARGS>
     ir::ExitIf* ExitIf(ir::If* i, ARGS&&... args) {
-        return Append(ir.instructions.Create<ir::ExitIf>(i, Values(std::forward<ARGS>(args)...)));
+        return Append(
+            ir.allocators.instructions.Create<ir::ExitIf>(i, Values(std::forward<ARGS>(args)...)));
     }
 
     /// Creates an exit instruction for the given control instruction
@@ -1322,8 +1329,8 @@
     ir::Access* Access(const core::type::Type* type, OBJ&& object, ARGS&&... indices) {
         CheckForNonDeterministicEvaluation<OBJ, ARGS...>();
         auto* obj_val = Value(std::forward<OBJ>(object));
-        return Append(ir.instructions.Create<ir::Access>(InstructionResult(type), obj_val,
-                                                         Values(std::forward<ARGS>(indices)...)));
+        return Append(ir.allocators.instructions.Create<ir::Access>(
+            InstructionResult(type), obj_val, Values(std::forward<ARGS>(indices)...)));
     }
 
     /// Creates a new `Access`
@@ -1345,8 +1352,8 @@
     template <typename OBJ>
     ir::Swizzle* Swizzle(const core::type::Type* type, OBJ&& object, VectorRef<uint32_t> indices) {
         auto* obj_val = Value(std::forward<OBJ>(object));
-        return Append(ir.instructions.Create<ir::Swizzle>(InstructionResult(type), obj_val,
-                                                          std::move(indices)));
+        return Append(ir.allocators.instructions.Create<ir::Swizzle>(InstructionResult(type),
+                                                                     obj_val, std::move(indices)));
     }
 
     /// Creates a new `Swizzle`
@@ -1370,8 +1377,8 @@
                          OBJ&& object,
                          std::initializer_list<uint32_t> indices) {
         auto* obj_val = Value(std::forward<OBJ>(object));
-        return Append(ir.instructions.Create<ir::Swizzle>(InstructionResult(type), obj_val,
-                                                          Vector<uint32_t, 4>(indices)));
+        return Append(ir.allocators.instructions.Create<ir::Swizzle>(
+            InstructionResult(type), obj_val, Vector<uint32_t, 4>(indices)));
     }
 
     /// Name names the value or instruction with @p name
@@ -1396,7 +1403,7 @@
     /// @param type the return type
     /// @returns the value
     ir::InstructionResult* InstructionResult(const core::type::Type* type) {
-        return ir.values.Create<ir::InstructionResult>(type);
+        return ir.allocators.values.Create<ir::InstructionResult>(type);
     }
 
     /// Creates a new runtime value
diff --git a/src/tint/lang/core/ir/construct.cc b/src/tint/lang/core/ir/construct.cc
index 8840b1f..d92ea6f 100644
--- a/src/tint/lang/core/ir/construct.cc
+++ b/src/tint/lang/core/ir/construct.cc
@@ -48,7 +48,7 @@
 Construct* Construct::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto args = ctx.Remap<Construct::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<Construct>(new_result, args);
+    return ctx.ir.allocators.instructions.Create<Construct>(new_result, args);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/continue.cc b/src/tint/lang/core/ir/continue.cc
index 66c9115..3b58584 100644
--- a/src/tint/lang/core/ir/continue.cc
+++ b/src/tint/lang/core/ir/continue.cc
@@ -58,7 +58,7 @@
     auto* loop = ctx.Remap(Loop());
     auto args = ctx.Remap<Continue::kDefaultNumOperands>(Args());
 
-    return ctx.ir.instructions.Create<Continue>(loop, args);
+    return ctx.ir.allocators.instructions.Create<Continue>(loop, args);
 }
 
 void Continue::SetLoop(ir::Loop* loop) {
diff --git a/src/tint/lang/core/ir/convert.cc b/src/tint/lang/core/ir/convert.cc
index 3c4ac56..985c47d 100644
--- a/src/tint/lang/core/ir/convert.cc
+++ b/src/tint/lang/core/ir/convert.cc
@@ -48,7 +48,7 @@
 Convert* Convert::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto* val = ctx.Remap(Args()[0]);
-    return ctx.ir.instructions.Create<Convert>(new_result, val);
+    return ctx.ir.allocators.instructions.Create<Convert>(new_result, val);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/core_binary.cc b/src/tint/lang/core/ir/core_binary.cc
index 0bab152..ef3cac4 100644
--- a/src/tint/lang/core/ir/core_binary.cc
+++ b/src/tint/lang/core/ir/core_binary.cc
@@ -46,7 +46,7 @@
     auto* new_result = ctx.Clone(Result(0));
     auto* lhs = ctx.Remap(LHS());
     auto* rhs = ctx.Remap(RHS());
-    return ctx.ir.instructions.Create<CoreBinary>(new_result, Op(), lhs, rhs);
+    return ctx.ir.allocators.instructions.Create<CoreBinary>(new_result, Op(), lhs, rhs);
 }
 
 const core::intrinsic::TableData& CoreBinary::TableData() const {
diff --git a/src/tint/lang/core/ir/core_builtin_call.cc b/src/tint/lang/core/ir/core_builtin_call.cc
index b8114ba..e095f1a 100644
--- a/src/tint/lang/core/ir/core_builtin_call.cc
+++ b/src/tint/lang/core/ir/core_builtin_call.cc
@@ -51,7 +51,7 @@
 CoreBuiltinCall* CoreBuiltinCall::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto args = ctx.Remap<CoreBuiltinCall::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<CoreBuiltinCall>(new_result, func_, args);
+    return ctx.ir.allocators.instructions.Create<CoreBuiltinCall>(new_result, func_, args);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/core_unary.cc b/src/tint/lang/core/ir/core_unary.cc
index 5ab130d..e01a261 100644
--- a/src/tint/lang/core/ir/core_unary.cc
+++ b/src/tint/lang/core/ir/core_unary.cc
@@ -44,7 +44,7 @@
 CoreUnary* CoreUnary::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto* val = ctx.Remap(Val());
-    return ctx.ir.instructions.Create<CoreUnary>(new_result, Op(), val);
+    return ctx.ir.allocators.instructions.Create<CoreUnary>(new_result, Op(), val);
 }
 
 const core::intrinsic::TableData& CoreUnary::TableData() const {
diff --git a/src/tint/lang/core/ir/discard.cc b/src/tint/lang/core/ir/discard.cc
index 25252d3..53cc77f 100644
--- a/src/tint/lang/core/ir/discard.cc
+++ b/src/tint/lang/core/ir/discard.cc
@@ -41,7 +41,7 @@
 Discard::~Discard() = default;
 
 Discard* Discard::Clone(CloneContext& ctx) {
-    return ctx.ir.instructions.Create<Discard>();
+    return ctx.ir.allocators.instructions.Create<Discard>();
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/exit_if.cc b/src/tint/lang/core/ir/exit_if.cc
index 5130495..b848a86 100644
--- a/src/tint/lang/core/ir/exit_if.cc
+++ b/src/tint/lang/core/ir/exit_if.cc
@@ -50,7 +50,7 @@
 ExitIf* ExitIf::Clone(CloneContext& ctx) {
     auto* if_ = ctx.Remap(If());
     auto args = ctx.Remap<ExitIf::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<ExitIf>(if_, args);
+    return ctx.ir.allocators.instructions.Create<ExitIf>(if_, args);
 }
 
 void ExitIf::SetIf(ir::If* i) {
diff --git a/src/tint/lang/core/ir/exit_loop.cc b/src/tint/lang/core/ir/exit_loop.cc
index 31289b5f..2222659 100644
--- a/src/tint/lang/core/ir/exit_loop.cc
+++ b/src/tint/lang/core/ir/exit_loop.cc
@@ -51,7 +51,7 @@
 ExitLoop* ExitLoop::Clone(CloneContext& ctx) {
     auto* loop = ctx.Remap(Loop());
     auto args = ctx.Remap<ExitLoop::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<ExitLoop>(loop, args);
+    return ctx.ir.allocators.instructions.Create<ExitLoop>(loop, args);
 }
 
 void ExitLoop::SetLoop(ir::Loop* l) {
diff --git a/src/tint/lang/core/ir/exit_switch.cc b/src/tint/lang/core/ir/exit_switch.cc
index 0ac5cc2..07363e7 100644
--- a/src/tint/lang/core/ir/exit_switch.cc
+++ b/src/tint/lang/core/ir/exit_switch.cc
@@ -50,7 +50,7 @@
 ExitSwitch* ExitSwitch::Clone(CloneContext& ctx) {
     auto* switch_ = ctx.Remap(Switch());
     auto args = ctx.Remap<ExitSwitch::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<ExitSwitch>(switch_, args);
+    return ctx.ir.allocators.instructions.Create<ExitSwitch>(switch_, args);
 }
 
 void ExitSwitch::SetSwitch(ir::Switch* s) {
diff --git a/src/tint/lang/core/ir/function.cc b/src/tint/lang/core/ir/function.cc
index ebe4211..51f7928 100644
--- a/src/tint/lang/core/ir/function.cc
+++ b/src/tint/lang/core/ir/function.cc
@@ -50,7 +50,8 @@
 Function::~Function() = default;
 
 Function* Function::Clone(CloneContext& ctx) {
-    auto* new_func = ctx.ir.values.Create<Function>(return_.type, pipeline_stage_, workgroup_size_);
+    auto* new_func =
+        ctx.ir.allocators.values.Create<Function>(return_.type, pipeline_stage_, workgroup_size_);
     new_func->block_ = ctx.ir.blocks.Create<ir::Block>();
     new_func->params_ = ctx.Clone<1>(params_.Slice());
     new_func->return_.builtin = return_.builtin;
diff --git a/src/tint/lang/core/ir/function_param.cc b/src/tint/lang/core/ir/function_param.cc
index 6dd00df..69b1dc0 100644
--- a/src/tint/lang/core/ir/function_param.cc
+++ b/src/tint/lang/core/ir/function_param.cc
@@ -42,7 +42,7 @@
 FunctionParam::~FunctionParam() = default;
 
 FunctionParam* FunctionParam::Clone(CloneContext& ctx) {
-    auto* out = ctx.ir.values.Create<FunctionParam>(type_);
+    auto* out = ctx.ir.allocators.values.Create<FunctionParam>(type_);
     out->builtin_ = builtin_;
     out->location_ = location_;
     out->binding_point_ = binding_point_;
diff --git a/src/tint/lang/core/ir/if.cc b/src/tint/lang/core/ir/if.cc
index f50abf0..be25a7c 100644
--- a/src/tint/lang/core/ir/if.cc
+++ b/src/tint/lang/core/ir/if.cc
@@ -68,7 +68,7 @@
     auto* new_true = ctx.ir.blocks.Create<ir::Block>();
     auto* new_false = ctx.ir.blocks.Create<ir::Block>();
 
-    auto* new_if = ctx.ir.instructions.Create<If>(cond, new_true, new_false);
+    auto* new_if = ctx.ir.allocators.instructions.Create<If>(cond, new_true, new_false);
     ctx.Replace(this, new_if);
 
     true_->CloneInto(ctx, new_true);
diff --git a/src/tint/lang/core/ir/instruction_result.cc b/src/tint/lang/core/ir/instruction_result.cc
index 2aa0fc7..6960792 100644
--- a/src/tint/lang/core/ir/instruction_result.cc
+++ b/src/tint/lang/core/ir/instruction_result.cc
@@ -51,7 +51,7 @@
 InstructionResult* InstructionResult::Clone(CloneContext& ctx) {
     // Do not clone the `Instruction`. It will be set when this result is placed in the new parent
     // instruction.
-    return ctx.ir.values.Create<InstructionResult>(type_);
+    return ctx.ir.allocators.values.Create<InstructionResult>(type_);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/let.cc b/src/tint/lang/core/ir/let.cc
index aeb185e..f0317e1 100644
--- a/src/tint/lang/core/ir/let.cc
+++ b/src/tint/lang/core/ir/let.cc
@@ -47,7 +47,7 @@
 Let* Let::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto* val = ctx.Remap(Value());
-    auto* new_let = ctx.ir.instructions.Create<Let>(new_result, val);
+    auto* new_let = ctx.ir.allocators.instructions.Create<Let>(new_result, val);
 
     auto name = ctx.ir.NameOf(this);
     ctx.ir.SetName(new_let, name.Name());
diff --git a/src/tint/lang/core/ir/load.cc b/src/tint/lang/core/ir/load.cc
index 1cdca31..77007cc 100644
--- a/src/tint/lang/core/ir/load.cc
+++ b/src/tint/lang/core/ir/load.cc
@@ -52,7 +52,7 @@
 Load* Load::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto* from = ctx.Remap(From());
-    return ctx.ir.instructions.Create<Load>(new_result, from);
+    return ctx.ir.allocators.instructions.Create<Load>(new_result, from);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/load_vector_element.cc b/src/tint/lang/core/ir/load_vector_element.cc
index 750a72e..29889e7 100644
--- a/src/tint/lang/core/ir/load_vector_element.cc
+++ b/src/tint/lang/core/ir/load_vector_element.cc
@@ -52,7 +52,7 @@
     auto* new_result = ctx.Clone(Result(0));
     auto* from = ctx.Remap(From());
     auto* index = ctx.Remap(Index());
-    return ctx.ir.instructions.Create<LoadVectorElement>(new_result, from, index);
+    return ctx.ir.allocators.instructions.Create<LoadVectorElement>(new_result, from, index);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/loop.cc b/src/tint/lang/core/ir/loop.cc
index 64a5a6c..a2c8e83 100644
--- a/src/tint/lang/core/ir/loop.cc
+++ b/src/tint/lang/core/ir/loop.cc
@@ -62,7 +62,8 @@
     auto* new_body = ctx.ir.blocks.Create<MultiInBlock>();
     auto* new_continuing = ctx.ir.blocks.Create<MultiInBlock>();
 
-    auto* new_loop = ctx.ir.instructions.Create<Loop>(new_init, new_body, new_continuing);
+    auto* new_loop =
+        ctx.ir.allocators.instructions.Create<Loop>(new_init, new_body, new_continuing);
     ctx.Replace(this, new_loop);
 
     initializer_->CloneInto(ctx, new_init);
diff --git a/src/tint/lang/core/ir/module.h b/src/tint/lang/core/ir/module.h
index fba89fb..f4a93a1 100644
--- a/src/tint/lang/core/ir/module.h
+++ b/src/tint/lang/core/ir/module.h
@@ -39,6 +39,7 @@
 #include "src/tint/lang/core/ir/value.h"
 #include "src/tint/lang/core/type/manager.h"
 #include "src/tint/utils/containers/const_propagating_ptr.h"
+#include "src/tint/utils/containers/filtered_iterator.h"
 #include "src/tint/utils/containers/vector.h"
 #include "src/tint/utils/diagnostic/source.h"
 #include "src/tint/utils/id/generation_id.h"
@@ -56,6 +57,12 @@
     /// Map of value to name
     Hashmap<const Value*, Symbol, 32> value_to_name_;
 
+    /// A predicate function that returns true if the instruction or value is alive.
+    struct IsAlive {
+        bool operator()(const Instruction* instruction) const { return instruction->Alive(); }
+        bool operator()(const Value* value) const { return value->Alive(); }
+    };
+
   public:
     /// Constructor
     Module();
@@ -102,17 +109,40 @@
     /// @return the type manager for the module
     const core::type::Manager& Types() const { return constant_values.types; }
 
+    /// @returns a iterable of all the alive instructions
+    FilteredIterable<IsAlive, BlockAllocator<Instruction>::View> Instructions() {
+        return {allocators.instructions.Objects()};
+    }
+
+    /// @returns a iterable of all the alive instructions
+    FilteredIterable<IsAlive, BlockAllocator<Instruction>::ConstView> Instructions() const {
+        return {allocators.instructions.Objects()};
+    }
+
+    /// @returns a iterable of all the alive values
+    FilteredIterable<IsAlive, BlockAllocator<Value>::View> Values() {
+        return {allocators.values.Objects()};
+    }
+
+    /// @returns a iterable of all the alive values
+    FilteredIterable<IsAlive, BlockAllocator<Value>::ConstView> Values() const {
+        return {allocators.values.Objects()};
+    }
+
     /// The block allocator
     BlockAllocator<Block> blocks;
 
     /// The constant value manager
     core::constant::Manager constant_values;
 
-    /// The instruction allocator
-    BlockAllocator<Instruction> instructions;
+    /// The various BlockAllocators for the module
+    struct {
+        /// The instruction allocator
+        BlockAllocator<Instruction> instructions;
 
-    /// The value allocator
-    BlockAllocator<Value> values;
+        /// The value allocator
+        BlockAllocator<Value> values;
+    } allocators;
 
     /// List of functions in the program
     Vector<ConstPropagatingPtr<Function>, 8> functions;
diff --git a/src/tint/lang/core/ir/next_iteration.cc b/src/tint/lang/core/ir/next_iteration.cc
index 5de66c7..5fc404f 100644
--- a/src/tint/lang/core/ir/next_iteration.cc
+++ b/src/tint/lang/core/ir/next_iteration.cc
@@ -57,7 +57,7 @@
 NextIteration* NextIteration::Clone(CloneContext& ctx) {
     auto* new_loop = ctx.Clone(loop_);
     auto args = ctx.Remap<NextIteration::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<NextIteration>(new_loop, args);
+    return ctx.ir.allocators.instructions.Create<NextIteration>(new_loop, args);
 }
 
 void NextIteration::SetLoop(ir::Loop* loop) {
diff --git a/src/tint/lang/core/ir/return.cc b/src/tint/lang/core/ir/return.cc
index fb6cd8a..d236eae 100644
--- a/src/tint/lang/core/ir/return.cc
+++ b/src/tint/lang/core/ir/return.cc
@@ -53,9 +53,9 @@
 Return* Return::Clone(CloneContext& ctx) {
     auto* fn = ctx.Remap(Func());
     if (auto* val = Value()) {
-        return ctx.ir.instructions.Create<Return>(fn, ctx.Remap(val));
+        return ctx.ir.allocators.instructions.Create<Return>(fn, ctx.Remap(val));
     }
-    return ctx.ir.instructions.Create<Return>(fn);
+    return ctx.ir.allocators.instructions.Create<Return>(fn);
 }
 
 Function* Return::Func() {
diff --git a/src/tint/lang/core/ir/store.cc b/src/tint/lang/core/ir/store.cc
index 2f6ab39..388646b 100644
--- a/src/tint/lang/core/ir/store.cc
+++ b/src/tint/lang/core/ir/store.cc
@@ -50,7 +50,7 @@
 Store* Store::Clone(CloneContext& ctx) {
     auto* to = ctx.Remap(To());
     auto* from = ctx.Remap(From());
-    return ctx.ir.instructions.Create<Store>(to, from);
+    return ctx.ir.allocators.instructions.Create<Store>(to, from);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/store_vector_element.cc b/src/tint/lang/core/ir/store_vector_element.cc
index 947c28e..16b6f87 100644
--- a/src/tint/lang/core/ir/store_vector_element.cc
+++ b/src/tint/lang/core/ir/store_vector_element.cc
@@ -52,7 +52,7 @@
     auto* to = ctx.Remap(To());
     auto* idx = ctx.Remap(Index());
     auto* val = ctx.Remap(Value());
-    return ctx.ir.instructions.Create<StoreVectorElement>(to, idx, val);
+    return ctx.ir.allocators.instructions.Create<StoreVectorElement>(to, idx, val);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/switch.cc b/src/tint/lang/core/ir/switch.cc
index 1b0d2bc..a124daa 100644
--- a/src/tint/lang/core/ir/switch.cc
+++ b/src/tint/lang/core/ir/switch.cc
@@ -55,7 +55,7 @@
 
 Switch* Switch::Clone(CloneContext& ctx) {
     auto* cond = ctx.Remap(Condition());
-    auto* new_switch = ctx.ir.instructions.Create<Switch>(cond);
+    auto* new_switch = ctx.ir.allocators.instructions.Create<Switch>(cond);
     ctx.Replace(this, new_switch);
 
     new_switch->cases_.Reserve(cases_.Length());
diff --git a/src/tint/lang/core/ir/swizzle.cc b/src/tint/lang/core/ir/swizzle.cc
index 184ddf6..4fd3c0c 100644
--- a/src/tint/lang/core/ir/swizzle.cc
+++ b/src/tint/lang/core/ir/swizzle.cc
@@ -57,7 +57,7 @@
 Swizzle* Swizzle::Clone(CloneContext& ctx) {
     auto* result = ctx.Clone(Result(0));
     auto* obj = ctx.Remap(Object());
-    return ctx.ir.instructions.Create<Swizzle>(result, obj, indices_);
+    return ctx.ir.allocators.instructions.Create<Swizzle>(result, obj, indices_);
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/terminate_invocation.cc b/src/tint/lang/core/ir/terminate_invocation.cc
index 6f87557..97bf084 100644
--- a/src/tint/lang/core/ir/terminate_invocation.cc
+++ b/src/tint/lang/core/ir/terminate_invocation.cc
@@ -37,7 +37,7 @@
 TerminateInvocation::~TerminateInvocation() = default;
 
 TerminateInvocation* TerminateInvocation::Clone(CloneContext& ctx) {
-    return ctx.ir.instructions.Create<TerminateInvocation>();
+    return ctx.ir.allocators.instructions.Create<TerminateInvocation>();
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/transform/binary_polyfill.cc b/src/tint/lang/core/ir/transform/binary_polyfill.cc
index 27d3b13..47c6861 100644
--- a/src/tint/lang/core/ir/transform/binary_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/binary_polyfill.cc
@@ -67,10 +67,7 @@
     void Process() {
         // Find the binary instructions that need to be polyfilled.
         Vector<ir::CoreBinary*, 64> worklist;
-        for (auto* inst : ir.instructions.Objects()) {
-            if (!inst->Alive()) {
-                continue;
-            }
+        for (auto* inst : ir.Instructions()) {
             if (auto* binary = inst->As<ir::CoreBinary>()) {
                 switch (binary->Op()) {
                     case BinaryOp::kDivide:
diff --git a/src/tint/lang/core/ir/transform/builtin_polyfill.cc b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
index 732fba8..2e91f34 100644
--- a/src/tint/lang/core/ir/transform/builtin_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/builtin_polyfill.cc
@@ -62,10 +62,7 @@
     void Process() {
         // Find the builtin call instructions that may need to be polyfilled.
         Vector<ir::CoreBuiltinCall*, 4> worklist;
-        for (auto* inst : ir.instructions.Objects()) {
-            if (!inst->Alive()) {
-                continue;
-            }
+        for (auto* inst : ir.Instructions()) {
             if (auto* builtin = inst->As<ir::CoreBuiltinCall>()) {
                 switch (builtin->Func()) {
                     case core::BuiltinFn::kClamp:
diff --git a/src/tint/lang/core/ir/transform/combine_access_instructions.cc b/src/tint/lang/core/ir/transform/combine_access_instructions.cc
index 7d82da4..1e9f6d0 100644
--- a/src/tint/lang/core/ir/transform/combine_access_instructions.cc
+++ b/src/tint/lang/core/ir/transform/combine_access_instructions.cc
@@ -48,8 +48,8 @@
     /// Process the module.
     void Process() {
         // Loop over every instruction looking for access instructions.
-        for (auto* inst : ir.instructions.Objects()) {
-            if (auto* access = inst->As<ir::Access>(); access && access->Alive()) {
+        for (auto* inst : ir.Instructions()) {
+            if (auto* access = inst->As<ir::Access>()) {
                 // Look for places where the result of this access instruction is used as a base
                 // pointer for another access instruction.
                 access->Result(0)->ForEachUse([&](Usage use) {
diff --git a/src/tint/lang/core/ir/transform/conversion_polyfill.cc b/src/tint/lang/core/ir/transform/conversion_polyfill.cc
index 5d4fdbc..c5f7bbd 100644
--- a/src/tint/lang/core/ir/transform/conversion_polyfill.cc
+++ b/src/tint/lang/core/ir/transform/conversion_polyfill.cc
@@ -68,10 +68,7 @@
     void Process() {
         // Find the conversion instructions that need to be polyfilled.
         Vector<ir::Convert*, 64> ftoi_worklist;
-        for (auto* inst : ir.instructions.Objects()) {
-            if (!inst->Alive()) {
-                continue;
-            }
+        for (auto* inst : ir.Instructions()) {
             if (auto* convert = inst->As<ir::Convert>()) {
                 auto* src_ty = convert->Args()[0]->Type();
                 auto* res_ty = convert->Result(0)->Type();
diff --git a/src/tint/lang/core/ir/transform/preserve_padding.cc b/src/tint/lang/core/ir/transform/preserve_padding.cc
index 79e6588..550d67d 100644
--- a/src/tint/lang/core/ir/transform/preserve_padding.cc
+++ b/src/tint/lang/core/ir/transform/preserve_padding.cc
@@ -61,8 +61,8 @@
     void Process() {
         // Find host-visible stores of types that contain padding bytes.
         Vector<Store*, 8> worklist;
-        for (auto inst : ir.instructions.Objects()) {
-            if (auto* store = inst->As<Store>(); store && store->Alive()) {
+        for (auto inst : ir.Instructions()) {
+            if (auto* store = inst->As<Store>()) {
                 auto* ptr = store->To()->Type()->As<core::type::Pointer>();
                 if (ptr->AddressSpace() == core::AddressSpace::kStorage &&
                     ContainsPadding(ptr->StoreType())) {
diff --git a/src/tint/lang/core/ir/transform/robustness.cc b/src/tint/lang/core/ir/transform/robustness.cc
index 5791b9b..43cd948 100644
--- a/src/tint/lang/core/ir/transform/robustness.cc
+++ b/src/tint/lang/core/ir/transform/robustness.cc
@@ -65,48 +65,46 @@
         Vector<ir::LoadVectorElement*, 64> vector_loads;
         Vector<ir::StoreVectorElement*, 64> vector_stores;
         Vector<ir::CoreBuiltinCall*, 64> texture_calls;
-        for (auto* inst : ir.instructions.Objects()) {
-            if (inst->Alive()) {
-                tint::Switch(
-                    inst,  //
-                    [&](ir::Access* access) {
-                        // Check if accesses into this object should be clamped.
-                        auto* ptr = access->Object()->Type()->As<type::Pointer>();
-                        if (ptr) {
-                            if (ShouldClamp(ptr->AddressSpace())) {
-                                accesses.Push(access);
-                            }
-                        } else {
-                            if (config.clamp_value) {
-                                accesses.Push(access);
-                            }
-                        }
-                    },
-                    [&](ir::LoadVectorElement* lve) {
-                        // Check if loads from this address space should be clamped.
-                        auto* ptr = lve->From()->Type()->As<type::Pointer>();
+        for (auto* inst : ir.Instructions()) {
+            tint::Switch(
+                inst,  //
+                [&](ir::Access* access) {
+                    // Check if accesses into this object should be clamped.
+                    auto* ptr = access->Object()->Type()->As<type::Pointer>();
+                    if (ptr) {
                         if (ShouldClamp(ptr->AddressSpace())) {
-                            vector_loads.Push(lve);
+                            accesses.Push(access);
                         }
-                    },
-                    [&](ir::StoreVectorElement* sve) {
-                        // Check if stores to this address space should be clamped.
-                        auto* ptr = sve->To()->Type()->As<type::Pointer>();
-                        if (ShouldClamp(ptr->AddressSpace())) {
-                            vector_stores.Push(sve);
+                    } else {
+                        if (config.clamp_value) {
+                            accesses.Push(access);
                         }
-                    },
-                    [&](ir::CoreBuiltinCall* call) {
-                        // Check if this is a texture builtin that needs to be clamped.
-                        if (config.clamp_texture) {
-                            if (call->Func() == core::BuiltinFn::kTextureDimensions ||
-                                call->Func() == core::BuiltinFn::kTextureLoad ||
-                                call->Func() == core::BuiltinFn::kTextureStore) {
-                                texture_calls.Push(call);
-                            }
+                    }
+                },
+                [&](ir::LoadVectorElement* lve) {
+                    // Check if loads from this address space should be clamped.
+                    auto* ptr = lve->From()->Type()->As<type::Pointer>();
+                    if (ShouldClamp(ptr->AddressSpace())) {
+                        vector_loads.Push(lve);
+                    }
+                },
+                [&](ir::StoreVectorElement* sve) {
+                    // Check if stores to this address space should be clamped.
+                    auto* ptr = sve->To()->Type()->As<type::Pointer>();
+                    if (ShouldClamp(ptr->AddressSpace())) {
+                        vector_stores.Push(sve);
+                    }
+                },
+                [&](ir::CoreBuiltinCall* call) {
+                    // Check if this is a texture builtin that needs to be clamped.
+                    if (config.clamp_texture) {
+                        if (call->Func() == core::BuiltinFn::kTextureDimensions ||
+                            call->Func() == core::BuiltinFn::kTextureLoad ||
+                            call->Func() == core::BuiltinFn::kTextureStore) {
+                            texture_calls.Push(call);
                         }
-                    });
-            }
+                    }
+                });
         }
 
         // Clamp access indices.
diff --git a/src/tint/lang/core/ir/transform/vectorize_scalar_matrix_constructors.cc b/src/tint/lang/core/ir/transform/vectorize_scalar_matrix_constructors.cc
index a9e2de6..b8c7703 100644
--- a/src/tint/lang/core/ir/transform/vectorize_scalar_matrix_constructors.cc
+++ b/src/tint/lang/core/ir/transform/vectorize_scalar_matrix_constructors.cc
@@ -55,8 +55,8 @@
     void Process() {
         // Find and replace matrix constructors that take scalar operands.
         Vector<Construct*, 8> worklist;
-        for (auto inst : ir.instructions.Objects()) {
-            if (auto* construct = inst->As<Construct>(); construct && construct->Alive()) {
+        for (auto inst : ir.Instructions()) {
+            if (auto* construct = inst->As<Construct>()) {
                 if (construct->Result(0)->Type()->As<type::Matrix>()) {
                     if (construct->Operands().Length() > 0 &&
                         construct->Operands()[0]->Type()->Is<type::Scalar>()) {
diff --git a/src/tint/lang/core/ir/unreachable.cc b/src/tint/lang/core/ir/unreachable.cc
index 47453f6..b4146ad 100644
--- a/src/tint/lang/core/ir/unreachable.cc
+++ b/src/tint/lang/core/ir/unreachable.cc
@@ -37,7 +37,7 @@
 Unreachable::~Unreachable() = default;
 
 Unreachable* Unreachable::Clone(CloneContext& ctx) {
-    return ctx.ir.instructions.Create<Unreachable>();
+    return ctx.ir.allocators.instructions.Create<Unreachable>();
 }
 
 }  // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/user_call.cc b/src/tint/lang/core/ir/user_call.cc
index 6cadb32..4da83b6 100644
--- a/src/tint/lang/core/ir/user_call.cc
+++ b/src/tint/lang/core/ir/user_call.cc
@@ -53,7 +53,7 @@
     auto* new_result = ctx.Clone(Result(0));
     auto* target = ctx.Remap(Target());
     auto args = ctx.Remap<UserCall::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<UserCall>(new_result, target, args);
+    return ctx.ir.allocators.instructions.Create<UserCall>(new_result, target, args);
 }
 
 void UserCall::SetArgs(VectorRef<Value*> arguments) {
diff --git a/src/tint/lang/core/ir/validator.cc b/src/tint/lang/core/ir/validator.cc
index 0a61365..3f1a681 100644
--- a/src/tint/lang/core/ir/validator.cc
+++ b/src/tint/lang/core/ir/validator.cc
@@ -357,8 +357,8 @@
 
     if (!diagnostics_.ContainsErrors()) {
         // Check for orphaned instructions.
-        for (auto* inst : mod_.instructions.Objects()) {
-            if (inst->Alive() && !visited_instructions_.Contains(inst)) {
+        for (auto* inst : mod_.Instructions()) {
+            if (!visited_instructions_.Contains(inst)) {
                 AddError(inst) << "orphaned instruction: " << inst->FriendlyName();
             }
         }
diff --git a/src/tint/lang/core/ir/validator_test.cc b/src/tint/lang/core/ir/validator_test.cc
index e049e4d..3e1a708 100644
--- a/src/tint/lang/core/ir/validator_test.cc
+++ b/src/tint/lang/core/ir/validator_test.cc
@@ -1145,7 +1145,7 @@
 }
 
 TEST_F(IR_ValidatorTest, Var_RootBlock_NullResult) {
-    auto* v = mod.instructions.Create<ir::Var>(nullptr);
+    auto* v = mod.allocators.instructions.Create<ir::Var>(nullptr);
     mod.root_block->Append(v);
 
     auto res = ir::Validate(mod);
@@ -1167,7 +1167,7 @@
 }
 
 TEST_F(IR_ValidatorTest, Var_Function_NullResult) {
-    auto* v = mod.instructions.Create<ir::Var>(nullptr);
+    auto* v = mod.allocators.instructions.Create<ir::Var>(nullptr);
 
     auto* f = b.Function("my_func", ty.void_());
 
@@ -1226,7 +1226,7 @@
 }
 
 TEST_F(IR_ValidatorTest, Let_NullResult) {
-    auto* v = mod.instructions.Create<ir::Let>(nullptr, b.Constant(1_i));
+    auto* v = mod.allocators.instructions.Create<ir::Let>(nullptr, b.Constant(1_i));
 
     auto* f = b.Function("my_func", ty.void_());
 
@@ -1255,7 +1255,7 @@
 }
 
 TEST_F(IR_ValidatorTest, Let_NullValue) {
-    auto* v = mod.instructions.Create<ir::Let>(b.InstructionResult(ty.f32()), nullptr);
+    auto* v = mod.allocators.instructions.Create<ir::Let>(b.InstructionResult(ty.f32()), nullptr);
 
     auto* f = b.Function("my_func", ty.void_());
 
@@ -1284,7 +1284,8 @@
 }
 
 TEST_F(IR_ValidatorTest, Let_WrongType) {
-    auto* v = mod.instructions.Create<ir::Let>(b.InstructionResult(ty.f32()), b.Constant(1_i));
+    auto* v =
+        mod.allocators.instructions.Create<ir::Let>(b.InstructionResult(ty.f32()), b.Constant(1_i));
 
     auto* f = b.Function("my_func", ty.void_());
 
@@ -1520,8 +1521,8 @@
 }
 
 TEST_F(IR_ValidatorTest, Binary_Result_Nullptr) {
-    auto* bin = mod.instructions.Create<ir::CoreBinary>(nullptr, BinaryOp::kAdd, b.Constant(3_i),
-                                                        b.Constant(2_i));
+    auto* bin = mod.allocators.instructions.Create<ir::CoreBinary>(
+        nullptr, BinaryOp::kAdd, b.Constant(3_i), b.Constant(2_i));
 
     auto* f = b.Function("my_func", ty.void_());
 
@@ -1577,8 +1578,8 @@
 }
 
 TEST_F(IR_ValidatorTest, Unary_Result_Nullptr) {
-    auto* bin =
-        mod.instructions.Create<ir::CoreUnary>(nullptr, UnaryOp::kNegation, b.Constant(2_i));
+    auto* bin = mod.allocators.instructions.Create<ir::CoreUnary>(nullptr, UnaryOp::kNegation,
+                                                                  b.Constant(2_i));
 
     auto* f = b.Function("my_func", ty.void_());
 
@@ -1651,7 +1652,7 @@
 
 TEST_F(IR_ValidatorTest, ExitIf_NullIf) {
     auto* if_ = b.If(true);
-    if_->True()->Append(mod.instructions.Create<ExitIf>(nullptr));
+    if_->True()->Append(mod.allocators.instructions.Create<ExitIf>(nullptr));
 
     auto* f = b.Function("my_func", ty.void_());
     auto sb = b.Append(f->Block());
@@ -2041,7 +2042,7 @@
     auto* switch_ = b.Switch(true);
 
     auto* def = b.DefaultCase(switch_);
-    def->Append(mod.instructions.Create<ExitSwitch>(nullptr));
+    def->Append(mod.allocators.instructions.Create<ExitSwitch>(nullptr));
 
     auto* f = b.Function("my_func", ty.void_());
     auto sb = b.Append(f->Block());
@@ -2421,7 +2422,7 @@
 TEST_F(IR_ValidatorTest, ExitLoop_NullLoop) {
     auto* loop = b.Loop();
     loop->Continuing()->Append(b.NextIteration(loop));
-    loop->Body()->Append(mod.instructions.Create<ExitLoop>(nullptr));
+    loop->Body()->Append(mod.allocators.instructions.Create<ExitLoop>(nullptr));
 
     auto* f = b.Function("my_func", ty.void_());
     auto sb = b.Append(f->Block());
@@ -3128,7 +3129,8 @@
     auto* f = b.Function("my_func", ty.void_());
 
     b.Append(f->Block(), [&] {
-        b.Append(mod.instructions.Create<ir::Load>(b.InstructionResult(ty.i32()), nullptr));
+        b.Append(
+            mod.allocators.instructions.Create<ir::Load>(b.InstructionResult(ty.i32()), nullptr));
         b.Return(f);
     });
 
@@ -3157,7 +3159,8 @@
 
     b.Append(f->Block(), [&] {
         auto* let = b.Let("l", 1_i);
-        b.Append(mod.instructions.Create<ir::Load>(b.InstructionResult(ty.f32()), let->Result(0)));
+        b.Append(mod.allocators.instructions.Create<ir::Load>(b.InstructionResult(ty.f32()),
+                                                              let->Result(0)));
         b.Return(f);
     });
 
@@ -3188,7 +3191,8 @@
 
     b.Append(f->Block(), [&] {
         auto* var = b.Var(ty.ptr<function, i32>());
-        b.Append(mod.instructions.Create<ir::Load>(b.InstructionResult(ty.f32()), var->Result(0)));
+        b.Append(mod.allocators.instructions.Create<ir::Load>(b.InstructionResult(ty.f32()),
+                                                              var->Result(0)));
         b.Return(f);
     });
 
@@ -3218,7 +3222,7 @@
     auto* f = b.Function("my_func", ty.void_());
 
     b.Append(f->Block(), [&] {
-        b.Append(mod.instructions.Create<ir::Store>(nullptr, b.Constant(42_i)));
+        b.Append(mod.allocators.instructions.Create<ir::Store>(nullptr, b.Constant(42_i)));
         b.Return(f);
     });
 
@@ -3247,7 +3251,7 @@
 
     b.Append(f->Block(), [&] {
         auto* var = b.Var(ty.ptr<function, i32>());
-        b.Append(mod.instructions.Create<ir::Store>(var->Result(0), nullptr));
+        b.Append(mod.allocators.instructions.Create<ir::Store>(var->Result(0), nullptr));
         b.Return(f);
     });
 
@@ -3277,7 +3281,7 @@
 
     b.Append(f->Block(), [&] {
         auto* let = b.Let("l", 1_i);
-        b.Append(mod.instructions.Create<ir::Store>(let->Result(0), b.Constant(42_u)));
+        b.Append(mod.allocators.instructions.Create<ir::Store>(let->Result(0), b.Constant(42_u)));
         b.Return(f);
     });
 
@@ -3308,7 +3312,7 @@
 
     b.Append(f->Block(), [&] {
         auto* var = b.Var(ty.ptr<function, i32>());
-        b.Append(mod.instructions.Create<ir::Store>(var->Result(0), b.Constant(42_u)));
+        b.Append(mod.allocators.instructions.Create<ir::Store>(var->Result(0), b.Constant(42_u)));
         b.Return(f);
     });
 
@@ -3339,8 +3343,8 @@
 
     b.Append(f->Block(), [&] {
         auto* var = b.Var(ty.ptr<function, vec3<f32>>());
-        b.Append(mod.instructions.Create<ir::LoadVectorElement>(nullptr, var->Result(0),
-                                                                b.Constant(1_i)));
+        b.Append(mod.allocators.instructions.Create<ir::LoadVectorElement>(nullptr, var->Result(0),
+                                                                           b.Constant(1_i)));
         b.Return(f);
     });
 
@@ -3370,8 +3374,8 @@
     auto* f = b.Function("my_func", ty.void_());
 
     b.Append(f->Block(), [&] {
-        b.Append(mod.instructions.Create<ir::LoadVectorElement>(b.InstructionResult(ty.f32()),
-                                                                nullptr, b.Constant(1_i)));
+        b.Append(mod.allocators.instructions.Create<ir::LoadVectorElement>(
+            b.InstructionResult(ty.f32()), nullptr, b.Constant(1_i)));
         b.Return(f);
     });
 
@@ -3400,8 +3404,8 @@
 
     b.Append(f->Block(), [&] {
         auto* var = b.Var(ty.ptr<function, vec3<f32>>());
-        b.Append(mod.instructions.Create<ir::LoadVectorElement>(b.InstructionResult(ty.f32()),
-                                                                var->Result(0), nullptr));
+        b.Append(mod.allocators.instructions.Create<ir::LoadVectorElement>(
+            b.InstructionResult(ty.f32()), var->Result(0), nullptr));
         b.Return(f);
     });
 
@@ -3430,8 +3434,8 @@
     auto* f = b.Function("my_func", ty.void_());
 
     b.Append(f->Block(), [&] {
-        b.Append(mod.instructions.Create<ir::StoreVectorElement>(nullptr, b.Constant(1_i),
-                                                                 b.Constant(2_i)));
+        b.Append(mod.allocators.instructions.Create<ir::StoreVectorElement>(
+            nullptr, b.Constant(1_i), b.Constant(2_i)));
         b.Return(f);
     });
 
@@ -3460,8 +3464,8 @@
 
     b.Append(f->Block(), [&] {
         auto* var = b.Var(ty.ptr<function, vec3<f32>>());
-        b.Append(mod.instructions.Create<ir::StoreVectorElement>(var->Result(0), nullptr,
-                                                                 b.Constant(2_i)));
+        b.Append(mod.allocators.instructions.Create<ir::StoreVectorElement>(var->Result(0), nullptr,
+                                                                            b.Constant(2_i)));
         b.Return(f);
     });
 
@@ -3499,8 +3503,8 @@
 
     b.Append(f->Block(), [&] {
         auto* var = b.Var(ty.ptr<function, vec3<f32>>());
-        b.Append(mod.instructions.Create<ir::StoreVectorElement>(var->Result(0), b.Constant(1_i),
-                                                                 nullptr));
+        b.Append(mod.allocators.instructions.Create<ir::StoreVectorElement>(
+            var->Result(0), b.Constant(1_i), nullptr));
         b.Return(f);
     });
 
diff --git a/src/tint/lang/core/ir/var.cc b/src/tint/lang/core/ir/var.cc
index 3788822..60e8c33 100644
--- a/src/tint/lang/core/ir/var.cc
+++ b/src/tint/lang/core/ir/var.cc
@@ -53,7 +53,7 @@
 
 Var* Var::Clone(CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
-    auto* new_var = ctx.ir.instructions.Create<Var>(new_result);
+    auto* new_var = ctx.ir.allocators.instructions.Create<Var>(new_result);
 
     new_var->binding_point_ = binding_point_;
     new_var->attributes_ = attributes_;
diff --git a/src/tint/lang/msl/ir/builtin_call.cc b/src/tint/lang/msl/ir/builtin_call.cc
index 8e20da3..65fcd1e 100644
--- a/src/tint/lang/msl/ir/builtin_call.cc
+++ b/src/tint/lang/msl/ir/builtin_call.cc
@@ -50,7 +50,7 @@
 BuiltinCall* BuiltinCall::Clone(core::ir::CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto new_args = ctx.Clone<BuiltinCall::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<BuiltinCall>(new_result, func_, new_args);
+    return ctx.ir.allocators.instructions.Create<BuiltinCall>(new_result, func_, new_args);
 }
 
 }  // namespace tint::msl::ir
diff --git a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
index ab0da1d..9788113 100644
--- a/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/msl/writer/raise/builtin_polyfill.cc
@@ -57,10 +57,7 @@
     void Process() {
         // Find the builtins that need replacing.
         Vector<core::ir::CoreBuiltinCall*, 4> worklist;
-        for (auto* inst : ir.instructions.Objects()) {
-            if (!inst->Alive()) {
-                continue;
-            }
+        for (auto* inst : ir.Instructions()) {
             if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) {
                 switch (builtin->Func()) {
                     case core::BuiltinFn::kStorageBarrier:
diff --git a/src/tint/lang/spirv/ir/builtin_call.cc b/src/tint/lang/spirv/ir/builtin_call.cc
index 5503e55..dfdb86d 100644
--- a/src/tint/lang/spirv/ir/builtin_call.cc
+++ b/src/tint/lang/spirv/ir/builtin_call.cc
@@ -50,7 +50,7 @@
 BuiltinCall* BuiltinCall::Clone(core::ir::CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto new_args = ctx.Clone<BuiltinCall::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<BuiltinCall>(new_result, func_, new_args);
+    return ctx.ir.allocators.instructions.Create<BuiltinCall>(new_result, func_, new_args);
 }
 
 }  // namespace tint::spirv::ir
diff --git a/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
index 1dee037..daf59df 100644
--- a/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
+++ b/src/tint/lang/spirv/reader/lower/vector_element_pointer.cc
@@ -62,10 +62,7 @@
     void Process() {
         // Find the access instructions that need to be replaced.
         Vector<Access, 8> worklist;
-        for (auto* inst : ir.instructions.Objects()) {
-            if (!inst->Alive()) {
-                continue;
-            }
+        for (auto* inst : ir.Instructions()) {
             if (auto* access = inst->As<core::ir::Access>()) {
                 auto* source_ty = access->Object()->Type();
                 if (!source_ty->Is<core::type::Pointer>()) {
diff --git a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
index 335b3b7..9af1437 100644
--- a/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
+++ b/src/tint/lang/spirv/writer/raise/builtin_polyfill.cc
@@ -68,10 +68,7 @@
     void Process() {
         // Find the builtins that need replacing.
         Vector<core::ir::CoreBuiltinCall*, 4> worklist;
-        for (auto* inst : ir.instructions.Objects()) {
-            if (!inst->Alive()) {
-                continue;
-            }
+        for (auto* inst : ir.Instructions()) {
             if (auto* builtin = inst->As<core::ir::CoreBuiltinCall>()) {
                 switch (builtin->Func()) {
                     case core::BuiltinFn::kArrayLength:
@@ -190,7 +187,7 @@
     /// @param value the literal value
     /// @returns the literal operand
     spirv::ir::LiteralOperand* Literal(u32 value) {
-        return ir.values.Create<spirv::ir::LiteralOperand>(b.ConstantValue(value));
+        return ir.allocators.values.Create<spirv::ir::LiteralOperand>(b.ConstantValue(value));
     }
 
     /// Handle an `arrayLength()` builtin.
diff --git a/src/tint/lang/spirv/writer/raise/expand_implicit_splats.cc b/src/tint/lang/spirv/writer/raise/expand_implicit_splats.cc
index 3385ad0..0a6ba91 100644
--- a/src/tint/lang/spirv/writer/raise/expand_implicit_splats.cc
+++ b/src/tint/lang/spirv/writer/raise/expand_implicit_splats.cc
@@ -48,10 +48,7 @@
     // to be replaced in a second pass.
     Vector<core::ir::CoreBinary*, 4> binary_worklist;
     Vector<core::ir::CoreBuiltinCall*, 4> builtin_worklist;
-    for (auto* inst : ir.instructions.Objects()) {
-        if (!inst->Alive()) {
-            continue;
-        }
+    for (auto* inst : ir.Instructions()) {
         if (auto* construct = inst->As<core::ir::Construct>()) {
             // A vector constructor with a single scalar argument needs to be modified to replicate
             // the argument N times.
diff --git a/src/tint/lang/spirv/writer/raise/handle_matrix_arithmetic.cc b/src/tint/lang/spirv/writer/raise/handle_matrix_arithmetic.cc
index 3e77d21..7cfd7e0 100644
--- a/src/tint/lang/spirv/writer/raise/handle_matrix_arithmetic.cc
+++ b/src/tint/lang/spirv/writer/raise/handle_matrix_arithmetic.cc
@@ -50,10 +50,7 @@
     // Find the instructions that need to be modified.
     Vector<core::ir::CoreBinary*, 4> binary_worklist;
     Vector<core::ir::Convert*, 4> convert_worklist;
-    for (auto* inst : ir.instructions.Objects()) {
-        if (!inst->Alive()) {
-            continue;
-        }
+    for (auto* inst : ir.Instructions()) {
         if (auto* binary = inst->As<core::ir::CoreBinary>()) {
             TINT_ASSERT(binary->Operands().Length() == 2);
             if (binary->LHS()->Type()->Is<core::type::Matrix>() ||
diff --git a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
index cd41b6e..901b4bb 100644
--- a/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
+++ b/src/tint/lang/spirv/writer/raise/var_for_dynamic_index.cc
@@ -125,8 +125,8 @@
 
     // Find the access instructions that need replacing.
     Vector<AccessToReplace, 4> worklist;
-    for (auto* inst : ir.instructions.Objects()) {
-        if (auto* access = inst->As<core::ir::Access>(); access && access->Alive()) {
+    for (auto* inst : ir.Instructions()) {
+        if (auto* access = inst->As<core::ir::Access>()) {
             if (auto to_replace = ShouldReplace(access)) {
                 worklist.Push(to_replace.value());
             }
diff --git a/src/tint/lang/wgsl/ir/builtin_call.cc b/src/tint/lang/wgsl/ir/builtin_call.cc
index b841a6d..15696df 100644
--- a/src/tint/lang/wgsl/ir/builtin_call.cc
+++ b/src/tint/lang/wgsl/ir/builtin_call.cc
@@ -50,7 +50,7 @@
 BuiltinCall* BuiltinCall::Clone(core::ir::CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto new_args = ctx.Clone<BuiltinCall::kDefaultNumOperands>(Args());
-    return ctx.ir.instructions.Create<BuiltinCall>(new_result, fn_, new_args);
+    return ctx.ir.allocators.instructions.Create<BuiltinCall>(new_result, fn_, new_args);
 }
 
 }  // namespace tint::wgsl::ir
diff --git a/src/tint/lang/wgsl/ir/unary.cc b/src/tint/lang/wgsl/ir/unary.cc
index f8c90f3..7d98bef 100644
--- a/src/tint/lang/wgsl/ir/unary.cc
+++ b/src/tint/lang/wgsl/ir/unary.cc
@@ -45,7 +45,7 @@
 Unary* Unary::Clone(core::ir::CloneContext& ctx) {
     auto* new_result = ctx.Clone(Result(0));
     auto* val = ctx.Remap(Val());
-    return ctx.ir.instructions.Create<Unary>(new_result, Op(), val);
+    return ctx.ir.allocators.instructions.Create<Unary>(new_result, Op(), val);
 }
 
 const core::intrinsic::TableData& Unary::TableData() const {
diff --git a/src/tint/lang/wgsl/reader/lower/lower.cc b/src/tint/lang/wgsl/reader/lower/lower.cc
index f9a6ed9..21782ee 100644
--- a/src/tint/lang/wgsl/reader/lower/lower.cc
+++ b/src/tint/lang/wgsl/reader/lower/lower.cc
@@ -181,10 +181,7 @@
 
     core::ir::Builder b{mod};
     core::type::Manager& ty{mod.Types()};
-    for (auto* inst : mod.instructions.Objects()) {
-        if (!inst->Alive()) {
-            continue;
-        }
+    for (auto* inst : mod.Instructions()) {
         if (auto* call = inst->As<wgsl::ir::BuiltinCall>()) {
             switch (call->Func()) {
                 case BuiltinFn::kWorkgroupUniformLoad: {
@@ -204,8 +201,9 @@
                 }
                 default: {
                     Vector<core::ir::Value*, 8> args(call->Args());
-                    auto* replacement = mod.instructions.Create<core::ir::CoreBuiltinCall>(
-                        call->Result(0), Convert(call->Func()), std::move(args));
+                    auto* replacement =
+                        mod.allocators.instructions.Create<core::ir::CoreBuiltinCall>(
+                            call->Result(0), Convert(call->Func()), std::move(args));
                     call->ReplaceWith(replacement);
                     call->ClearResults();
                     break;
diff --git a/src/tint/lang/wgsl/reader/lower/lower_test.cc b/src/tint/lang/wgsl/reader/lower/lower_test.cc
index d8ec462..ba73688 100644
--- a/src/tint/lang/wgsl/reader/lower/lower_test.cc
+++ b/src/tint/lang/wgsl/reader/lower/lower_test.cc
@@ -45,11 +45,12 @@
     auto* f = b.Function("f", ty.void_());
     b.Append(f->Block(), [&] {  //
         auto* result = b.InstructionResult(ty.i32());
-        b.Append(b.ir.instructions.Create<wgsl::ir::BuiltinCall>(result, wgsl::BuiltinFn::kMax,
-                                                                 Vector{
-                                                                     b.Value(i32(1)),
-                                                                     b.Value(i32(2)),
-                                                                 }));
+        b.Append(b.ir.allocators.instructions.Create<wgsl::ir::BuiltinCall>(result,
+                                                                            wgsl::BuiltinFn::kMax,
+                                                                            Vector{
+                                                                                b.Value(i32(1)),
+                                                                                b.Value(i32(2)),
+                                                                            }));
         b.Return(f);
     });
 
@@ -84,7 +85,7 @@
     auto* f = b.Function("f", ty.i32());
     b.Append(f->Block(), [&] {  //
         auto* result = b.InstructionResult(ty.i32());
-        b.Append(b.ir.instructions.Create<wgsl::ir::BuiltinCall>(
+        b.Append(b.ir.allocators.instructions.Create<wgsl::ir::BuiltinCall>(
             result, wgsl::BuiltinFn::kWorkgroupUniformLoad, Vector{wgvar->Result(0)}));
         b.Return(f, result);
     });
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
index b838c02..4feffb8 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir.cc
@@ -1045,8 +1045,9 @@
                         inst = impl.builder_.Bitcast(ty, args[0]);
                     } else {
                         auto* res = impl.builder_.InstructionResult(ty);
-                        inst = impl.builder_.ir.instructions.Create<wgsl::ir::BuiltinCall>(
-                            res, b->Fn(), std::move(args));
+                        inst =
+                            impl.builder_.ir.allocators.instructions.Create<wgsl::ir::BuiltinCall>(
+                                res, b->Fn(), std::move(args));
                     }
                 } else if (sem->Target()->As<sem::ValueConstructor>()) {
                     inst = impl.builder_.Construct(ty, std::move(args));
diff --git a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
index 8ae7aa1..f8f2d0a 100644
--- a/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
+++ b/src/tint/lang/wgsl/reader/program_to_ir/program_to_ir_test.cc
@@ -51,7 +51,7 @@
 T* FindSingleInstruction(core::ir::Module& mod) {
     T* found = nullptr;
     size_t count = 0;
-    for (auto* node : mod.instructions.Objects()) {
+    for (auto* node : mod.Instructions()) {
         if (auto* as = node->As<T>()) {
             count++;
             if (!found) {
diff --git a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
index 6e78226..10e88ca 100644
--- a/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
+++ b/src/tint/lang/wgsl/writer/ir_to_program/ir_to_program_test.cc
@@ -2545,7 +2545,7 @@
 TEST_F(IRToProgramTest, Enable_ChromiumExperimentalSubgroups_SubgroupBallot) {
     auto* fn = b.Function("f", ty.void_());
     b.Append(fn->Block(), [&] {
-        auto* call = b.Append(mod.instructions.Create<wgsl::ir::BuiltinCall>(
+        auto* call = b.Append(mod.allocators.instructions.Create<wgsl::ir::BuiltinCall>(
             b.InstructionResult(ty.vec4<u32>()), wgsl::BuiltinFn::kSubgroupBallot, Empty));
         b.Let("v", call);
         b.Return(fn);
@@ -2564,7 +2564,7 @@
     auto* fn = b.Function("f", ty.void_());
     b.Append(fn->Block(), [&] {
         auto* one = b.Value(1_u);
-        auto* call = b.Append(mod.instructions.Create<wgsl::ir::BuiltinCall>(
+        auto* call = b.Append(mod.allocators.instructions.Create<wgsl::ir::BuiltinCall>(
             b.InstructionResult(ty.u32()), wgsl::BuiltinFn::kSubgroupBroadcast, Vector{one, one}));
         b.Let("v", call);
         b.Return(fn);
diff --git a/src/tint/lang/wgsl/writer/raise/ptr_to_ref.cc b/src/tint/lang/wgsl/writer/raise/ptr_to_ref.cc
index 916a65d..d6765c6 100644
--- a/src/tint/lang/wgsl/writer/raise/ptr_to_ref.cc
+++ b/src/tint/lang/wgsl/writer/raise/ptr_to_ref.cc
@@ -100,7 +100,8 @@
         TINT_ASSERT(operand);
         if (auto* ref_ty = As<core::type::Reference>(operand->Type())) {
             auto* as_ptr = b.InstructionResult(RefToPtr(ref_ty));
-            mod.instructions.Create<wgsl::ir::Unary>(as_ptr, core::UnaryOp::kAddressOf, operand)
+            mod.allocators.instructions
+                .Create<wgsl::ir::Unary>(as_ptr, core::UnaryOp::kAddressOf, operand)
                 ->InsertBefore(use.instruction);
             use.instruction->SetOperand(use.operand_index, as_ptr);
         }
@@ -115,7 +116,8 @@
         auto* operand = use.instruction->Operand(use.operand_index);
         if (auto* ptr_ty = As<core::type::Pointer>(operand->Type())) {
             auto* as_ptr = b.InstructionResult(PtrToRef(ptr_ty));
-            mod.instructions.Create<wgsl::ir::Unary>(as_ptr, core::UnaryOp::kIndirection, operand)
+            mod.allocators.instructions
+                .Create<wgsl::ir::Unary>(as_ptr, core::UnaryOp::kIndirection, operand)
                 ->InsertBefore(use.instruction);
             use.instruction->SetOperand(use.operand_index, as_ptr);
         }
diff --git a/src/tint/lang/wgsl/writer/raise/raise.cc b/src/tint/lang/wgsl/writer/raise/raise.cc
index 2678f4c..46cfe7e 100644
--- a/src/tint/lang/wgsl/writer/raise/raise.cc
+++ b/src/tint/lang/wgsl/writer/raise/raise.cc
@@ -171,7 +171,7 @@
 
 void ReplaceBuiltinFnCall(core::ir::Module& mod, core::ir::CoreBuiltinCall* call) {
     Vector<core::ir::Value*, 8> args(call->Args());
-    auto* replacement = mod.instructions.Create<wgsl::ir::BuiltinCall>(
+    auto* replacement = mod.allocators.instructions.Create<wgsl::ir::BuiltinCall>(
         call->Result(0), Convert(call->Func()), std::move(args));
     call->ReplaceWith(replacement);
     call->ClearResults();
@@ -206,7 +206,7 @@
     call->Destroy();
 
     // Replace load with workgroupUniformLoad
-    auto* replacement = mod.instructions.Create<wgsl::ir::BuiltinCall>(
+    auto* replacement = mod.allocators.instructions.Create<wgsl::ir::BuiltinCall>(
         load->Result(0), wgsl::BuiltinFn::kWorkgroupUniformLoad, Vector{load->From()});
     load->ReplaceWith(replacement);
     load->ClearResults();
@@ -216,10 +216,7 @@
 }  // namespace
 
 Result<SuccessType> Raise(core::ir::Module& mod) {
-    for (auto* inst : mod.instructions.Objects()) {
-        if (!inst->Alive()) {
-            continue;
-        }
+    for (auto* inst : mod.Instructions()) {
         if (auto* call = inst->As<core::ir::CoreBuiltinCall>()) {
             switch (call->Func()) {
                 case core::BuiltinFn::kWorkgroupBarrier: