[tint][ir] Use the FilteredIterator to avoid dead instructions
The common pattern by many IR transforms is to iterate over the ir::Module::instructions.Objects(), which is the list of allocated instructions.
These include instructions that have been destroyed, and must not be used.
We've hit the same bug time and time again, where we've forgotten to check for `Instruction::Alive()` before using the instruction.
To avoid this mistake, add new ir::Module::Instructions() and ir::Module::Values() method that returns an iterator filtered to the alive instructions / values.
Move the allocators into a new 'allocators' structure member to discourage use of these for object iteration.
Change-Id: Ic0043c8a4b82b9ec10b2b8a3479f33c6e5d286ff
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/182122
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: James Price <jrprice@google.com>
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 31289b5..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:
