Import Tint changes from Dawn
Changes:
- 4c7b09fbb942a291f60c73b2e519be203a53b4ca [tint][ir] Use the FilteredIterator to avoid dead instruc... by Ben Clayton <bclayton@google.com>
GitOrigin-RevId: 4c7b09fbb942a291f60c73b2e519be203a53b4ca
Change-Id: Iaff6e8ff345944099e7bd8f0f200f752265ccdfc
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/182700
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@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 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:
