[ir] Add ability to clone in the IR.
Add support for cloning in the IR.
This CL is based off of https://dawn-review.googlesource.com/c/dawn/+/152400
Bug: tint:1718
Change-Id: I8fce30872ad2c02c59363d2cd8b5771917d97342
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/152221
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/lang/core/ir/BUILD.bazel b/src/tint/lang/core/ir/BUILD.bazel
index 7d7188c..195dbfd 100644
--- a/src/tint/lang/core/ir/BUILD.bazel
+++ b/src/tint/lang/core/ir/BUILD.bazel
@@ -35,6 +35,7 @@
"builder.cc",
"builtin_call.cc",
"call.cc",
+ "clone_context.cc",
"constant.cc",
"construct.cc",
"continue.cc",
@@ -85,6 +86,7 @@
"builder.h",
"builtin_call.h",
"call.h",
+ "clone_context.h",
"constant.h",
"construct.h",
"continue.h",
@@ -188,7 +190,9 @@
"store_vector_element_test.cc",
"switch_test.cc",
"swizzle_test.cc",
+ "terminate_invocation_test.cc",
"unary_test.cc",
+ "unreachable_test.cc",
"user_call_test.cc",
"validator_test.cc",
"value_test.cc",
diff --git a/src/tint/lang/core/ir/BUILD.cmake b/src/tint/lang/core/ir/BUILD.cmake
index 3e21f9e..8af569e 100644
--- a/src/tint/lang/core/ir/BUILD.cmake
+++ b/src/tint/lang/core/ir/BUILD.cmake
@@ -46,6 +46,8 @@
lang/core/ir/builtin_call.h
lang/core/ir/call.cc
lang/core/ir/call.h
+ lang/core/ir/clone_context.cc
+ lang/core/ir/clone_context.h
lang/core/ir/constant.cc
lang/core/ir/constant.h
lang/core/ir/construct.cc
@@ -188,7 +190,9 @@
lang/core/ir/store_vector_element_test.cc
lang/core/ir/switch_test.cc
lang/core/ir/swizzle_test.cc
+ lang/core/ir/terminate_invocation_test.cc
lang/core/ir/unary_test.cc
+ lang/core/ir/unreachable_test.cc
lang/core/ir/user_call_test.cc
lang/core/ir/validator_test.cc
lang/core/ir/value_test.cc
diff --git a/src/tint/lang/core/ir/BUILD.gn b/src/tint/lang/core/ir/BUILD.gn
index 4a844c9..661272f 100644
--- a/src/tint/lang/core/ir/BUILD.gn
+++ b/src/tint/lang/core/ir/BUILD.gn
@@ -49,6 +49,8 @@
"builtin_call.h",
"call.cc",
"call.h",
+ "clone_context.cc",
+ "clone_context.h",
"constant.cc",
"constant.h",
"construct.cc",
@@ -189,7 +191,9 @@
"store_vector_element_test.cc",
"switch_test.cc",
"swizzle_test.cc",
+ "terminate_invocation_test.cc",
"unary_test.cc",
+ "unreachable_test.cc",
"user_call_test.cc",
"validator_test.cc",
"value_test.cc",
diff --git a/src/tint/lang/core/ir/access.cc b/src/tint/lang/core/ir/access.cc
index 34051ac..73682b9 100644
--- a/src/tint/lang/core/ir/access.cc
+++ b/src/tint/lang/core/ir/access.cc
@@ -16,6 +16,9 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Access);
namespace tint::core::ir {
@@ -28,6 +31,13 @@
}
Access::~Access() = default;
+
+Access* Access::Clone(CloneContext& ctx) {
+ auto new_result = ctx.Clone(Result());
+ auto new_obj = ctx.Clone(Object());
+ auto new_indices = ctx.Clone<Access::kDefaultNumOperands>(Indices());
+ return ctx.ir.instructions.Create<Access>(new_result, new_obj, new_indices);
+}
//! @endcond
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/access.h b/src/tint/lang/core/ir/access.h
index 4c76121..0221f50 100644
--- a/src/tint/lang/core/ir/access.h
+++ b/src/tint/lang/core/ir/access.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// An access instruction in the IR.
-class Access : public Castable<Access, OperandInstruction<3, 1>> {
+class Access final : public Castable<Access, OperandInstruction<3, 1>> {
public:
/// The offset in Operands() for the object being accessed
static constexpr size_t kObjectOperandOffset = 0;
@@ -38,6 +38,9 @@
Access(InstructionResult* result, Value* object, VectorRef<Value*> indices);
~Access() override;
+ /// @copydoc Instruction::Clone()
+ Access* Clone(CloneContext& ctx) override;
+
/// @returns the object used for the access
Value* Object() { return operands_[kObjectOperandOffset]; }
diff --git a/src/tint/lang/core/ir/access_test.cc b/src/tint/lang/core/ir/access_test.cc
index 8b3c310..7d9c4d6 100644
--- a/src/tint/lang/core/ir/access_test.cc
+++ b/src/tint/lang/core/ir/access_test.cc
@@ -18,7 +18,8 @@
#include "gtest/gtest-spi.h"
#include "src/tint/lang/core/ir/ir_helper_test.h"
-using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::fluent_types; // NOLINT
+using namespace tint::core::number_suffixes; // NOLINT
namespace tint::core::ir {
namespace {
@@ -60,5 +61,43 @@
"");
}
+TEST_F(IR_AccessTest, Clone) {
+ auto* type = ty.ptr<function, i32>();
+ auto* var = b.Var(type);
+ auto* idx1 = b.Constant(u32(1));
+ auto* idx2 = b.Constant(u32(2));
+ auto* a = b.Access(type, var, idx1, idx2);
+
+ auto* new_a = clone_ctx.Clone(a);
+
+ EXPECT_NE(a, new_a);
+
+ EXPECT_NE(a->Result(), new_a->Result());
+ EXPECT_EQ(type, new_a->Result()->Type());
+
+ EXPECT_NE(nullptr, new_a->Object());
+ EXPECT_NE(a->Object(), new_a->Object());
+
+ auto indices = new_a->Indices();
+ EXPECT_EQ(2u, indices.Length());
+
+ auto* val0 = indices[0]->As<Constant>()->Value();
+ EXPECT_EQ(1_u, val0->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto* val1 = indices[1]->As<Constant>()->Value();
+ EXPECT_EQ(2_u, val1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_AccessTest, CloneNoIndices) {
+ auto* type = ty.ptr<function, i32>();
+ auto* var = b.Var(type);
+ auto* a = b.Access(type, var);
+
+ auto* new_a = clone_ctx.Clone(a);
+
+ auto indices = new_a->Indices();
+ EXPECT_EQ(0u, indices.Length());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/binary.cc b/src/tint/lang/core/ir/binary.cc
index 1b1bdcf..251d587 100644
--- a/src/tint/lang/core/ir/binary.cc
+++ b/src/tint/lang/core/ir/binary.cc
@@ -14,6 +14,9 @@
#include "src/tint/lang/core/ir/binary.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Binary);
namespace tint::core::ir {
@@ -26,6 +29,13 @@
Binary::~Binary() = default;
+Binary* Binary::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_lhs = ctx.Clone(LHS());
+ auto* new_rhs = ctx.Clone(RHS());
+ return ctx.ir.instructions.Create<Binary>(new_result, kind_, new_lhs, new_rhs);
+}
+
std::string_view ToString(enum Binary::Kind kind) {
switch (kind) {
case Binary::Kind::kAdd:
diff --git a/src/tint/lang/core/ir/binary.h b/src/tint/lang/core/ir/binary.h
index 53ddf1a..3f4c79a 100644
--- a/src/tint/lang/core/ir/binary.h
+++ b/src/tint/lang/core/ir/binary.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A binary instruction in the IR.
-class Binary : public Castable<Binary, OperandInstruction<2, 1>> {
+class Binary final : public Castable<Binary, OperandInstruction<2, 1>> {
public:
/// The offset in Operands() for the LHS
static constexpr size_t kLhsOperandOffset = 0;
@@ -62,6 +62,9 @@
Binary(InstructionResult* result, enum Kind kind, Value* lhs, Value* rhs);
~Binary() override;
+ /// @copydoc Instruction::Clone()
+ Binary* Clone(CloneContext& ctx) override;
+
/// @returns the kind of the binary instruction
enum Kind Kind() { return kind_; }
diff --git a/src/tint/lang/core/ir/binary_test.cc b/src/tint/lang/core/ir/binary_test.cc
index 721bf45..bf8553f 100644
--- a/src/tint/lang/core/ir/binary_test.cc
+++ b/src/tint/lang/core/ir/binary_test.cc
@@ -374,5 +374,26 @@
EXPECT_THAT(rhs_b->Usages(), testing::UnorderedElementsAre(Usage{inst, 1u}));
}
+TEST_F(IR_BinaryTest, Clone) {
+ auto* lhs = b.Constant(2_i);
+ auto* rhs = b.Constant(4_i);
+ auto* inst = b.And(mod.Types().i32(), lhs, rhs);
+
+ auto* c = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, c);
+
+ EXPECT_EQ(mod.Types().i32(), c->Result()->Type());
+ EXPECT_EQ(Binary::Kind::kAnd, c->Kind());
+
+ auto new_lhs = c->LHS()->As<Constant>()->Value();
+ ASSERT_TRUE(new_lhs->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(2_i, new_lhs->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+
+ auto new_rhs = c->RHS()->As<Constant>()->Value();
+ ASSERT_TRUE(new_rhs->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(4_i, new_rhs->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/bitcast.cc b/src/tint/lang/core/ir/bitcast.cc
index cdaf695..f1fc737 100644
--- a/src/tint/lang/core/ir/bitcast.cc
+++ b/src/tint/lang/core/ir/bitcast.cc
@@ -14,6 +14,9 @@
#include "src/tint/lang/core/ir/bitcast.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Bitcast);
namespace tint::core::ir {
@@ -25,4 +28,10 @@
Bitcast::~Bitcast() = default;
+Bitcast* Bitcast::Clone(CloneContext& ctx) {
+ auto* new_res = ctx.Clone(Result());
+ auto* new_val = ctx.Clone(Val());
+ return ctx.ir.instructions.Create<Bitcast>(new_res, new_val);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/bitcast.h b/src/tint/lang/core/ir/bitcast.h
index d5f7425..3c35280 100644
--- a/src/tint/lang/core/ir/bitcast.h
+++ b/src/tint/lang/core/ir/bitcast.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A bitcast instruction in the IR.
-class Bitcast : public Castable<Bitcast, Call> {
+class Bitcast final : public Castable<Bitcast, Call> {
public:
/// The offset in Operands() for the value
static constexpr size_t kValueOperandOffset = 0;
@@ -34,6 +34,9 @@
Bitcast(InstructionResult* result, Value* val);
~Bitcast() override;
+ /// @copydoc Instruction::Clone()
+ Bitcast* Clone(CloneContext& ctx) override;
+
/// @returns the operand value
Value* Val() { return operands_[kValueOperandOffset]; }
diff --git a/src/tint/lang/core/ir/bitcast_test.cc b/src/tint/lang/core/ir/bitcast_test.cc
index 02b190d..b6d14ff 100644
--- a/src/tint/lang/core/ir/bitcast_test.cc
+++ b/src/tint/lang/core/ir/bitcast_test.cc
@@ -70,5 +70,19 @@
"");
}
+TEST_F(IR_BitcastTest, Clone) {
+ auto* inst = b.Bitcast(mod.Types().i32(), 4_i);
+
+ auto* n = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, n);
+
+ EXPECT_EQ(mod.Types().i32(), n->Result()->Type());
+
+ auto new_val = n->Val()->As<Constant>()->Value();
+ ASSERT_TRUE(new_val->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(4_i, new_val->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/block.cc b/src/tint/lang/core/ir/block.cc
index f49c9d0..00654b9 100644
--- a/src/tint/lang/core/ir/block.cc
+++ b/src/tint/lang/core/ir/block.cc
@@ -13,6 +13,10 @@
// limitations under the License.
#include "src/tint/lang/core/ir/block.h"
+
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/control_instruction.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Block);
@@ -23,6 +27,30 @@
Block::~Block() = default;
+Block* Block::Clone(CloneContext&) {
+ TINT_UNREACHABLE() << "blocks must be cloned with CloneInto";
+ return nullptr;
+}
+
+void Block::CloneInto(CloneContext& ctx, Block* out) {
+ // Note, the `parent_` is not cloned here. Doing so can end up in infinite loops as we try to
+ // clone a control instruction and the blocks inside of it. The `parent_` pointer should be set
+ // by the control instructions constructor.
+
+ for (auto* inst_in : *this) {
+ auto* inst_out = inst_in->Clone(ctx);
+ auto results_out = inst_out->Results();
+ auto results_in = inst_in->Results();
+ TINT_ASSERT(results_out.Length() == results_in.Length());
+
+ size_t len = results_out.Length();
+ for (size_t i = 0; i < len; ++i) {
+ ctx.Replace(results_in[i], results_out[i]);
+ }
+ out->Append(inst_out);
+ }
+}
+
Instruction* Block::Prepend(Instruction* inst) {
TINT_ASSERT_OR_RETURN_VALUE(inst, inst);
TINT_ASSERT_OR_RETURN_VALUE(inst->Block() == nullptr, inst);
diff --git a/src/tint/lang/core/ir/block.h b/src/tint/lang/core/ir/block.h
index 1badce2..a01e8a0 100644
--- a/src/tint/lang/core/ir/block.h
+++ b/src/tint/lang/core/ir/block.h
@@ -36,6 +36,15 @@
Block();
~Block() override;
+ /// @param ctx the CloneContext used to clone this block
+ /// @returns a clone of this block
+ virtual Block* Clone(CloneContext& ctx);
+
+ /// Clones the block contents into the given block
+ /// @param ctx the CloneContext used to clone
+ /// @param out the block to clone into
+ virtual void CloneInto(CloneContext& ctx, Block* out);
+
/// @returns true if this is block has a terminator instruction
bool HasTerminator() {
return instructions_.last != nullptr && instructions_.last->Is<ir::Terminator>();
diff --git a/src/tint/lang/core/ir/block_param.cc b/src/tint/lang/core/ir/block_param.cc
index b2a992e..1993417 100644
--- a/src/tint/lang/core/ir/block_param.cc
+++ b/src/tint/lang/core/ir/block_param.cc
@@ -13,6 +13,9 @@
// limitations under the License.
#include "src/tint/lang/core/ir/block_param.h"
+
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::BlockParam);
@@ -25,4 +28,14 @@
BlockParam::~BlockParam() = default;
+BlockParam* BlockParam::Clone(CloneContext& ctx) {
+ auto* new_bp = ctx.ir.values.Create<BlockParam>(type_);
+
+ auto name = ctx.ir.NameOf(this);
+ if (name.IsValid()) {
+ ctx.ir.SetName(new_bp, ctx.ir.NameOf(this).Name());
+ }
+ return new_bp;
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/block_param.h b/src/tint/lang/core/ir/block_param.h
index 315aba7..778a6f2 100644
--- a/src/tint/lang/core/ir/block_param.h
+++ b/src/tint/lang/core/ir/block_param.h
@@ -31,6 +31,9 @@
/// @returns the type of the var
const core::type::Type* Type() override { return type_; }
+ /// @copydoc Instruction::Clone()
+ BlockParam* Clone(CloneContext& ctx) override;
+
private:
/// the result type of the instruction
const core::type::Type* type_ = nullptr;
diff --git a/src/tint/lang/core/ir/block_param_test.cc b/src/tint/lang/core/ir/block_param_test.cc
index 8677c83..ca29c98 100644
--- a/src/tint/lang/core/ir/block_param_test.cc
+++ b/src/tint/lang/core/ir/block_param_test.cc
@@ -12,8 +12,10 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/lang/core/ir/block_param.h"
+#include <string>
+
#include "gtest/gtest-spi.h"
+#include "src/tint/lang/core/ir/block_param.h"
#include "src/tint/lang/core/ir/ir_helper_test.h"
namespace tint::core::ir {
@@ -32,5 +34,23 @@
"");
}
+TEST_F(IR_BlockParamTest, Clone) {
+ auto* inst = b.BlockParam(mod.Types().i32());
+
+ auto* new_inst = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, new_inst);
+ EXPECT_EQ(mod.Types().i32(), new_inst->Type());
+}
+
+TEST_F(IR_BlockParamTest, CloneWithName) {
+ auto* inst = b.BlockParam("p", mod.Types().i32());
+
+ auto* new_inst = clone_ctx.Clone(inst);
+ EXPECT_EQ(mod.Types().i32(), new_inst->Type());
+
+ EXPECT_EQ(std::string("p"), mod.NameOf(new_inst).Name());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/break_if.cc b/src/tint/lang/core/ir/break_if.cc
index 56ce784..bf4a6c3 100644
--- a/src/tint/lang/core/ir/break_if.cc
+++ b/src/tint/lang/core/ir/break_if.cc
@@ -17,7 +17,9 @@
#include <utility>
#include "src/tint/lang/core/ir/block.h"
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/loop.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/utils/ice/ice.h"
@@ -38,4 +40,11 @@
BreakIf::~BreakIf() = default;
+BreakIf* BreakIf::Clone(CloneContext& ctx) {
+ auto* new_loop = ctx.Clone(loop_);
+ auto* new_cond = ctx.Clone(Condition());
+ auto new_args = ctx.Clone<BreakIf::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<BreakIf>(new_cond, new_loop, new_args);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/break_if.h b/src/tint/lang/core/ir/break_if.h
index 1f73c2f..e672b90 100644
--- a/src/tint/lang/core/ir/break_if.h
+++ b/src/tint/lang/core/ir/break_if.h
@@ -29,7 +29,7 @@
namespace tint::core::ir {
/// A break-if iteration instruction.
-class BreakIf : public Castable<BreakIf, Terminator> {
+class BreakIf final : public Castable<BreakIf, Terminator> {
public:
/// The offset in Operands() for the condition
static constexpr size_t kConditionOperandOffset = 0;
@@ -44,6 +44,9 @@
BreakIf(Value* condition, ir::Loop* loop, VectorRef<Value*> args = tint::Empty);
~BreakIf() override;
+ /// @copydoc Instruction::Clone()
+ BreakIf* Clone(CloneContext& ctx) override;
+
/// @returns the MultiInBlock arguments
tint::Slice<Value* const> Args() override {
return operands_.Slice().Offset(kArgsOperandOffset);
diff --git a/src/tint/lang/core/ir/break_if_test.cc b/src/tint/lang/core/ir/break_if_test.cc
index 644e2dd7..18e5d2d 100644
--- a/src/tint/lang/core/ir/break_if_test.cc
+++ b/src/tint/lang/core/ir/break_if_test.cc
@@ -58,5 +58,49 @@
"");
}
+TEST_F(IR_BreakIfTest, Clone) {
+ auto* loop = b.Loop();
+ auto* cond = b.Constant(true);
+ auto* arg1 = b.Constant(1_u);
+ auto* arg2 = b.Constant(2_u);
+
+ auto* brk = b.BreakIf(loop, cond, arg1, arg2);
+
+ auto* new_loop = clone_ctx.Clone(loop);
+ clone_ctx.Replace(loop, new_loop);
+
+ auto* new_brk = clone_ctx.Clone(brk);
+
+ EXPECT_NE(brk, new_brk);
+
+ EXPECT_EQ(new_loop, new_brk->Loop());
+
+ auto args = new_brk->Args();
+ EXPECT_EQ(2u, args.Length());
+
+ auto new_cond = new_brk->Condition()->As<Constant>()->Value();
+ ASSERT_TRUE(new_cond->Is<core::constant::Scalar<bool>>());
+ EXPECT_TRUE(new_cond->As<core::constant::Scalar<bool>>()->ValueAs<bool>());
+
+ auto new_arg0 = args[0]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg0->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(1_u, new_arg0->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto new_arg1 = args[1]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg1->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(2_u, new_arg1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_BreakIfTest, CloneNoArgs) {
+ auto* loop = b.Loop();
+ auto* cond = b.Constant(true);
+
+ auto* brk = b.BreakIf(loop, cond);
+ auto* new_brk = clone_ctx.Clone(brk);
+
+ auto args = new_brk->Args();
+ EXPECT_EQ(0u, args.Length());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/builder.h b/src/tint/lang/core/ir/builder.h
index 90608d2..7088ee0 100644
--- a/src/tint/lang/core/ir/builder.h
+++ b/src/tint/lang/core/ir/builder.h
@@ -793,7 +793,6 @@
return nullptr;
}
auto* let = Append(ir.instructions.Create<ir::Let>(InstructionResult(val->Type()), val));
- ir.SetName(let, name);
ir.SetName(let->Result(), name);
return let;
}
diff --git a/src/tint/lang/core/ir/builtin_call.h b/src/tint/lang/core/ir/builtin_call.h
index f152f7d..fe35766 100644
--- a/src/tint/lang/core/ir/builtin_call.h
+++ b/src/tint/lang/core/ir/builtin_call.h
@@ -21,7 +21,7 @@
namespace tint::core::ir {
-/// A builtin call instruction in the IR.
+/// The base class for builtin call instructions in the IR.
class BuiltinCall : public Castable<BuiltinCall, Call> {
public:
/// The base offset in Operands() for the args
diff --git a/src/tint/lang/core/ir/clone_context.cc b/src/tint/lang/core/ir/clone_context.cc
new file mode 100644
index 0000000..d5faa58
--- /dev/null
+++ b/src/tint/lang/core/ir/clone_context.cc
@@ -0,0 +1,24 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/lang/core/ir/clone_context.h"
+
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/let.h"
+
+namespace tint::core::ir {
+
+CloneContext::CloneContext(Module& module) : ir(module) {}
+
+} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/clone_context.h b/src/tint/lang/core/ir/clone_context.h
new file mode 100644
index 0000000..b5c375b
--- /dev/null
+++ b/src/tint/lang/core/ir/clone_context.h
@@ -0,0 +1,111 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#ifndef SRC_TINT_LANG_CORE_IR_CLONE_CONTEXT_H_
+#define SRC_TINT_LANG_CORE_IR_CLONE_CONTEXT_H_
+
+#include "src/tint/utils/containers/hashmap.h"
+#include "src/tint/utils/containers/transform.h"
+#include "src/tint/utils/traits/traits.h"
+
+namespace tint::core::ir {
+class Block;
+class Instruction;
+class Module;
+class Value;
+} // namespace tint::core::ir
+
+namespace tint::core::ir {
+
+/// Constant in the IR.
+class CloneContext {
+ public:
+ /// @param module the IR module
+ explicit CloneContext(Module& module);
+
+ /// The IR module
+ Module& ir;
+
+ /// Performs a clone of @p what.
+ /// @param what the item to clone
+ /// @return the cloned item
+ template <typename T>
+ T* Clone(T* what) {
+ if (auto replacement = replacements_.Get(what)) {
+ auto* cast = As<T>((*replacement)());
+ TINT_ASSERT(cast);
+ return cast;
+ }
+ auto* result = what->Clone(*this)->template As<T>();
+ Replace(what, result);
+ return result;
+ }
+
+ /// Performs a clone of all the elements in @p what.
+ /// @param what the elements to clone
+ /// @return the cloned elements
+ template <size_t N, typename T>
+ Vector<T*, N> Clone(Slice<T* const> what) {
+ return Transform<N>(what, [&](T* const p) { return Clone(p); });
+ }
+
+ /// Performs a clone of all the elements in @p what.
+ /// @param what the elements to clone
+ /// @return the cloned elements
+ template <size_t N, typename T>
+ Vector<T*, N> Clone(Slice<T*> what) {
+ return Transform<N>(what, [&](T* p) { return Clone(p); });
+ }
+
+ /// Performs a clone of all the elements in @p what.
+ /// @param what the elements to clone
+ /// @return the cloned elements
+ template <size_t N, typename T>
+ Vector<T*, N> Clone(Vector<T*, N> what) {
+ return Transform(what, [&](T* p) { return Clone(p); });
+ }
+
+ /// Registers the replacement of `what` with `with`
+ /// @param what the value or instruction to replace
+ /// @param with either a pointer to a replacement instruction, or a function with the signature
+ /// `T*(T*)` used to build the replacement
+ template <typename WHAT, typename WITH>
+ void Replace(WHAT* what, WITH&& with) {
+ using T = std::decay_t<WHAT>;
+ using F = std::decay_t<WITH>;
+
+ constexpr bool T_is_value = traits::IsTypeOrDerived<T, Value>;
+ constexpr bool T_is_instruction = traits::IsTypeOrDerived<T, Instruction>;
+ static_assert(T_is_value || T_is_instruction);
+
+ constexpr bool F_is_pointer = std::is_pointer_v<F>;
+ constexpr bool F_is_function = std::is_function_v<F>;
+ static_assert(F_is_pointer || F_is_function);
+
+ if constexpr (F_is_pointer) {
+ replacements_.Add(what, [with]() { return with; });
+ } else if constexpr (F_is_function) {
+ static_assert(std::is_same_v<traits::ParameterType<F, 0>, T*>);
+ static_assert(std::is_same_v<traits::ReturnType<F>, T*>);
+ replacements_.Add(what, [what, with]() { return with(what); });
+ }
+ }
+
+ private:
+ Hashmap<CastableBase*, std::function<CastableBase*()>, 8> replacements_;
+};
+
+} // namespace tint::core::ir
+
+#endif // SRC_TINT_LANG_CORE_IR_CLONE_CONTEXT_H_
diff --git a/src/tint/lang/core/ir/constant.cc b/src/tint/lang/core/ir/constant.cc
index 2510a39..a9cc837 100644
--- a/src/tint/lang/core/ir/constant.cc
+++ b/src/tint/lang/core/ir/constant.cc
@@ -25,4 +25,8 @@
Constant::~Constant() = default;
+Constant* Constant::Clone(CloneContext&) {
+ return this; // Constants are immutable so can just return ourselves.
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/constant.h b/src/tint/lang/core/ir/constant.h
index a316e24..ebf584b 100644
--- a/src/tint/lang/core/ir/constant.h
+++ b/src/tint/lang/core/ir/constant.h
@@ -34,6 +34,9 @@
/// @returns the type of the constant
const core::type::Type* Type() override { return value_->Type(); }
+ /// @copydoc Value::Clone()
+ Constant* Clone(CloneContext& ctx) override;
+
private:
const core::constant::Value* const value_ = nullptr;
};
diff --git a/src/tint/lang/core/ir/constant_test.cc b/src/tint/lang/core/ir/constant_test.cc
index 73b4282..1bdcfb9 100644
--- a/src/tint/lang/core/ir/constant_test.cc
+++ b/src/tint/lang/core/ir/constant_test.cc
@@ -113,5 +113,12 @@
"");
}
+TEST_F(IR_ConstantTest, Clone) {
+ auto* c = b.Constant(2_u);
+ auto* new_c = clone_ctx.Clone(c);
+
+ EXPECT_EQ(c, new_c);
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/construct.cc b/src/tint/lang/core/ir/construct.cc
index bbdb1db..f9ae628 100644
--- a/src/tint/lang/core/ir/construct.cc
+++ b/src/tint/lang/core/ir/construct.cc
@@ -16,6 +16,9 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Construct);
namespace tint::core::ir {
@@ -27,4 +30,10 @@
Construct::~Construct() = default;
+Construct* Construct::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto new_args = ctx.Clone<Construct::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<Construct>(new_result, new_args);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/construct.h b/src/tint/lang/core/ir/construct.h
index 61c9a76..29fe22c 100644
--- a/src/tint/lang/core/ir/construct.h
+++ b/src/tint/lang/core/ir/construct.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A constructor instruction in the IR.
-class Construct : public Castable<Construct, Call> {
+class Construct final : public Castable<Construct, Call> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
@@ -34,6 +34,9 @@
explicit Construct(InstructionResult* result, VectorRef<Value*> args = tint::Empty);
~Construct() override;
+ /// @copydoc Instruction::Clone()
+ Construct* Clone(CloneContext& ctx) override;
+
/// @returns the friendly name for the instruction
std::string FriendlyName() override { return "construct"; }
};
diff --git a/src/tint/lang/core/ir/construct_test.cc b/src/tint/lang/core/ir/construct_test.cc
index 044fb87..0340892 100644
--- a/src/tint/lang/core/ir/construct_test.cc
+++ b/src/tint/lang/core/ir/construct_test.cc
@@ -54,5 +54,35 @@
"");
}
+TEST_F(IR_ConstructTest, Clone) {
+ auto* arg1 = b.Constant(true);
+ auto* arg2 = b.Constant(false);
+ auto* c = b.Construct(mod.Types().f32(), arg1, arg2);
+
+ auto* new_c = clone_ctx.Clone(c);
+
+ EXPECT_NE(c, new_c);
+ EXPECT_NE(c->Result(), new_c->Result());
+ EXPECT_EQ(mod.Types().f32(), new_c->Result()->Type());
+
+ auto args = new_c->Args();
+ EXPECT_EQ(2u, args.Length());
+
+ auto* val0 = args[0]->As<Constant>()->Value();
+ EXPECT_TRUE(val0->As<core::constant::Scalar<bool>>()->ValueAs<bool>());
+
+ auto* val1 = args[1]->As<Constant>()->Value();
+ EXPECT_FALSE(val1->As<core::constant::Scalar<bool>>()->ValueAs<bool>());
+}
+
+TEST_F(IR_ConstructTest, CloneEmpty) {
+ auto* c = b.Construct(mod.Types().f32());
+
+ auto* new_c = clone_ctx.Clone(c);
+ EXPECT_NE(c->Result(), new_c->Result());
+ EXPECT_EQ(mod.Types().f32(), new_c->Result()->Type());
+ EXPECT_TRUE(new_c->Args().IsEmpty());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/continue.cc b/src/tint/lang/core/ir/continue.cc
index 6fb585b..7ddb1a3 100644
--- a/src/tint/lang/core/ir/continue.cc
+++ b/src/tint/lang/core/ir/continue.cc
@@ -17,7 +17,9 @@
#include <utility>
#include "src/tint/lang/core/ir/block.h"
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/loop.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/utils/ice/ice.h"
@@ -37,4 +39,11 @@
Continue::~Continue() = default;
+Continue* Continue::Clone(CloneContext& ctx) {
+ auto* new_loop = ctx.Clone(Loop());
+ auto new_args = ctx.Clone<Continue::kDefaultNumOperands>(Args());
+
+ return ctx.ir.instructions.Create<Continue>(new_loop, new_args);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/continue.h b/src/tint/lang/core/ir/continue.h
index 1c393e1..b4a0178 100644
--- a/src/tint/lang/core/ir/continue.h
+++ b/src/tint/lang/core/ir/continue.h
@@ -28,7 +28,7 @@
namespace tint::core::ir {
/// A continue instruction.
-class Continue : public Castable<Continue, Terminator> {
+class Continue final : public Castable<Continue, Terminator> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
@@ -39,6 +39,9 @@
explicit Continue(ir::Loop* loop, VectorRef<Value*> args = tint::Empty);
~Continue() override;
+ /// @copydoc Instruction::Clone()
+ Continue* Clone(CloneContext& ctx) override;
+
/// @returns the loop owning the continue block
ir::Loop* Loop() { return loop_; }
diff --git a/src/tint/lang/core/ir/continue_test.cc b/src/tint/lang/core/ir/continue_test.cc
index 0dfe024..88c25ec 100644
--- a/src/tint/lang/core/ir/continue_test.cc
+++ b/src/tint/lang/core/ir/continue_test.cc
@@ -56,5 +56,38 @@
"");
}
+TEST_F(IR_ContinueTest, Clone) {
+ auto* loop = b.Loop();
+ auto* cont = b.Continue(loop);
+
+ auto* new_loop = clone_ctx.Clone(loop);
+ clone_ctx.Replace(loop, new_loop);
+
+ auto* new_c = clone_ctx.Clone(cont);
+
+ EXPECT_NE(cont, new_c);
+ EXPECT_EQ(new_loop, new_c->Loop());
+ EXPECT_TRUE(new_c->Args().IsEmpty());
+}
+
+TEST_F(IR_ContinueTest, CloneWithArgs) {
+ auto* loop = b.Loop();
+ auto* arg1 = b.Constant(1_u);
+ auto* arg2 = b.Constant(2_u);
+
+ auto* cont = b.Continue(loop, arg1, arg2);
+
+ auto* new_c = clone_ctx.Clone(cont);
+
+ auto args = new_c->Args();
+ EXPECT_EQ(2u, args.Length());
+
+ auto* val0 = args[0]->As<Constant>()->Value();
+ EXPECT_EQ(1_u, val0->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto* val1 = args[1]->As<Constant>()->Value();
+ EXPECT_EQ(2_u, val1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/convert.cc b/src/tint/lang/core/ir/convert.cc
index 3f3a70b..c9a5696 100644
--- a/src/tint/lang/core/ir/convert.cc
+++ b/src/tint/lang/core/ir/convert.cc
@@ -16,6 +16,9 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Convert);
namespace tint::core::ir {
@@ -27,4 +30,10 @@
Convert::~Convert() = default;
+Convert* Convert::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_val = ctx.Clone(Args()[0]);
+ return ctx.ir.instructions.Create<Convert>(new_result, new_val);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/convert.h b/src/tint/lang/core/ir/convert.h
index 1c58228..6699c18 100644
--- a/src/tint/lang/core/ir/convert.h
+++ b/src/tint/lang/core/ir/convert.h
@@ -24,7 +24,7 @@
namespace tint::core::ir {
/// A value conversion instruction in the IR.
-class Convert : public Castable<Convert, Call> {
+class Convert final : public Castable<Convert, Call> {
public:
/// The offset in Operands() for the value
static constexpr size_t kValueOperandOffset = 0;
@@ -35,6 +35,9 @@
Convert(InstructionResult* result, Value* value);
~Convert() override;
+ /// @copydoc Instruction::Clone()
+ Convert* Clone(CloneContext& ctx) override;
+
/// @returns the friendly name for the instruction
std::string FriendlyName() override { return "convert"; }
};
diff --git a/src/tint/lang/core/ir/convert_test.cc b/src/tint/lang/core/ir/convert_test.cc
index d239dcc..3732b8f 100644
--- a/src/tint/lang/core/ir/convert_test.cc
+++ b/src/tint/lang/core/ir/convert_test.cc
@@ -41,5 +41,21 @@
EXPECT_EQ(c->Result()->Source(), c);
}
+TEST_F(IR_ConvertTest, Clone) {
+ auto* c = b.Convert(mod.Types().f32(), 1_u);
+
+ auto* new_c = clone_ctx.Clone(c);
+
+ EXPECT_NE(c, new_c);
+ EXPECT_NE(c->Result(), new_c->Result());
+ EXPECT_EQ(mod.Types().f32(), new_c->Result()->Type());
+
+ auto args = new_c->Args();
+ EXPECT_EQ(1u, args.Length());
+
+ auto* val0 = args[0]->As<Constant>()->Value();
+ EXPECT_EQ(1_u, val0->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/core_builtin_call.cc b/src/tint/lang/core/ir/core_builtin_call.cc
index 68a49c8..ada7835 100644
--- a/src/tint/lang/core/ir/core_builtin_call.cc
+++ b/src/tint/lang/core/ir/core_builtin_call.cc
@@ -16,6 +16,8 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::CoreBuiltinCall);
@@ -32,4 +34,10 @@
CoreBuiltinCall::~CoreBuiltinCall() = default;
+CoreBuiltinCall* CoreBuiltinCall::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto new_args = ctx.Clone<CoreBuiltinCall::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<CoreBuiltinCall>(new_result, func_, new_args);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/core_builtin_call.h b/src/tint/lang/core/ir/core_builtin_call.h
index bc1907b..c887b4b 100644
--- a/src/tint/lang/core/ir/core_builtin_call.h
+++ b/src/tint/lang/core/ir/core_builtin_call.h
@@ -26,7 +26,7 @@
namespace tint::core::ir {
/// A core builtin call instruction in the IR.
-class CoreBuiltinCall : public Castable<CoreBuiltinCall, BuiltinCall> {
+class CoreBuiltinCall final : public Castable<CoreBuiltinCall, BuiltinCall> {
public:
/// Constructor
/// @param result the result value
@@ -37,6 +37,9 @@
VectorRef<Value*> args = tint::Empty);
~CoreBuiltinCall() override;
+ /// @copydoc Instruction::Clone()
+ CoreBuiltinCall* Clone(CloneContext& ctx) override;
+
/// @returns the builtin function
core::Function Func() { return func_; }
diff --git a/src/tint/lang/core/ir/core_builtin_call_test.cc b/src/tint/lang/core/ir/core_builtin_call_test.cc
index 1a201e4..e05e564 100644
--- a/src/tint/lang/core/ir/core_builtin_call_test.cc
+++ b/src/tint/lang/core/ir/core_builtin_call_test.cc
@@ -73,5 +73,39 @@
"");
}
+TEST_F(IR_CoreBuiltinCallTest, Clone) {
+ auto* builtin = b.Call(mod.Types().f32(), core::Function::kAbs, 1_u, 2_u);
+
+ auto* new_b = clone_ctx.Clone(builtin);
+
+ EXPECT_NE(builtin, new_b);
+ EXPECT_NE(builtin->Result(), new_b->Result());
+ EXPECT_EQ(mod.Types().f32(), new_b->Result()->Type());
+
+ EXPECT_EQ(core::Function::kAbs, new_b->Func());
+
+ auto args = new_b->Args();
+ EXPECT_EQ(2u, args.Length());
+
+ auto* val0 = args[0]->As<Constant>()->Value();
+ EXPECT_EQ(1_u, val0->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto* val1 = args[1]->As<Constant>()->Value();
+ EXPECT_EQ(2_u, val1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_CoreBuiltinCallTest, CloneNoArgs) {
+ auto* builtin = b.Call(mod.Types().f32(), core::Function::kAbs);
+
+ auto* new_b = clone_ctx.Clone(builtin);
+ EXPECT_NE(builtin->Result(), new_b->Result());
+ EXPECT_EQ(mod.Types().f32(), new_b->Result()->Type());
+
+ EXPECT_EQ(core::Function::kAbs, new_b->Func());
+
+ auto args = new_b->Args();
+ EXPECT_TRUE(args.IsEmpty());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/discard.cc b/src/tint/lang/core/ir/discard.cc
index a4aaa19..7112f01 100644
--- a/src/tint/lang/core/ir/discard.cc
+++ b/src/tint/lang/core/ir/discard.cc
@@ -13,7 +13,9 @@
// limitations under the License.
#include "src/tint/lang/core/ir/discard.h"
-#include "src/tint/lang/core/type/void.h"
+
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Discard);
@@ -23,4 +25,8 @@
Discard::~Discard() = default;
+Discard* Discard::Clone(CloneContext& ctx) {
+ return ctx.ir.instructions.Create<Discard>();
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/discard.h b/src/tint/lang/core/ir/discard.h
index b98c753..b9d3ee9 100644
--- a/src/tint/lang/core/ir/discard.h
+++ b/src/tint/lang/core/ir/discard.h
@@ -23,12 +23,15 @@
namespace tint::core::ir {
/// A discard instruction in the IR.
-class Discard : public Castable<Discard, Call> {
+class Discard final : public Castable<Discard, Call> {
public:
/// Constructor
Discard();
~Discard() override;
+ /// @copydoc Instruction::Clone()
+ Discard* Clone(CloneContext& ctx) override;
+
/// @returns the friendly name for the instruction
std::string FriendlyName() override { return "discard"; }
};
diff --git a/src/tint/lang/core/ir/discard_test.cc b/src/tint/lang/core/ir/discard_test.cc
index 2ebb80b..cb3bf9be 100644
--- a/src/tint/lang/core/ir/discard_test.cc
+++ b/src/tint/lang/core/ir/discard_test.cc
@@ -34,5 +34,13 @@
EXPECT_FALSE(inst->HasMultiResults());
}
+TEST_F(IR_DiscardTest, Clone) {
+ auto* d = b.Discard();
+ auto* new_d = clone_ctx.Clone(d);
+
+ EXPECT_NE(d, new_d);
+ EXPECT_NE(nullptr, new_d);
+}
+
} // namespace
} // 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 641df8c..8bb6539 100644
--- a/src/tint/lang/core/ir/exit_if.cc
+++ b/src/tint/lang/core/ir/exit_if.cc
@@ -16,7 +16,9 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/if.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::ExitIf);
@@ -30,6 +32,12 @@
ExitIf::~ExitIf() = default;
+ExitIf* ExitIf::Clone(CloneContext& ctx) {
+ auto* new_if = ctx.Clone(If());
+ auto new_args = ctx.Clone<ExitIf::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<ExitIf>(new_if, new_args);
+}
+
void ExitIf::SetIf(ir::If* i) {
SetControlInstruction(i);
}
diff --git a/src/tint/lang/core/ir/exit_if.h b/src/tint/lang/core/ir/exit_if.h
index 31e6d55..57950ee 100644
--- a/src/tint/lang/core/ir/exit_if.h
+++ b/src/tint/lang/core/ir/exit_if.h
@@ -28,7 +28,7 @@
namespace tint::core::ir {
/// A exit if instruction.
-class ExitIf : public Castable<ExitIf, Exit> {
+class ExitIf final : public Castable<ExitIf, Exit> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
@@ -39,6 +39,9 @@
explicit ExitIf(ir::If* i, VectorRef<Value*> args = tint::Empty);
~ExitIf() override;
+ /// @copydoc Instruction::Clone()
+ ExitIf* Clone(CloneContext& ctx) override;
+
/// Re-associates the exit with the given if instruction
/// @param i the new If to exit from
void SetIf(ir::If* i);
diff --git a/src/tint/lang/core/ir/exit_if_test.cc b/src/tint/lang/core/ir/exit_if_test.cc
index a74ae52..46b70e1 100644
--- a/src/tint/lang/core/ir/exit_if_test.cc
+++ b/src/tint/lang/core/ir/exit_if_test.cc
@@ -53,5 +53,40 @@
EXPECT_FALSE(exit->Alive());
}
+TEST_F(IR_ExitIfTest, Clone) {
+ auto* arg1 = b.Constant(1_u);
+ auto* arg2 = b.Constant(2_u);
+ auto* if_ = b.If(true);
+ auto* e = b.ExitIf(if_, arg1, arg2);
+
+ auto* new_if = clone_ctx.Clone(if_);
+ auto* new_exit = clone_ctx.Clone(e);
+
+ EXPECT_NE(e, new_exit);
+ EXPECT_EQ(new_if, new_exit->If());
+
+ auto args = new_exit->Args();
+ ASSERT_EQ(2u, args.Length());
+
+ auto new_arg1 = args[0]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg1->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(1_u, new_arg1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto new_arg2 = args[1]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg2->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(2_u, new_arg2->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_ExitIfTest, CloneNoArgs) {
+ auto* if_ = b.If(true);
+ auto* e = b.ExitIf(if_);
+
+ auto* new_if = clone_ctx.Clone(if_);
+ auto* new_exit = clone_ctx.Clone(e);
+
+ EXPECT_EQ(new_if, new_exit->If());
+ EXPECT_TRUE(new_exit->Args().IsEmpty());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/exit_loop.cc b/src/tint/lang/core/ir/exit_loop.cc
index 674ce88..8e8d504 100644
--- a/src/tint/lang/core/ir/exit_loop.cc
+++ b/src/tint/lang/core/ir/exit_loop.cc
@@ -17,7 +17,9 @@
#include <utility>
#include "src/tint/lang/core/ir/block.h"
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/loop.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::ExitLoop);
@@ -31,6 +33,12 @@
ExitLoop::~ExitLoop() = default;
+ExitLoop* ExitLoop::Clone(CloneContext& ctx) {
+ auto* new_loop = ctx.Clone(Loop());
+ auto new_args = ctx.Clone<ExitLoop::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<ExitLoop>(new_loop, new_args);
+}
+
void ExitLoop::SetLoop(ir::Loop* l) {
SetControlInstruction(l);
}
diff --git a/src/tint/lang/core/ir/exit_loop.h b/src/tint/lang/core/ir/exit_loop.h
index 6e173df..e50c7f7 100644
--- a/src/tint/lang/core/ir/exit_loop.h
+++ b/src/tint/lang/core/ir/exit_loop.h
@@ -28,7 +28,7 @@
namespace tint::core::ir {
/// A exit loop instruction.
-class ExitLoop : public Castable<ExitLoop, Exit> {
+class ExitLoop final : public Castable<ExitLoop, Exit> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
@@ -39,6 +39,9 @@
explicit ExitLoop(ir::Loop* loop, VectorRef<Value*> args = tint::Empty);
~ExitLoop() override;
+ /// @copydoc Instruction::Clone()
+ ExitLoop* Clone(CloneContext& ctx) override;
+
/// Re-associates the exit with the given loop instruction
/// @param l the new loop to exit from
void SetLoop(ir::Loop* l);
diff --git a/src/tint/lang/core/ir/exit_loop_test.cc b/src/tint/lang/core/ir/exit_loop_test.cc
index 0853d6b1..79ce43a 100644
--- a/src/tint/lang/core/ir/exit_loop_test.cc
+++ b/src/tint/lang/core/ir/exit_loop_test.cc
@@ -43,5 +43,40 @@
EXPECT_FALSE(exit->Alive());
}
+TEST_F(IR_ExitLoopTest, Clone) {
+ auto* arg1 = b.Constant(1_u);
+ auto* arg2 = b.Constant(2_u);
+ auto* loop = b.Loop();
+ auto* e = b.ExitLoop(loop, arg1, arg2);
+
+ auto* new_loop = clone_ctx.Clone(loop);
+ auto* new_exit = clone_ctx.Clone(e);
+
+ EXPECT_NE(e, new_exit);
+ EXPECT_EQ(new_loop, new_exit->Loop());
+
+ auto args = new_exit->Args();
+ ASSERT_EQ(2u, args.Length());
+
+ auto new_arg1 = args[0]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg1->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(1_u, new_arg1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto new_arg2 = args[1]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg2->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(2_u, new_arg2->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_ExitLoopTest, CloneNoArgs) {
+ auto* loop = b.Loop();
+ auto* e = b.ExitLoop(loop);
+
+ auto* new_loop = clone_ctx.Clone(loop);
+ auto* new_exit = clone_ctx.Clone(e);
+
+ EXPECT_EQ(new_loop, new_exit->Loop());
+ EXPECT_TRUE(new_exit->Args().IsEmpty());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/exit_switch.cc b/src/tint/lang/core/ir/exit_switch.cc
index 7a70073..2216c8f 100644
--- a/src/tint/lang/core/ir/exit_switch.cc
+++ b/src/tint/lang/core/ir/exit_switch.cc
@@ -16,6 +16,8 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/lang/core/ir/switch.h"
@@ -30,6 +32,12 @@
ExitSwitch::~ExitSwitch() = default;
+ExitSwitch* ExitSwitch::Clone(CloneContext& ctx) {
+ auto* new_switch = ctx.Clone(Switch());
+ auto new_args = ctx.Clone<ExitSwitch::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<ExitSwitch>(new_switch, new_args);
+}
+
void ExitSwitch::SetSwitch(ir::Switch* s) {
SetControlInstruction(s);
}
diff --git a/src/tint/lang/core/ir/exit_switch.h b/src/tint/lang/core/ir/exit_switch.h
index 45a2733..c73ad6b 100644
--- a/src/tint/lang/core/ir/exit_switch.h
+++ b/src/tint/lang/core/ir/exit_switch.h
@@ -28,7 +28,7 @@
namespace tint::core::ir {
/// A exit switch instruction.
-class ExitSwitch : public Castable<ExitSwitch, Exit> {
+class ExitSwitch final : public Castable<ExitSwitch, Exit> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
@@ -39,6 +39,9 @@
explicit ExitSwitch(ir::Switch* sw, VectorRef<Value*> args = tint::Empty);
~ExitSwitch() override;
+ /// @copydoc Instruction::Clone()
+ ExitSwitch* Clone(CloneContext& ctx) override;
+
/// Re-associates the exit with the given switch instruction
/// @param s the new switch to exit from
void SetSwitch(ir::Switch* s);
diff --git a/src/tint/lang/core/ir/exit_switch_test.cc b/src/tint/lang/core/ir/exit_switch_test.cc
index d476cda..23f2f9c 100644
--- a/src/tint/lang/core/ir/exit_switch_test.cc
+++ b/src/tint/lang/core/ir/exit_switch_test.cc
@@ -53,5 +53,40 @@
EXPECT_FALSE(exit->Alive());
}
+TEST_F(IR_ExitSwitchTest, Clone) {
+ auto* arg1 = b.Constant(1_u);
+ auto* arg2 = b.Constant(2_u);
+ auto* switch_ = b.Switch(true);
+ auto* e = b.ExitSwitch(switch_, arg1, arg2);
+
+ auto* new_switch = clone_ctx.Clone(switch_);
+ auto* new_exit = clone_ctx.Clone(e);
+
+ EXPECT_NE(e, new_exit);
+ EXPECT_EQ(new_switch, new_exit->Switch());
+
+ auto args = new_exit->Args();
+ ASSERT_EQ(2u, args.Length());
+
+ auto new_arg1 = args[0]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg1->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(1_u, new_arg1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto new_arg2 = args[1]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg2->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(2_u, new_arg2->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_ExitSwitchTest, CloneNoArgs) {
+ auto* switch_ = b.Switch(true);
+ auto* e = b.ExitSwitch(switch_);
+
+ auto* new_switch = clone_ctx.Clone(switch_);
+ auto* new_exit = clone_ctx.Clone(e);
+
+ EXPECT_EQ(new_switch, new_exit->Switch());
+ EXPECT_TRUE(new_exit->Args().IsEmpty());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/function.cc b/src/tint/lang/core/ir/function.cc
index 28a23f2..57309c0 100644
--- a/src/tint/lang/core/ir/function.cc
+++ b/src/tint/lang/core/ir/function.cc
@@ -14,6 +14,8 @@
#include "src/tint/lang/core/ir/function.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/containers/predicates.h"
#include "src/tint/utils/ice/ice.h"
@@ -32,6 +34,22 @@
Function::~Function() = default;
+Function* Function::Clone(CloneContext& ctx) {
+ auto* new_func = ctx.ir.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;
+ new_func->return_.location = return_.location;
+ new_func->return_.invariant = return_.invariant;
+
+ ctx.Replace(this, new_func);
+ block_->CloneInto(ctx, new_func->block_);
+
+ ctx.ir.SetName(new_func, ctx.ir.NameOf(this).Name());
+ ctx.ir.functions.Push(new_func);
+ return new_func;
+}
+
void Function::SetParams(VectorRef<FunctionParam*> params) {
params_ = std::move(params);
TINT_ASSERT(!params_.Any(IsNull));
diff --git a/src/tint/lang/core/ir/function.h b/src/tint/lang/core/ir/function.h
index c24db41..34a9276 100644
--- a/src/tint/lang/core/ir/function.h
+++ b/src/tint/lang/core/ir/function.h
@@ -67,6 +67,9 @@
std::optional<std::array<uint32_t, 3>> wg_size = {});
~Function() override;
+ /// @copydoc Instruction::Clone()
+ Function* Clone(CloneContext& ctx) override;
+
/// Sets the function stage
/// @param stage the stage to set
void SetStage(PipelineStage stage) { pipeline_stage_ = stage; }
diff --git a/src/tint/lang/core/ir/function_param.cc b/src/tint/lang/core/ir/function_param.cc
index 877364e..8d667ee 100644
--- a/src/tint/lang/core/ir/function_param.cc
+++ b/src/tint/lang/core/ir/function_param.cc
@@ -14,6 +14,8 @@
#include "src/tint/lang/core/ir/function_param.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::FunctionParam);
@@ -58,4 +60,18 @@
return "<unknown>";
}
+FunctionParam* FunctionParam::Clone(CloneContext& ctx) {
+ auto* out = ctx.ir.values.Create<FunctionParam>(type_);
+ out->builtin_ = builtin_;
+ out->location_ = location_;
+ out->binding_point_ = binding_point_;
+ out->invariant_ = invariant_;
+
+ auto name = ctx.ir.NameOf(this);
+ if (name.IsValid()) {
+ ctx.ir.SetName(out, name);
+ }
+ return out;
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/function_param.h b/src/tint/lang/core/ir/function_param.h
index d9bf75c..f3b83b7 100644
--- a/src/tint/lang/core/ir/function_param.h
+++ b/src/tint/lang/core/ir/function_param.h
@@ -67,6 +67,9 @@
/// @returns the type of the var
const core::type::Type* Type() override { return type_; }
+ /// @copydoc Value::Clone()
+ FunctionParam* Clone(CloneContext& ctx) override;
+
/// Sets the builtin information. Note, it is currently an error if the builtin is already set.
/// @param val the builtin to set
void SetBuiltin(FunctionParam::Builtin val) {
diff --git a/src/tint/lang/core/ir/function_param_test.cc b/src/tint/lang/core/ir/function_param_test.cc
index 1ca35d7..eb1b4c1 100644
--- a/src/tint/lang/core/ir/function_param_test.cc
+++ b/src/tint/lang/core/ir/function_param_test.cc
@@ -12,8 +12,10 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/lang/core/ir/function_param.h"
+#include <string>
+
#include "gtest/gtest-spi.h"
+#include "src/tint/lang/core/ir/function_param.h"
#include "src/tint/lang/core/ir/ir_helper_test.h"
namespace tint::core::ir {
@@ -44,5 +46,53 @@
"");
}
+TEST_F(IR_FunctionParamTest, CloneEmpty) {
+ auto* fp = b.FunctionParam(mod.Types().f32());
+
+ auto* new_fp = clone_ctx.Clone(fp);
+ EXPECT_EQ(new_fp->Type(), mod.Types().f32());
+ EXPECT_FALSE(new_fp->Builtin().has_value());
+ EXPECT_FALSE(new_fp->Location().has_value());
+ EXPECT_FALSE(new_fp->BindingPoint().has_value());
+ EXPECT_FALSE(new_fp->Invariant());
+}
+
+TEST_F(IR_FunctionParamTest, Clone) {
+ auto* fp = b.FunctionParam(mod.Types().f32());
+ fp->SetBuiltin(FunctionParam::Builtin::kVertexIndex);
+ fp->SetLocation(
+ 1, Interpolation{core::InterpolationType::kFlat, core::InterpolationSampling::kCentroid});
+ fp->SetInvariant(true);
+ fp->SetBindingPoint(1, 2);
+
+ auto* new_fp = clone_ctx.Clone(fp);
+
+ EXPECT_NE(fp, new_fp);
+ EXPECT_EQ(new_fp->Type(), mod.Types().f32());
+
+ EXPECT_TRUE(new_fp->Builtin().has_value());
+ EXPECT_EQ(FunctionParam::Builtin::kVertexIndex, new_fp->Builtin().value());
+
+ EXPECT_TRUE(new_fp->Location().has_value());
+ auto loc = new_fp->Location();
+ EXPECT_EQ(1u, loc->value);
+ EXPECT_EQ(core::InterpolationType::kFlat, loc->interpolation->type);
+ EXPECT_EQ(core::InterpolationSampling::kCentroid, loc->interpolation->sampling);
+
+ EXPECT_TRUE(new_fp->BindingPoint().has_value());
+ auto bp = new_fp->BindingPoint();
+ EXPECT_EQ(1u, bp->group);
+ EXPECT_EQ(2u, bp->binding);
+
+ EXPECT_TRUE(new_fp->Invariant());
+}
+
+TEST_F(IR_FunctionParamTest, CloneWithName) {
+ auto* fp = b.FunctionParam("fp", mod.Types().f32());
+ auto* new_fp = clone_ctx.Clone(fp);
+
+ EXPECT_EQ(std::string("fp"), mod.NameOf(new_fp).Name());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/function_test.cc b/src/tint/lang/core/ir/function_test.cc
index 862c9ab..6739793 100644
--- a/src/tint/lang/core/ir/function_test.cc
+++ b/src/tint/lang/core/ir/function_test.cc
@@ -12,8 +12,10 @@
// See the License for the specific language governing permissions and
// limitations under the License.
-#include "src/tint/lang/core/ir/function.h"
+#include <string>
+
#include "gtest/gtest-spi.h"
+#include "src/tint/lang/core/ir/function.h"
#include "src/tint/lang/core/ir/ir_helper_test.h"
namespace tint::core::ir {
@@ -66,5 +68,61 @@
"");
}
+TEST_F(IR_FunctionTest, Clone) {
+ auto* f =
+ b.Function("my_func", mod.Types().i32(), Function::PipelineStage::kCompute, {{2, 3, 4}});
+ f->SetReturnBuiltin(Function::ReturnBuiltin::kFragDepth);
+ f->SetReturnLocation(
+ 1, Interpolation{core::InterpolationType::kFlat, core::InterpolationSampling::kCentroid});
+ f->SetReturnInvariant(true);
+
+ auto* param1 = b.FunctionParam("a", mod.Types().i32());
+ auto* param2 = b.FunctionParam("b", mod.Types().f32());
+ f->SetParams({param1, param2});
+
+ auto* new_param1 = clone_ctx.Clone(param1);
+ auto* new_param2 = clone_ctx.Clone(param2);
+ auto* new_f = clone_ctx.Clone(f);
+
+ EXPECT_NE(f, new_f);
+ EXPECT_EQ(std::string("my_func"), mod.NameOf(new_f).Name());
+
+ EXPECT_EQ(Function::PipelineStage::kCompute, new_f->Stage());
+ EXPECT_TRUE(new_f->WorkgroupSize().has_value());
+ auto wg = new_f->WorkgroupSize().value();
+ EXPECT_EQ(2u, wg[0]);
+ EXPECT_EQ(3u, wg[1]);
+ EXPECT_EQ(4u, wg[2]);
+
+ EXPECT_EQ(mod.Types().i32(), new_f->ReturnType());
+
+ EXPECT_TRUE(new_f->ReturnBuiltin().has_value());
+ EXPECT_EQ(Function::ReturnBuiltin::kFragDepth, new_f->ReturnBuiltin().value());
+
+ EXPECT_TRUE(new_f->ReturnLocation().has_value());
+ auto loc = new_f->ReturnLocation().value();
+ EXPECT_EQ(1u, loc.value);
+ EXPECT_EQ(core::InterpolationType::kFlat, loc.interpolation->type);
+ EXPECT_EQ(core::InterpolationSampling::kCentroid, loc.interpolation->sampling);
+
+ EXPECT_TRUE(new_f->ReturnInvariant());
+
+ EXPECT_EQ(2u, new_f->Params().Length());
+ EXPECT_EQ(new_param1, new_f->Params()[0]);
+ EXPECT_EQ(new_param2, new_f->Params()[1]);
+
+ EXPECT_EQ(new_f, mod.functions.Back());
+}
+
+TEST_F(IR_FunctionTest, CloneWithExits) {
+ auto* f = b.Function("my_func", mod.Types().void_());
+ b.Append(f->Block(), [&] { b.Return(f); });
+
+ auto* new_f = clone_ctx.Clone(f);
+ EXPECT_EQ(1u, new_f->Block()->Length());
+ EXPECT_TRUE(new_f->Block()->Front()->Is<Return>());
+ EXPECT_EQ(new_f, new_f->Block()->Front()->As<Return>()->Func());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/if.cc b/src/tint/lang/core/ir/if.cc
index a57adf5..6368db2 100644
--- a/src/tint/lang/core/ir/if.cc
+++ b/src/tint/lang/core/ir/if.cc
@@ -16,6 +16,8 @@
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::If);
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/utils/ice/ice.h"
@@ -46,4 +48,17 @@
}
}
+If* If::Clone(CloneContext& ctx) {
+ auto* new_cond = ctx.Clone(Condition());
+ 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>(new_cond, new_true, new_false);
+ ctx.Replace(this, new_if);
+
+ true_->CloneInto(ctx, new_true);
+ false_->CloneInto(ctx, new_false);
+ return new_if;
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/if.h b/src/tint/lang/core/ir/if.h
index aa68508..d4f031b 100644
--- a/src/tint/lang/core/ir/if.h
+++ b/src/tint/lang/core/ir/if.h
@@ -42,7 +42,7 @@
/// ▼
/// out
/// ```
-class If : public Castable<If, ControlInstruction> {
+class If final : public Castable<If, ControlInstruction> {
public:
/// The index of the condition operand
static constexpr size_t kConditionOperandOffset = 0;
@@ -54,6 +54,9 @@
If(Value* cond, ir::Block* t, ir::Block* f);
~If() override;
+ /// @copydoc Instruction::Clone()
+ If* Clone(CloneContext& ctx) override;
+
/// @copydoc ControlInstruction::ForeachBlock
void ForeachBlock(const std::function<void(ir::Block*)>& cb) override;
diff --git a/src/tint/lang/core/ir/if_test.cc b/src/tint/lang/core/ir/if_test.cc
index 8ecb742..01a8538 100644
--- a/src/tint/lang/core/ir/if_test.cc
+++ b/src/tint/lang/core/ir/if_test.cc
@@ -63,5 +63,34 @@
"");
}
+TEST_F(IR_IfTest, Clone) {
+ auto* if_ = b.If(b.Constant(true));
+ auto* new_if = clone_ctx.Clone(if_);
+
+ EXPECT_NE(if_, new_if);
+
+ auto new_cond = new_if->Condition()->As<Constant>()->Value();
+ ASSERT_TRUE(new_cond->Is<core::constant::Scalar<bool>>());
+ EXPECT_TRUE(new_cond->As<core::constant::Scalar<bool>>()->ValueAs<bool>());
+
+ EXPECT_NE(nullptr, new_if->True());
+ EXPECT_NE(nullptr, new_if->False());
+ EXPECT_NE(if_->True(), new_if->True());
+ EXPECT_NE(if_->False(), new_if->False());
+}
+
+TEST_F(IR_IfTest, CloneWithExits) {
+ If* new_if = nullptr;
+ {
+ auto* if_ = b.If(true);
+ b.Append(if_->True(), [&] { b.ExitIf(if_); });
+ new_if = clone_ctx.Clone(if_);
+ }
+
+ ASSERT_EQ(1u, new_if->True()->Length());
+ EXPECT_TRUE(new_if->True()->Front()->Is<ExitIf>());
+ EXPECT_EQ(new_if, new_if->True()->Front()->As<ExitIf>()->If());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/instruction.h b/src/tint/lang/core/ir/instruction.h
index d2503d2..7a6f790 100644
--- a/src/tint/lang/core/ir/instruction.h
+++ b/src/tint/lang/core/ir/instruction.h
@@ -25,6 +25,7 @@
// Forward declarations
namespace tint::core::ir {
class Block;
+class CloneContext;
} // namespace tint::core::ir
namespace tint::core::ir {
@@ -62,6 +63,10 @@
/// @returns the friendly name for the instruction
virtual std::string FriendlyName() = 0;
+ /// @param ctx the CloneContext used to clone this instruction
+ /// @returns a clone of this instruction
+ virtual Instruction* Clone(CloneContext& ctx) = 0;
+
/// @returns true if the Instruction has not been destroyed with Destroy()
bool Alive() const { return !flags_.Contains(Flag::kDead); }
diff --git a/src/tint/lang/core/ir/instruction_result.cc b/src/tint/lang/core/ir/instruction_result.cc
index cc78b33..01e8d61 100644
--- a/src/tint/lang/core/ir/instruction_result.cc
+++ b/src/tint/lang/core/ir/instruction_result.cc
@@ -14,8 +14,10 @@
#include "src/tint/lang/core/ir/instruction_result.h"
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/constant.h"
#include "src/tint/lang/core/ir/instruction.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::InstructionResult);
@@ -33,4 +35,10 @@
Base::Destroy();
}
+InstructionResult* InstructionResult::Clone(CloneContext& ctx) {
+ // Do not clone the `Source`. It will be set when this result is placed in the new parent
+ // instruction.
+ return ctx.ir.values.Create<InstructionResult>(type_);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/instruction_result.h b/src/tint/lang/core/ir/instruction_result.h
index d139a3c..c3b747c 100644
--- a/src/tint/lang/core/ir/instruction_result.h
+++ b/src/tint/lang/core/ir/instruction_result.h
@@ -36,6 +36,9 @@
/// @returns the type of the value
const core::type::Type* Type() override { return type_; }
+ /// @copydoc Value::Clone()
+ InstructionResult* Clone(CloneContext& ctx) override;
+
/// Sets the type of the value to @p type
/// @param type the new type of the value
void SetType(const core::type::Type* type) { type_ = type; }
diff --git a/src/tint/lang/core/ir/instruction_result_test.cc b/src/tint/lang/core/ir/instruction_result_test.cc
index 68951b7..0991ea3 100644
--- a/src/tint/lang/core/ir/instruction_result_test.cc
+++ b/src/tint/lang/core/ir/instruction_result_test.cc
@@ -35,5 +35,14 @@
"");
}
+TEST_F(IR_InstructionResultTest, Clone) {
+ auto* val = b.Add(mod.Types().i32(), 1_i, 2_i)->Result();
+ auto* new_res = clone_ctx.Clone(val);
+
+ EXPECT_NE(val, new_res);
+ EXPECT_EQ(nullptr, new_res->Source());
+ EXPECT_EQ(mod.Types().i32(), new_res->Type());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/ir_helper_test.h b/src/tint/lang/core/ir/ir_helper_test.h
index 9e13af9..8d749c2 100644
--- a/src/tint/lang/core/ir/ir_helper_test.h
+++ b/src/tint/lang/core/ir/ir_helper_test.h
@@ -17,6 +17,7 @@
#include "gtest/gtest.h"
#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/module.h"
namespace tint::core::ir {
@@ -34,6 +35,9 @@
Builder b{mod};
/// The type manager
core::type::Manager& ty{mod.Types()};
+
+ /// CloneContext
+ CloneContext clone_ctx{mod};
};
using IRTestHelper = IRTestHelperBase<testing::Test>;
diff --git a/src/tint/lang/core/ir/let.cc b/src/tint/lang/core/ir/let.cc
index 4e7236a..b843ee5 100644
--- a/src/tint/lang/core/ir/let.cc
+++ b/src/tint/lang/core/ir/let.cc
@@ -13,6 +13,9 @@
// limitations under the License.
#include "src/tint/lang/core/ir/let.h"
+
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/store.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Let);
@@ -26,4 +29,15 @@
Let::~Let() = default;
+Let* Let::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_val = ctx.Clone(Value());
+ auto* new_let = ctx.ir.instructions.Create<Let>(new_result, new_val);
+
+ auto name = ctx.ir.NameOf(this);
+ ctx.ir.SetName(new_let, name.Name());
+
+ return new_let;
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/let.h b/src/tint/lang/core/ir/let.h
index 5c0355b..4077534 100644
--- a/src/tint/lang/core/ir/let.h
+++ b/src/tint/lang/core/ir/let.h
@@ -22,7 +22,7 @@
namespace tint::core::ir {
/// A no-op instruction in the IR, used to position and name a value
-class Let : public Castable<Let, OperandInstruction<1, 1>> {
+class Let final : public Castable<Let, OperandInstruction<1, 1>> {
public:
/// The offset in Operands() for the value
static constexpr size_t kValueOperandOffset = 0;
@@ -33,6 +33,9 @@
Let(InstructionResult* result, Value* value);
~Let() override;
+ /// @copydoc Instruction::Clone()
+ Let* Clone(CloneContext& ctx) override;
+
/// @returns the value
ir::Value* Value() { return operands_[kValueOperandOffset]; }
diff --git a/src/tint/lang/core/ir/let_test.cc b/src/tint/lang/core/ir/let_test.cc
index b0f716c..6769734 100644
--- a/src/tint/lang/core/ir/let_test.cc
+++ b/src/tint/lang/core/ir/let_test.cc
@@ -49,5 +49,23 @@
EXPECT_EQ(let->Result()->Type(), value->Type());
}
+TEST_F(IR_LetTest, Clone) {
+ auto* value = b.Constant(4_f);
+ auto* let = b.Let("l", value);
+
+ auto* new_let = clone_ctx.Clone(let);
+
+ EXPECT_NE(let, new_let);
+ EXPECT_NE(nullptr, new_let->Result());
+ EXPECT_NE(let->Result(), new_let->Result());
+
+ auto new_val = new_let->Value()->As<Constant>()->Value();
+ ASSERT_TRUE(new_val->Is<core::constant::Scalar<f32>>());
+ EXPECT_FLOAT_EQ(4_f, new_val->As<core::constant::Scalar<f32>>()->ValueAs<f32>());
+
+ EXPECT_EQ(std::string("l"), mod.NameOf(new_let).Name());
+ EXPECT_EQ(std::string("l"), mod.NameOf(new_let->Result()).Name());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/load.cc b/src/tint/lang/core/ir/load.cc
index d60e151..361643b 100644
--- a/src/tint/lang/core/ir/load.cc
+++ b/src/tint/lang/core/ir/load.cc
@@ -14,6 +14,8 @@
#include "src/tint/lang/core/ir/load.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/utils/ice/ice.h"
@@ -33,4 +35,10 @@
Load::~Load() = default;
+Load* Load::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_from = ctx.Clone(From());
+ return ctx.ir.instructions.Create<Load>(new_result, new_from);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/load.h b/src/tint/lang/core/ir/load.h
index 55fb983..1bf1ca1 100644
--- a/src/tint/lang/core/ir/load.h
+++ b/src/tint/lang/core/ir/load.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A load instruction in the IR.
-class Load : public Castable<Load, OperandInstruction<1, 1>> {
+class Load final : public Castable<Load, OperandInstruction<1, 1>> {
public:
/// The offset in Operands() for the from value
static constexpr size_t kFromOperandOffset = 0;
@@ -35,6 +35,9 @@
~Load() override;
+ /// @copydoc Instruction::Clone()
+ Load* Clone(CloneContext& ctx) override;
+
/// @returns the value being loaded from
Value* From() { return operands_[kFromOperandOffset]; }
diff --git a/src/tint/lang/core/ir/load_test.cc b/src/tint/lang/core/ir/load_test.cc
index fe4c6fe..1a07476 100644
--- a/src/tint/lang/core/ir/load_test.cc
+++ b/src/tint/lang/core/ir/load_test.cc
@@ -69,5 +69,19 @@
"");
}
+TEST_F(IR_LoadTest, Clone) {
+ auto* var = b.Var(ty.ptr<function, i32>());
+ auto* inst = b.Load(var);
+
+ auto* new_var = clone_ctx.Clone(var);
+ auto* new_inst = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, new_inst);
+ EXPECT_NE(nullptr, new_inst->Result());
+ EXPECT_NE(inst->Result(), new_inst->Result());
+
+ EXPECT_EQ(new_var->Result(), new_inst->From());
+}
+
} // namespace
} // 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 e0969b9..aebf4d5 100644
--- a/src/tint/lang/core/ir/load_vector_element.cc
+++ b/src/tint/lang/core/ir/load_vector_element.cc
@@ -14,6 +14,9 @@
#include "src/tint/lang/core/ir/load_vector_element.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::LoadVectorElement);
namespace tint::core::ir {
@@ -28,4 +31,11 @@
LoadVectorElement::~LoadVectorElement() = default;
+LoadVectorElement* LoadVectorElement::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_from = ctx.Clone(From());
+ auto* new_index = ctx.Clone(Index());
+ return ctx.ir.instructions.Create<LoadVectorElement>(new_result, new_from, new_index);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/load_vector_element.h b/src/tint/lang/core/ir/load_vector_element.h
index cd7e206..c141a0b 100644
--- a/src/tint/lang/core/ir/load_vector_element.h
+++ b/src/tint/lang/core/ir/load_vector_element.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A load instruction for a single vector element in the IR.
-class LoadVectorElement : public Castable<LoadVectorElement, OperandInstruction<3, 0>> {
+class LoadVectorElement final : public Castable<LoadVectorElement, OperandInstruction<3, 0>> {
public:
/// The offset in Operands() for the `from` value
static constexpr size_t kFromOperandOffset = 0;
@@ -38,6 +38,9 @@
LoadVectorElement(InstructionResult* result, ir::Value* from, ir::Value* index);
~LoadVectorElement() override;
+ /// @copydoc Instruction::Clone()
+ LoadVectorElement* Clone(CloneContext& ctx) override;
+
/// @returns the vector pointer value
ir::Value* From() { return operands_[kFromOperandOffset]; }
diff --git a/src/tint/lang/core/ir/load_vector_element_test.cc b/src/tint/lang/core/ir/load_vector_element_test.cc
index 35a755d..744cf2b 100644
--- a/src/tint/lang/core/ir/load_vector_element_test.cc
+++ b/src/tint/lang/core/ir/load_vector_element_test.cc
@@ -58,5 +58,23 @@
EXPECT_FALSE(inst->HasMultiResults());
}
+TEST_F(IR_LoadVectorElementTest, Clone) {
+ auto* from = b.Var(ty.ptr<private_, vec3<i32>>());
+ auto* inst = b.LoadVectorElement(from, 2_i);
+
+ auto* new_from = clone_ctx.Clone(from);
+ auto* new_inst = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, new_inst);
+ EXPECT_NE(nullptr, new_inst->Result());
+ EXPECT_NE(inst->Result(), new_inst->Result());
+
+ EXPECT_EQ(new_from->Result(), new_inst->From());
+
+ auto new_idx = new_inst->Index()->As<Constant>()->Value();
+ ASSERT_TRUE(new_idx->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(2_i, new_idx->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/loop.cc b/src/tint/lang/core/ir/loop.cc
index 5112718..ada2516 100644
--- a/src/tint/lang/core/ir/loop.cc
+++ b/src/tint/lang/core/ir/loop.cc
@@ -16,6 +16,8 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/utils/ice/ice.h"
@@ -42,6 +44,21 @@
Loop::~Loop() = default;
+Loop* Loop::Clone(CloneContext& ctx) {
+ auto* new_init = ctx.ir.blocks.Create<MultiInBlock>();
+ 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);
+ ctx.Replace(this, new_loop);
+
+ initializer_->CloneInto(ctx, new_init);
+ body_->CloneInto(ctx, new_body);
+ continuing_->CloneInto(ctx, new_continuing);
+
+ return new_loop;
+}
+
void Loop::ForeachBlock(const std::function<void(ir::Block*)>& cb) {
if (initializer_) {
cb(initializer_);
diff --git a/src/tint/lang/core/ir/loop.h b/src/tint/lang/core/ir/loop.h
index 689b9e2..5e97226 100644
--- a/src/tint/lang/core/ir/loop.h
+++ b/src/tint/lang/core/ir/loop.h
@@ -56,7 +56,7 @@
/// out
///
/// ```
-class Loop : public Castable<Loop, ControlInstruction> {
+class Loop final : public Castable<Loop, ControlInstruction> {
public:
/// Constructor
/// @param i the initializer block
@@ -65,6 +65,9 @@
Loop(ir::Block* i, ir::MultiInBlock* b, ir::MultiInBlock* c);
~Loop() override;
+ /// @copydoc Instruction::Clone()
+ Loop* Clone(CloneContext& ctx) override;
+
/// @copydoc ControlInstruction::ForeachBlock
void ForeachBlock(const std::function<void(ir::Block*)>& cb) override;
diff --git a/src/tint/lang/core/ir/loop_test.cc b/src/tint/lang/core/ir/loop_test.cc
index 1a4beb5..89699f0 100644
--- a/src/tint/lang/core/ir/loop_test.cc
+++ b/src/tint/lang/core/ir/loop_test.cc
@@ -65,5 +65,57 @@
"");
}
+TEST_F(IR_LoopTest, Clone) {
+ auto* loop = b.Loop();
+ auto* new_loop = clone_ctx.Clone(loop);
+
+ EXPECT_NE(loop, new_loop);
+ EXPECT_FALSE(new_loop->HasResults());
+ EXPECT_EQ(0u, new_loop->Exits().Count());
+ EXPECT_NE(nullptr, new_loop->Initializer());
+ EXPECT_NE(loop->Initializer(), new_loop->Initializer());
+
+ EXPECT_NE(nullptr, new_loop->Body());
+ EXPECT_NE(loop->Body(), new_loop->Body());
+
+ EXPECT_NE(nullptr, new_loop->Continuing());
+ EXPECT_NE(loop->Continuing(), new_loop->Continuing());
+}
+
+TEST_F(IR_LoopTest, CloneWithExits) {
+ Loop* new_loop = nullptr;
+ {
+ auto* loop = b.Loop();
+ b.Append(loop->Body(), [&] {
+ auto* if_ = b.If(true);
+ b.Append(if_->True(), [&] { b.Continue(loop); });
+ b.Append(if_->False(), [&] { b.ExitLoop(loop); });
+ b.Append(loop->Continuing(), [&] { b.BreakIf(loop, false); });
+
+ b.NextIteration(loop);
+ });
+ new_loop = clone_ctx.Clone(loop);
+ }
+
+ ASSERT_EQ(2u, new_loop->Body()->Length());
+ EXPECT_TRUE(new_loop->Body()->Front()->Is<If>());
+
+ auto* new_if = new_loop->Body()->Front()->As<If>();
+ ASSERT_EQ(1u, new_if->True()->Length());
+ EXPECT_TRUE(new_if->True()->Front()->Is<Continue>());
+ EXPECT_EQ(new_loop, new_if->True()->Front()->As<Continue>()->Loop());
+
+ ASSERT_EQ(1u, new_if->False()->Length());
+ EXPECT_TRUE(new_if->False()->Front()->Is<ExitLoop>());
+ EXPECT_EQ(new_loop, new_if->False()->Front()->As<ExitLoop>()->Loop());
+
+ ASSERT_EQ(1u, new_loop->Continuing()->Length());
+ EXPECT_TRUE(new_loop->Continuing()->Front()->Is<BreakIf>());
+ EXPECT_EQ(new_loop, new_loop->Continuing()->Front()->As<BreakIf>()->Loop());
+
+ EXPECT_TRUE(new_loop->Body()->Back()->Is<NextIteration>());
+ EXPECT_EQ(new_loop, new_loop->Body()->Back()->As<NextIteration>()->Loop());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/multi_in_block.cc b/src/tint/lang/core/ir/multi_in_block.cc
index 6916512..676691b 100644
--- a/src/tint/lang/core/ir/multi_in_block.cc
+++ b/src/tint/lang/core/ir/multi_in_block.cc
@@ -14,6 +14,9 @@
#include "src/tint/lang/core/ir/multi_in_block.h"
+#include "src/tint/lang/core/ir/block_param.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/containers/predicates.h"
#include "src/tint/utils/ice/ice.h"
@@ -25,6 +28,19 @@
MultiInBlock::~MultiInBlock() = default;
+MultiInBlock* MultiInBlock::Clone(CloneContext&) {
+ TINT_UNREACHABLE() << "blocks must be cloned with CloneInto";
+ return nullptr;
+}
+
+void MultiInBlock::CloneInto(CloneContext& ctx, Block* out) {
+ TINT_ASSERT(out->Is<MultiInBlock>());
+
+ auto new_params = ctx.Clone(params_);
+ out->As<MultiInBlock>()->SetParams(new_params);
+ Block::CloneInto(ctx, out);
+}
+
void MultiInBlock::SetParams(VectorRef<BlockParam*> params) {
params_ = std::move(params);
}
diff --git a/src/tint/lang/core/ir/multi_in_block.h b/src/tint/lang/core/ir/multi_in_block.h
index 1fecb6d..5934dca 100644
--- a/src/tint/lang/core/ir/multi_in_block.h
+++ b/src/tint/lang/core/ir/multi_in_block.h
@@ -35,6 +35,12 @@
MultiInBlock();
~MultiInBlock() override;
+ /// @copydoc Block::Clone()
+ MultiInBlock* Clone(CloneContext& ctx) override;
+
+ /// @copydoc Block::CloneInto()
+ void CloneInto(CloneContext& ctx, Block* out) override;
+
/// Sets the params to the block
/// @param params the params for the block
void SetParams(VectorRef<BlockParam*> params);
diff --git a/src/tint/lang/core/ir/multi_in_block_test.cc b/src/tint/lang/core/ir/multi_in_block_test.cc
index f2145af..2e7552e 100644
--- a/src/tint/lang/core/ir/multi_in_block_test.cc
+++ b/src/tint/lang/core/ir/multi_in_block_test.cc
@@ -35,5 +35,43 @@
"");
}
+TEST_F(IR_MultiInBlockTest, CloneInto) {
+ auto* loop = b.Loop();
+
+ auto* blk = b.MultiInBlock();
+ auto* add = b.Add(mod.Types().i32(), 1_i, 2_i);
+ blk->Append(add);
+ blk->SetParams({b.BlockParam(mod.Types().i32()), b.BlockParam(mod.Types().f32())});
+ blk->SetParent(loop);
+
+ auto* terminate = b.TerminateInvocation();
+ blk->AddInboundSiblingBranch(terminate);
+
+ auto* new_blk = b.MultiInBlock();
+ blk->CloneInto(clone_ctx, new_blk);
+
+ EXPECT_EQ(0u, new_blk->InboundSiblingBranches().Length());
+
+ EXPECT_EQ(2u, new_blk->Params().Length());
+ EXPECT_EQ(mod.Types().i32(), new_blk->Params()[0]->Type());
+ EXPECT_EQ(mod.Types().f32(), new_blk->Params()[1]->Type());
+
+ EXPECT_EQ(nullptr, new_blk->Parent());
+
+ EXPECT_EQ(1u, new_blk->Length());
+ EXPECT_NE(add, new_blk->Front());
+ EXPECT_TRUE(new_blk->Front()->Is<Binary>());
+ EXPECT_EQ(Binary::Kind::kAdd, new_blk->Front()->As<Binary>()->Kind());
+}
+
+TEST_F(IR_MultiInBlockTest, CloneEmpty) {
+ auto* blk = b.MultiInBlock();
+ auto* new_blk = b.MultiInBlock();
+ blk->CloneInto(clone_ctx, new_blk);
+
+ EXPECT_EQ(0u, new_blk->InboundSiblingBranches().Length());
+ EXPECT_EQ(0u, new_blk->Params().Length());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/next_iteration.cc b/src/tint/lang/core/ir/next_iteration.cc
index 1aeb9f4..1c3fe65 100644
--- a/src/tint/lang/core/ir/next_iteration.cc
+++ b/src/tint/lang/core/ir/next_iteration.cc
@@ -16,7 +16,9 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/loop.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/multi_in_block.h"
#include "src/tint/utils/ice/ice.h"
@@ -37,4 +39,10 @@
NextIteration::~NextIteration() = default;
+NextIteration* NextIteration::Clone(CloneContext& ctx) {
+ auto* new_loop = ctx.Clone(loop_);
+ auto new_args = ctx.Clone<NextIteration::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<NextIteration>(new_loop, new_args);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/next_iteration.h b/src/tint/lang/core/ir/next_iteration.h
index 87de579..f845dee 100644
--- a/src/tint/lang/core/ir/next_iteration.h
+++ b/src/tint/lang/core/ir/next_iteration.h
@@ -28,7 +28,7 @@
namespace tint::core::ir {
/// A next iteration instruction.
-class NextIteration : public Castable<NextIteration, Terminator> {
+class NextIteration final : public Castable<NextIteration, Terminator> {
public:
/// The base offset in Operands() for the args
static constexpr size_t kArgsOperandOffset = 0;
@@ -39,6 +39,9 @@
explicit NextIteration(ir::Loop* loop, VectorRef<Value*> args = tint::Empty);
~NextIteration() override;
+ /// @copydoc Instruction::Clone()
+ NextIteration* Clone(CloneContext& ctx) override;
+
/// @returns the loop being iterated
ir::Loop* Loop() { return loop_; }
diff --git a/src/tint/lang/core/ir/next_iteration_test.cc b/src/tint/lang/core/ir/next_iteration_test.cc
index 3d81331..f1efc9e 100644
--- a/src/tint/lang/core/ir/next_iteration_test.cc
+++ b/src/tint/lang/core/ir/next_iteration_test.cc
@@ -39,5 +39,38 @@
EXPECT_FALSE(inst->HasMultiResults());
}
+TEST_F(IR_NextIterationTest, Clone) {
+ auto* arg1 = b.Constant(1_u);
+ auto* arg2 = b.Constant(2_u);
+
+ auto* loop = b.Loop();
+ auto* inst = b.NextIteration(loop, arg1, arg2);
+
+ auto* new_loop = clone_ctx.Clone(loop);
+ auto* new_inst = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, new_inst);
+ EXPECT_EQ(new_loop, new_inst->Loop());
+
+ auto args = new_inst->Args();
+ EXPECT_EQ(2u, args.Length());
+
+ auto* val0 = args[0]->As<Constant>()->Value();
+ EXPECT_EQ(1_u, val0->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto* val1 = args[1]->As<Constant>()->Value();
+ EXPECT_EQ(2_u, val1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_NextIterationTest, CloneNoArgs) {
+ auto* loop = b.Loop();
+ auto* inst = b.NextIteration(loop);
+
+ auto* new_loop = clone_ctx.Clone(loop);
+ auto* new_inst = clone_ctx.Clone(inst);
+
+ EXPECT_EQ(new_loop, new_inst->Loop());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/operand_instruction.h b/src/tint/lang/core/ir/operand_instruction.h
index 49c2c85..980c696 100644
--- a/src/tint/lang/core/ir/operand_instruction.h
+++ b/src/tint/lang/core/ir/operand_instruction.h
@@ -134,6 +134,9 @@
Vector<ir::Value*, N> operands_;
/// The results of this instruction.
Vector<ir::InstructionResult*, R> results_;
+
+ /// The default number of operands
+ static constexpr size_t kDefaultNumOperands = N;
};
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/return.cc b/src/tint/lang/core/ir/return.cc
index ab6e841..49b5939 100644
--- a/src/tint/lang/core/ir/return.cc
+++ b/src/tint/lang/core/ir/return.cc
@@ -16,7 +16,9 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
#include "src/tint/lang/core/ir/function.h"
+#include "src/tint/lang/core/ir/module.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Return);
@@ -33,6 +35,12 @@
Return::~Return() = default;
+Return* Return::Clone(CloneContext& ctx) {
+ auto* new_func = ctx.Clone(Func());
+ auto new_val = Value() ? ctx.Clone(Value()) : nullptr;
+ return ctx.ir.instructions.Create<Return>(new_func, new_val);
+}
+
Function* Return::Func() const {
return tint::As<Function>(operands_[kFunctionOperandOffset]);
}
diff --git a/src/tint/lang/core/ir/return.h b/src/tint/lang/core/ir/return.h
index f5b8c8a..0f7c434 100644
--- a/src/tint/lang/core/ir/return.h
+++ b/src/tint/lang/core/ir/return.h
@@ -28,7 +28,7 @@
namespace tint::core::ir {
/// A return instruction.
-class Return : public Castable<Return, Terminator> {
+class Return final : public Castable<Return, Terminator> {
public:
/// The offset in Operands() for the function being returned
static constexpr size_t kFunctionOperandOffset = 0;
@@ -47,6 +47,9 @@
~Return() override;
+ /// @copydoc Instruction::Clone()
+ Return* Clone(CloneContext& ctx) override;
+
/// @returns the function being returned
Function* Func() const;
diff --git a/src/tint/lang/core/ir/return_test.cc b/src/tint/lang/core/ir/return_test.cc
index f08c4b5..dc68fd6 100644
--- a/src/tint/lang/core/ir/return_test.cc
+++ b/src/tint/lang/core/ir/return_test.cc
@@ -62,5 +62,33 @@
}
}
+TEST_F(IR_ReturnTest, Clone) {
+ auto* func = b.Function("func", ty.i32());
+ auto* ret = b.Return(func, b.Constant(1_i));
+
+ auto* new_func = clone_ctx.Clone(func);
+ auto* new_ret = clone_ctx.Clone(ret);
+
+ EXPECT_NE(ret, new_ret);
+ EXPECT_EQ(new_func, new_ret->Func());
+
+ EXPECT_EQ(1u, new_ret->Args().Length());
+
+ auto new_val = new_ret->Value()->As<Constant>()->Value();
+ ASSERT_TRUE(new_val->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(1_i, new_val->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+}
+
+TEST_F(IR_ReturnTest, CloneWithoutArgs) {
+ auto* func = b.Function("func", ty.i32());
+ auto* ret = b.Return(func);
+
+ auto* new_func = clone_ctx.Clone(func);
+ auto* new_ret = clone_ctx.Clone(ret);
+
+ EXPECT_EQ(new_func, new_ret->Func());
+ EXPECT_EQ(nullptr, new_ret->Value());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/store.cc b/src/tint/lang/core/ir/store.cc
index c333ede..8088087 100644
--- a/src/tint/lang/core/ir/store.cc
+++ b/src/tint/lang/core/ir/store.cc
@@ -14,6 +14,9 @@
#include "src/tint/lang/core/ir/store.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Store);
namespace tint::core::ir {
@@ -27,4 +30,10 @@
Store::~Store() = default;
+Store* Store::Clone(CloneContext& ctx) {
+ auto* new_to = ctx.Clone(To());
+ auto* new_from = ctx.Clone(From());
+ return ctx.ir.instructions.Create<Store>(new_to, new_from);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/store.h b/src/tint/lang/core/ir/store.h
index 924a6b8..3c5071e 100644
--- a/src/tint/lang/core/ir/store.h
+++ b/src/tint/lang/core/ir/store.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A store instruction in the IR.
-class Store : public Castable<Store, OperandInstruction<2, 0>> {
+class Store final : public Castable<Store, OperandInstruction<2, 0>> {
public:
/// The offset in Operands() for the `to` value
static constexpr size_t kToOperandOffset = 0;
@@ -37,6 +37,9 @@
Store(Value* to, Value* from);
~Store() override;
+ /// @copydoc Instruction::Clone()
+ Store* Clone(CloneContext& ctx) override;
+
/// @returns the value being stored too
Value* To() { return operands_[kToOperandOffset]; }
diff --git a/src/tint/lang/core/ir/store_test.cc b/src/tint/lang/core/ir/store_test.cc
index f13049b..c5677f0 100644
--- a/src/tint/lang/core/ir/store_test.cc
+++ b/src/tint/lang/core/ir/store_test.cc
@@ -58,5 +58,20 @@
EXPECT_FALSE(inst->HasMultiResults());
}
+TEST_F(IR_StoreTest, Clone) {
+ auto* v = b.Var("a", mod.Types().ptr<private_, i32>());
+ auto* s = b.Store(v, b.Constant(1_i));
+
+ auto* new_v = clone_ctx.Clone(v);
+ auto* new_s = clone_ctx.Clone(s);
+
+ EXPECT_NE(s, new_s);
+ EXPECT_EQ(new_v->Result(), new_s->To());
+
+ auto new_from = new_s->From()->As<Constant>()->Value();
+ ASSERT_TRUE(new_from->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(1_i, new_from->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+}
+
} // namespace
} // 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 c133211..7d4a45d 100644
--- a/src/tint/lang/core/ir/store_vector_element.cc
+++ b/src/tint/lang/core/ir/store_vector_element.cc
@@ -14,6 +14,9 @@
#include "src/tint/lang/core/ir/store_vector_element.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::StoreVectorElement);
namespace tint::core::ir {
@@ -28,4 +31,11 @@
StoreVectorElement::~StoreVectorElement() = default;
+StoreVectorElement* StoreVectorElement::Clone(CloneContext& ctx) {
+ auto* new_to = ctx.Clone(To());
+ auto* new_idx = ctx.Clone(Index());
+ auto* new_val = ctx.Clone(Value());
+ return ctx.ir.instructions.Create<StoreVectorElement>(new_to, new_idx, new_val);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/store_vector_element.h b/src/tint/lang/core/ir/store_vector_element.h
index a4b31a8..5f89adb 100644
--- a/src/tint/lang/core/ir/store_vector_element.h
+++ b/src/tint/lang/core/ir/store_vector_element.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A store instruction for a single vector element in the IR.
-class StoreVectorElement : public Castable<StoreVectorElement, OperandInstruction<3, 0>> {
+class StoreVectorElement final : public Castable<StoreVectorElement, OperandInstruction<3, 0>> {
public:
/// The offset in Operands() for the `to` value
static constexpr size_t kToOperandOffset = 0;
@@ -41,6 +41,9 @@
StoreVectorElement(ir::Value* to, ir::Value* index, ir::Value* value);
~StoreVectorElement() override;
+ /// @copydoc Instruction::Clone()
+ StoreVectorElement* Clone(CloneContext& ctx) override;
+
/// @returns the vector pointer value
ir::Value* To() { return operands_[kToOperandOffset]; }
diff --git a/src/tint/lang/core/ir/store_vector_element_test.cc b/src/tint/lang/core/ir/store_vector_element_test.cc
index 03e147e..28f417e 100644
--- a/src/tint/lang/core/ir/store_vector_element_test.cc
+++ b/src/tint/lang/core/ir/store_vector_element_test.cc
@@ -66,5 +66,24 @@
EXPECT_FALSE(inst->HasMultiResults());
}
+TEST_F(IR_StoreVectorElementTest, Clone) {
+ auto* to = b.Var(ty.ptr<private_, vec3<i32>>());
+ auto* inst = b.StoreVectorElement(to, 2_i, 4_i);
+
+ auto* new_to = clone_ctx.Clone(to);
+ auto* new_inst = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, new_inst);
+ EXPECT_EQ(new_to->Result(), new_inst->To());
+
+ auto new_idx = new_inst->Index()->As<Constant>()->Value();
+ ASSERT_TRUE(new_idx->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(2_i, new_idx->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+
+ auto new_val = new_inst->Value()->As<Constant>()->Value();
+ ASSERT_TRUE(new_val->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(4_i, new_val->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/switch.cc b/src/tint/lang/core/ir/switch.cc
index df353ab..d260921 100644
--- a/src/tint/lang/core/ir/switch.cc
+++ b/src/tint/lang/core/ir/switch.cc
@@ -14,6 +14,8 @@
#include "src/tint/lang/core/ir/switch.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Switch);
@@ -34,4 +36,25 @@
}
}
+Switch* Switch::Clone(CloneContext& ctx) {
+ auto* new_cond = ctx.Clone(Condition());
+ auto* new_switch = ctx.ir.instructions.Create<Switch>(new_cond);
+ ctx.Replace(this, new_switch);
+
+ new_switch->cases_.Reserve(cases_.Length());
+ for (const auto& cse : cases_) {
+ Switch::Case new_case{};
+ new_case.block = ctx.ir.blocks.Create<ir::Block>();
+ cse.block->CloneInto(ctx, new_case.block);
+
+ new_case.selectors.Reserve(cse.selectors.Length());
+ for (const auto& sel : cse.selectors) {
+ auto* new_val = sel.val ? ctx.Clone(sel.val) : nullptr;
+ new_case.selectors.Push(Switch::CaseSelector{new_val});
+ }
+ new_switch->cases_.Push(new_case);
+ }
+ return new_switch;
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/switch.h b/src/tint/lang/core/ir/switch.h
index 1b450ac..44f0e6a 100644
--- a/src/tint/lang/core/ir/switch.h
+++ b/src/tint/lang/core/ir/switch.h
@@ -43,7 +43,7 @@
/// ▼
/// out
/// ```
-class Switch : public Castable<Switch, ControlInstruction> {
+class Switch final : public Castable<Switch, ControlInstruction> {
public:
/// The offset in Operands() for the condition
static constexpr size_t kConditionOperandOffset = 0;
@@ -73,6 +73,9 @@
explicit Switch(Value* cond);
~Switch() override;
+ /// @copydoc Instruction::Clone()
+ Switch* Clone(CloneContext& ctx) override;
+
/// @copydoc ControlInstruction::ForeachBlock
void ForeachBlock(const std::function<void(ir::Block*)>& cb) override;
diff --git a/src/tint/lang/core/ir/switch_test.cc b/src/tint/lang/core/ir/switch_test.cc
index 5d1442c..83cc223 100644
--- a/src/tint/lang/core/ir/switch_test.cc
+++ b/src/tint/lang/core/ir/switch_test.cc
@@ -22,6 +22,7 @@
namespace {
using namespace tint::core::number_suffixes; // NOLINT
+
using IR_SwitchTest = IRTestHelper;
TEST_F(IR_SwitchTest, Usage) {
@@ -43,5 +44,62 @@
EXPECT_THAT(switch_->Cases().Front().Block()->Parent(), switch_);
}
+TEST_F(IR_SwitchTest, Clone) {
+ auto* switch_ = b.Switch(1_i);
+ switch_->Cases().Push(
+ Switch::Case{{Switch::CaseSelector{}, Switch::CaseSelector{b.Constant(2_i)}}, b.Block()});
+ switch_->Cases().Push(Switch::Case{{Switch::CaseSelector{b.Constant(3_i)}}, b.Block()});
+
+ auto* new_switch = clone_ctx.Clone(switch_);
+
+ EXPECT_NE(switch_, new_switch);
+
+ auto new_cond = new_switch->Condition()->As<Constant>()->Value();
+ ASSERT_TRUE(new_cond->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(1_i, new_cond->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+
+ auto& cases = new_switch->Cases();
+ ASSERT_EQ(2u, cases.Length());
+
+ {
+ auto& case1 = cases[0];
+ EXPECT_NE(nullptr, case1.block);
+ EXPECT_NE(switch_->Cases()[0].block, case1.block);
+
+ ASSERT_EQ(2u, case1.selectors.Length());
+ EXPECT_EQ(nullptr, case1.selectors[0].val);
+ auto val = case1.selectors[1].val->Value();
+ ASSERT_TRUE(val->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(2_i, val->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+ }
+
+ {
+ auto& case2 = cases[1];
+ EXPECT_NE(nullptr, case2.block);
+ EXPECT_NE(switch_->Cases()[1].block, case2.block);
+
+ ASSERT_EQ(1u, case2.selectors.Length());
+ auto val = case2.selectors[0].val->Value();
+ ASSERT_TRUE(val->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(3_i, val->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+ }
+}
+
+TEST_F(IR_SwitchTest, CloneWithExits) {
+ Switch* new_switch = nullptr;
+ {
+ auto* switch_ = b.Switch(1_i);
+
+ auto* blk = b.Block();
+ b.Append(blk, [&] { b.ExitSwitch(switch_); });
+ switch_->Cases().Push(Switch::Case{{Switch::CaseSelector{b.Constant(3_i)}}, blk});
+ new_switch = clone_ctx.Clone(switch_);
+ }
+
+ auto& case_ = new_switch->Cases().Front();
+ ASSERT_TRUE(case_.block->Front()->Is<ExitSwitch>());
+ EXPECT_EQ(new_switch, case_.block->Front()->As<ExitSwitch>()->Switch());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/swizzle.cc b/src/tint/lang/core/ir/swizzle.cc
index 3789021..0180d48 100644
--- a/src/tint/lang/core/ir/swizzle.cc
+++ b/src/tint/lang/core/ir/swizzle.cc
@@ -16,6 +16,8 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/utils/ice/ice.h"
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Swizzle);
@@ -37,4 +39,10 @@
Swizzle::~Swizzle() = default;
+Swizzle* Swizzle::Clone(CloneContext& ctx) {
+ auto* result = ctx.Clone(Result());
+ auto* new_obj = ctx.Clone(Object());
+ return ctx.ir.instructions.Create<Swizzle>(result, new_obj, indices_);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/swizzle.h b/src/tint/lang/core/ir/swizzle.h
index 48a838e..880e5c6 100644
--- a/src/tint/lang/core/ir/swizzle.h
+++ b/src/tint/lang/core/ir/swizzle.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A swizzle instruction in the IR.
-class Swizzle : public Castable<Swizzle, OperandInstruction<1, 1>> {
+class Swizzle final : public Castable<Swizzle, OperandInstruction<1, 1>> {
public:
/// The offset in Operands() for the object being swizzled
static constexpr size_t kObjectOperandOffset = 0;
@@ -35,6 +35,9 @@
Swizzle(InstructionResult* result, Value* object, VectorRef<uint32_t> indices);
~Swizzle() override;
+ /// @copydoc Instruction::Clone()
+ Swizzle* Clone(CloneContext& ctx) override;
+
/// @returns the object used for the access
Value* Object() { return operands_[kObjectOperandOffset]; }
diff --git a/src/tint/lang/core/ir/swizzle_test.cc b/src/tint/lang/core/ir/swizzle_test.cc
index 3f435ef..09eb7e9 100644
--- a/src/tint/lang/core/ir/swizzle_test.cc
+++ b/src/tint/lang/core/ir/swizzle_test.cc
@@ -86,5 +86,22 @@
"");
}
+TEST_F(IR_SwizzleTest, Clone) {
+ auto* var = b.Var(ty.ptr<function, i32>());
+ auto* s = b.Swizzle(mod.Types().i32(), var, {2u});
+
+ auto* new_var = clone_ctx.Clone(var);
+ auto* new_s = clone_ctx.Clone(s);
+
+ EXPECT_NE(s, new_s);
+ EXPECT_NE(nullptr, new_s->Result());
+ EXPECT_NE(s->Result(), new_s->Result());
+
+ EXPECT_EQ(new_var->Result(), new_s->Object());
+
+ EXPECT_EQ(1u, new_s->Indices().Length());
+ EXPECT_EQ(2u, new_s->Indices().Front());
+}
+
} // namespace
} // 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 a7ab28e..9fa5b54 100644
--- a/src/tint/lang/core/ir/terminate_invocation.cc
+++ b/src/tint/lang/core/ir/terminate_invocation.cc
@@ -14,10 +14,17 @@
#include "src/tint/lang/core/ir/terminate_invocation.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::TerminateInvocation);
namespace tint::core::ir {
TerminateInvocation::~TerminateInvocation() = default;
+TerminateInvocation* TerminateInvocation::Clone(CloneContext& ctx) {
+ return ctx.ir.instructions.Create<TerminateInvocation>();
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/terminate_invocation.h b/src/tint/lang/core/ir/terminate_invocation.h
index 8edf03e..097145c 100644
--- a/src/tint/lang/core/ir/terminate_invocation.h
+++ b/src/tint/lang/core/ir/terminate_invocation.h
@@ -22,10 +22,13 @@
namespace tint::core::ir {
/// An terminate invocation instruction in the IR.
-class TerminateInvocation : public Castable<TerminateInvocation, Terminator> {
+class TerminateInvocation final : public Castable<TerminateInvocation, Terminator> {
public:
~TerminateInvocation() override;
+ /// @copydoc Instruction::Clone()
+ TerminateInvocation* Clone(CloneContext& ctx) override;
+
/// @returns the friendly name for the instruction
std::string FriendlyName() override { return "terminate_invocation"; }
};
diff --git a/src/tint/lang/core/ir/terminate_invocation_test.cc b/src/tint/lang/core/ir/terminate_invocation_test.cc
new file mode 100644
index 0000000..459aae0
--- /dev/null
+++ b/src/tint/lang/core/ir/terminate_invocation_test.cc
@@ -0,0 +1,34 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "src/tint/lang/core/ir/terminate_invocation.h"
+#include "gtest/gtest.h"
+#include "src/tint/lang/core/ir/ir_helper_test.h"
+
+namespace tint::core::ir {
+namespace {
+
+using IR_TerminateInvocationTest = IRTestHelper;
+
+TEST_F(IR_TerminateInvocationTest, Clone) {
+ auto* ti = b.TerminateInvocation();
+ auto* new_ti = clone_ctx.Clone(ti);
+
+ EXPECT_NE(ti, new_ti);
+ EXPECT_NE(nullptr, new_ti);
+ EXPECT_NE(ti, new_ti);
+}
+
+} // namespace
+} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/unary.cc b/src/tint/lang/core/ir/unary.cc
index d6ce129..69b6857 100644
--- a/src/tint/lang/core/ir/unary.cc
+++ b/src/tint/lang/core/ir/unary.cc
@@ -14,6 +14,9 @@
#include "src/tint/lang/core/ir/unary.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Unary);
namespace tint::core::ir {
@@ -25,4 +28,10 @@
Unary::~Unary() = default;
+Unary* Unary::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_val = ctx.Clone(Val());
+ return ctx.ir.instructions.Create<Unary>(new_result, kind_, new_val);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/unary.h b/src/tint/lang/core/ir/unary.h
index 2e55cc3..15fdfa3 100644
--- a/src/tint/lang/core/ir/unary.h
+++ b/src/tint/lang/core/ir/unary.h
@@ -23,7 +23,7 @@
namespace tint::core::ir {
/// A unary instruction in the IR.
-class Unary : public Castable<Unary, OperandInstruction<1, 1>> {
+class Unary final : public Castable<Unary, OperandInstruction<1, 1>> {
public:
/// The offset in Operands() for the value
static constexpr size_t kValueOperandOffset = 0;
@@ -41,6 +41,9 @@
Unary(InstructionResult* result, enum Kind kind, Value* val);
~Unary() override;
+ /// @copydoc Instruction::Clone()
+ Unary* Clone(CloneContext& ctx) override;
+
/// @returns the value for the instruction
Value* Val() { return operands_[kValueOperandOffset]; }
diff --git a/src/tint/lang/core/ir/unary_test.cc b/src/tint/lang/core/ir/unary_test.cc
index 95d30cf..ab90266 100644
--- a/src/tint/lang/core/ir/unary_test.cc
+++ b/src/tint/lang/core/ir/unary_test.cc
@@ -78,5 +78,20 @@
"");
}
+TEST_F(IR_UnaryTest, Clone) {
+ auto* inst = b.Complement(mod.Types().i32(), 4_i);
+ auto* new_inst = clone_ctx.Clone(inst);
+
+ EXPECT_NE(inst, new_inst);
+ EXPECT_NE(nullptr, new_inst->Result());
+ EXPECT_NE(inst->Result(), new_inst->Result());
+
+ EXPECT_EQ(Unary::Kind::kComplement, new_inst->Kind());
+
+ auto new_val = new_inst->Val()->As<Constant>()->Value();
+ ASSERT_TRUE(new_val->Is<core::constant::Scalar<i32>>());
+ EXPECT_EQ(4_i, new_val->As<core::constant::Scalar<i32>>()->ValueAs<i32>());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/unreachable.cc b/src/tint/lang/core/ir/unreachable.cc
index 0011530..da73d70 100644
--- a/src/tint/lang/core/ir/unreachable.cc
+++ b/src/tint/lang/core/ir/unreachable.cc
@@ -14,10 +14,17 @@
#include "src/tint/lang/core/ir/unreachable.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::Unreachable);
namespace tint::core::ir {
Unreachable::~Unreachable() = default;
+Unreachable* Unreachable::Clone(CloneContext& ctx) {
+ return ctx.ir.instructions.Create<Unreachable>();
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/unreachable.h b/src/tint/lang/core/ir/unreachable.h
index fd92538..9a9d264 100644
--- a/src/tint/lang/core/ir/unreachable.h
+++ b/src/tint/lang/core/ir/unreachable.h
@@ -22,10 +22,13 @@
namespace tint::core::ir {
/// An unreachable instruction in the IR.
-class Unreachable : public Castable<Unreachable, Terminator> {
+class Unreachable final : public Castable<Unreachable, Terminator> {
public:
~Unreachable() override;
+ /// @copydoc Instruction::Clone()
+ Unreachable* Clone(CloneContext& ctx) override;
+
/// @returns the friendly name for the instruction
std::string FriendlyName() override { return "unreachable"; }
};
diff --git a/src/tint/lang/core/ir/unreachable_test.cc b/src/tint/lang/core/ir/unreachable_test.cc
new file mode 100644
index 0000000..5e172bd
--- /dev/null
+++ b/src/tint/lang/core/ir/unreachable_test.cc
@@ -0,0 +1,46 @@
+// Copyright 2023 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include "gtest/gtest-spi.h"
+#include "src/tint/lang/core/ir/builder.h"
+#include "src/tint/lang/core/ir/instruction.h"
+#include "src/tint/lang/core/ir/ir_helper_test.h"
+
+namespace tint::core::ir {
+namespace {
+
+using IR_UnreachableTest = IRTestHelper;
+
+TEST_F(IR_UnreachableTest, Unreachable) {
+ auto* inst = b.Unreachable();
+ ASSERT_TRUE(inst->Is<ir::Unreachable>());
+}
+
+TEST_F(IR_UnreachableTest, Result) {
+ auto* inst = b.Unreachable();
+
+ EXPECT_FALSE(inst->HasResults());
+ EXPECT_FALSE(inst->HasMultiResults());
+}
+
+TEST_F(IR_UnreachableTest, Clone) {
+ auto* d = b.Unreachable();
+ auto* new_d = clone_ctx.Clone(d);
+
+ EXPECT_NE(d, new_d);
+ EXPECT_NE(nullptr, new_d);
+}
+
+} // namespace
+} // 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 1a41472..0f20129 100644
--- a/src/tint/lang/core/ir/user_call.cc
+++ b/src/tint/lang/core/ir/user_call.cc
@@ -16,6 +16,9 @@
#include <utility>
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
+
TINT_INSTANTIATE_TYPEINFO(tint::core::ir::UserCall);
namespace tint::core::ir {
@@ -28,4 +31,11 @@
UserCall::~UserCall() = default;
+UserCall* UserCall::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_target = ctx.Clone(Target());
+ auto new_args = ctx.Clone<UserCall::kDefaultNumOperands>(Args());
+ return ctx.ir.instructions.Create<UserCall>(new_result, new_target, new_args);
+}
+
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/user_call.h b/src/tint/lang/core/ir/user_call.h
index f6155ab..f6688e2 100644
--- a/src/tint/lang/core/ir/user_call.h
+++ b/src/tint/lang/core/ir/user_call.h
@@ -24,7 +24,7 @@
namespace tint::core::ir {
/// A user call instruction in the IR.
-class UserCall : public Castable<UserCall, Call> {
+class UserCall final : public Castable<UserCall, Call> {
public:
/// The offset in Operands() for the function being called
static constexpr size_t kFunctionOperandOffset = 0;
@@ -39,6 +39,9 @@
UserCall(InstructionResult* result, Function* func, VectorRef<Value*> args);
~UserCall() override;
+ /// @copydoc Instruction::Clone()
+ UserCall* Clone(CloneContext& ctx) override;
+
/// @returns the call arguments
tint::Slice<Value*> Args() override { return operands_.Slice().Offset(kArgsOperandOffset); }
diff --git a/src/tint/lang/core/ir/user_call_test.cc b/src/tint/lang/core/ir/user_call_test.cc
index 51f99eb..efbe64c 100644
--- a/src/tint/lang/core/ir/user_call_test.cc
+++ b/src/tint/lang/core/ir/user_call_test.cc
@@ -56,5 +56,39 @@
"");
}
+TEST_F(IR_UserCallTest, Clone) {
+ auto* func = b.Function("myfunc", mod.Types().void_());
+ auto* e = b.Call(mod.Types().void_(), func, Vector{b.Constant(1_u), b.Constant(2_u)});
+
+ auto* new_func = clone_ctx.Clone(func);
+ auto* new_e = clone_ctx.Clone(e);
+
+ EXPECT_NE(e, new_e);
+ EXPECT_NE(nullptr, new_e->Result());
+ EXPECT_NE(e->Result(), new_e->Result());
+
+ EXPECT_EQ(new_func, new_e->Target());
+
+ auto args = new_e->Args();
+ EXPECT_EQ(2u, args.Length());
+
+ auto new_arg1 = args[0]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg1->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(1_u, new_arg1->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+
+ auto new_arg2 = args[1]->As<Constant>()->Value();
+ ASSERT_TRUE(new_arg2->Is<core::constant::Scalar<u32>>());
+ EXPECT_EQ(2_u, new_arg2->As<core::constant::Scalar<u32>>()->ValueAs<u32>());
+}
+
+TEST_F(IR_UserCallTest, CloneWithoutArgs) {
+ auto* func = b.Function("myfunc", mod.Types().void_());
+ auto* e = b.Call(mod.Types().void_(), func);
+
+ auto* new_e = clone_ctx.Clone(e);
+
+ EXPECT_EQ(0u, new_e->Args().Length());
+}
+
} // namespace
} // namespace tint::core::ir
diff --git a/src/tint/lang/core/ir/value.h b/src/tint/lang/core/ir/value.h
index 6393bc9..176ad0e 100644
--- a/src/tint/lang/core/ir/value.h
+++ b/src/tint/lang/core/ir/value.h
@@ -21,6 +21,7 @@
// Forward declarations
namespace tint::core::ir {
+class CloneContext;
class Instruction;
} // namespace tint::core::ir
@@ -63,6 +64,10 @@
/// The Value must not be in use by any instruction.
virtual void Destroy();
+ /// @param ctx the CloneContext used to clone this value
+ /// @returns a clone of this value
+ virtual Value* Clone(CloneContext& ctx) = 0;
+
/// @returns true if the Value has not been destroyed with Destroy()
bool Alive() const { return !flags_.Contains(Flag::kDead); }
diff --git a/src/tint/lang/core/ir/var.cc b/src/tint/lang/core/ir/var.cc
index 501d6af..3ea21f9 100644
--- a/src/tint/lang/core/ir/var.cc
+++ b/src/tint/lang/core/ir/var.cc
@@ -14,6 +14,8 @@
#include "src/tint/lang/core/ir/var.h"
+#include "src/tint/lang/core/ir/clone_context.h"
+#include "src/tint/lang/core/ir/module.h"
#include "src/tint/lang/core/ir/store.h"
#include "src/tint/lang/core/type/pointer.h"
#include "src/tint/utils/ice/ice.h"
@@ -34,6 +36,20 @@
Var::~Var() = default;
+Var* Var::Clone(CloneContext& ctx) {
+ auto* new_result = ctx.Clone(Result());
+ auto* new_var = ctx.ir.instructions.Create<Var>(new_result);
+
+ new_var->binding_point_ = binding_point_;
+ new_var->attributes_ = attributes_;
+
+ auto name = ctx.ir.NameOf(this);
+ if (name.IsValid()) {
+ ctx.ir.SetName(new_var, name.Name());
+ }
+ return new_var;
+}
+
void Var::SetInitializer(Value* initializer) {
SetOperand(Var::kInitializerOperandOffset, initializer);
}
diff --git a/src/tint/lang/core/ir/var.h b/src/tint/lang/core/ir/var.h
index 384f7b1..01e8c38 100644
--- a/src/tint/lang/core/ir/var.h
+++ b/src/tint/lang/core/ir/var.h
@@ -40,7 +40,7 @@
};
/// A var instruction in the IR.
-class Var : public Castable<Var, OperandInstruction<1, 1>> {
+class Var final : public Castable<Var, OperandInstruction<1, 1>> {
public:
/// The offset in Operands() for the initializer
static constexpr size_t kInitializerOperandOffset = 0;
@@ -50,6 +50,9 @@
explicit Var(InstructionResult* result);
~Var() override;
+ /// @copydoc Instruction::Clone()
+ Var* Clone(CloneContext& ctx) override;
+
/// Sets the var initializer
/// @param initializer the initializer
void SetInitializer(Value* initializer);
diff --git a/src/tint/lang/core/ir/var_test.cc b/src/tint/lang/core/ir/var_test.cc
index 24fb70e..c147dc3 100644
--- a/src/tint/lang/core/ir/var_test.cc
+++ b/src/tint/lang/core/ir/var_test.cc
@@ -58,5 +58,54 @@
EXPECT_TRUE(init->Usages().IsEmpty());
}
+TEST_F(IR_VarTest, Clone) {
+ auto* v = b.Var(mod.Types().ptr(core::AddressSpace::kFunction, mod.Types().f32()));
+ v->SetInitializer(b.Constant(4_f));
+ v->SetBindingPoint(1, 2);
+ v->SetAttributes(IOAttributes{
+ 3, 4, core::BuiltinValue::kFragDepth,
+ Interpolation{core::InterpolationType::kFlat, core::InterpolationSampling::kCentroid},
+ true});
+
+ auto* new_v = clone_ctx.Clone(v);
+
+ EXPECT_NE(v, new_v);
+ ASSERT_NE(nullptr, new_v->Result());
+ EXPECT_NE(v->Result(), new_v->Result());
+ EXPECT_EQ(new_v->Result()->Type(),
+ mod.Types().ptr(core::AddressSpace::kFunction, mod.Types().f32()));
+
+ auto new_val = v->Initializer()->As<Constant>()->Value();
+ ASSERT_TRUE(new_val->Is<core::constant::Scalar<f32>>());
+ EXPECT_FLOAT_EQ(4_f, new_val->As<core::constant::Scalar<f32>>()->ValueAs<f32>());
+
+ EXPECT_TRUE(new_v->BindingPoint().has_value());
+ EXPECT_EQ(1u, new_v->BindingPoint()->group);
+ EXPECT_EQ(2u, new_v->BindingPoint()->binding);
+
+ auto& attrs = new_v->Attributes();
+ EXPECT_TRUE(attrs.location.has_value());
+ EXPECT_EQ(3u, attrs.location.value());
+
+ EXPECT_TRUE(attrs.index.has_value());
+ EXPECT_EQ(4u, attrs.index.value());
+
+ EXPECT_TRUE(attrs.builtin.has_value());
+ EXPECT_EQ(core::BuiltinValue::kFragDepth, attrs.builtin.value());
+
+ EXPECT_TRUE(attrs.interpolation.has_value());
+ EXPECT_EQ(core::InterpolationType::kFlat, attrs.interpolation->type);
+ EXPECT_EQ(core::InterpolationSampling::kCentroid, attrs.interpolation->sampling);
+
+ EXPECT_TRUE(attrs.invariant);
+}
+
+TEST_F(IR_VarTest, CloneWithName) {
+ auto* v = b.Var("v", mod.Types().ptr(core::AddressSpace::kFunction, mod.Types().f32()));
+ auto* new_v = clone_ctx.Clone(v);
+
+ EXPECT_EQ(std::string("v"), mod.NameOf(new_v).Name());
+}
+
} // namespace
} // namespace tint::core::ir
