Import Tint changes from Dawn
Changes:
- 24cb81116d30d0eee37ed5897c5a65df4baf31a0 [ir] Shift back to accessors. by dan sinclair <dsinclair@chromium.org>
- 97744832bc6e24eead7416d26f35c6ee9320981a [ir] Deduplicate constants by dan sinclair <dsinclair@chromium.org>
- 92151b238b7a3c71ae5e27f9eaa769b618b3f6ac [ir][spirv-writer] Emit load instructions by James Price <jrprice@google.com>
- 84d750e9823a70aee01d40b19eb0d4a92ecaa40d [ir] Add function parameters. by dan sinclair <dsinclair@chromium.org>
GitOrigin-RevId: 24cb81116d30d0eee37ed5897c5a65df4baf31a0
Change-Id: I7607901715e20e152804c5ef21101760636346c3
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/133680
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 47efd52..64fde4c 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -1226,6 +1226,8 @@
"ir/flow_node.h",
"ir/function.cc",
"ir/function.h",
+ "ir/function_param.cc",
+ "ir/function_param.h",
"ir/function_terminator.cc",
"ir/function_terminator.h",
"ir/if.cc",
@@ -2263,6 +2265,7 @@
"ir/store_test.cc",
"ir/test_helper.h",
"ir/to_program_roundtrip_test.cc",
+ "ir/transform/add_empty_entry_point_test.cc",
"ir/unary_test.cc",
]
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index d0c7487..8ec573a 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -742,6 +742,8 @@
ir/flow_node.h
ir/function.cc
ir/function.h
+ ir/function_param.cc
+ ir/function_param.h
ir/function_terminator.cc
ir/function_terminator.h
ir/if.cc
@@ -1474,6 +1476,7 @@
ir/module_test.cc
ir/store_test.cc
ir/test_helper.h
+ ir/transform/add_empty_entry_point_test.cc
ir/unary_test.cc
)
endif()
diff --git a/src/tint/ir/binary.cc b/src/tint/ir/binary.cc
index af28306..2b179ac 100644
--- a/src/tint/ir/binary.cc
+++ b/src/tint/ir/binary.cc
@@ -19,8 +19,8 @@
namespace tint::ir {
-Binary::Binary(Kind k, const type::Type* res_ty, Value* lhs, Value* rhs)
- : kind(k), result_type(res_ty), lhs_(lhs), rhs_(rhs) {
+Binary::Binary(enum Kind kind, const type::Type* res_ty, Value* lhs, Value* rhs)
+ : kind_(kind), result_type_(res_ty), lhs_(lhs), rhs_(rhs) {
TINT_ASSERT(IR, lhs);
TINT_ASSERT(IR, rhs);
lhs_->AddUsage(this);
diff --git a/src/tint/ir/binary.h b/src/tint/ir/binary.h
index a6dc722..7ea6c08 100644
--- a/src/tint/ir/binary.h
+++ b/src/tint/ir/binary.h
@@ -51,7 +51,7 @@
/// @param type the result type
/// @param lhs the lhs of the instruction
/// @param rhs the rhs of the instruction
- Binary(Kind kind, const type::Type* type, Value* lhs, Value* rhs);
+ Binary(enum Kind kind, const type::Type* type, Value* lhs, Value* rhs);
Binary(const Binary& inst) = delete;
Binary(Binary&& inst) = delete;
~Binary() override;
@@ -59,8 +59,11 @@
Binary& operator=(const Binary& inst) = delete;
Binary& operator=(Binary&& inst) = delete;
+ /// @returns the kind of the binary instruction
+ enum Kind Kind() const { return kind_; }
+
/// @returns the type of the value
- const type::Type* Type() const override { return result_type; }
+ const type::Type* Type() const override { return result_type_; }
/// @returns the left-hand-side value for the instruction
const Value* LHS() const { return lhs_; }
@@ -68,15 +71,11 @@
/// @returns the right-hand-side value for the instruction
const Value* RHS() const { return rhs_; }
- /// the kind of binary instruction
- Kind kind = Kind::kAdd;
-
- /// the result type of the instruction
- const type::Type* result_type = nullptr;
-
private:
- Value* lhs_ = nullptr;
- Value* rhs_ = nullptr;
+ enum Kind kind_;
+ const type::Type* result_type_;
+ Value* lhs_;
+ Value* rhs_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/binary_test.cc b/src/tint/ir/binary_test.cc
index 2968d28..dc35978 100644
--- a/src/tint/ir/binary_test.cc
+++ b/src/tint/ir/binary_test.cc
@@ -30,17 +30,16 @@
const auto* inst = b.And(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
- ASSERT_NE(inst->result_type, nullptr);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kAnd);
ASSERT_NE(inst->Type(), nullptr);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -52,15 +51,15 @@
const auto* inst = b.Or(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kOr);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kOr);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -72,15 +71,15 @@
const auto* inst = b.Xor(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kXor);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kXor);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -92,15 +91,15 @@
const auto* inst = b.Equal(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kEqual);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -112,15 +111,15 @@
const auto* inst = b.NotEqual(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kNotEqual);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kNotEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -132,15 +131,15 @@
const auto* inst = b.LessThan(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kLessThan);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kLessThan);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -153,15 +152,15 @@
b.GreaterThan(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThan);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kGreaterThan);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -174,15 +173,15 @@
b.LessThanEqual(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kLessThanEqual);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kLessThanEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -195,15 +194,15 @@
b.GreaterThanEqual(b.ir.types.Get<type::Bool>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThanEqual);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kGreaterThanEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -214,15 +213,15 @@
const auto* inst = b.Not(b.ir.types.Get<type::Bool>(), b.Constant(true));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kEqual);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kEqual);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<bool>>());
EXPECT_TRUE(lhs->As<constant::Scalar<bool>>()->ValueAs<bool>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<bool>>());
EXPECT_FALSE(rhs->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
@@ -234,15 +233,15 @@
const auto* inst = b.ShiftLeft(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kShiftLeft);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kShiftLeft);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -254,15 +253,15 @@
const auto* inst = b.ShiftRight(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kShiftRight);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kShiftRight);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -274,15 +273,15 @@
const auto* inst = b.Add(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kAdd);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kAdd);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -294,15 +293,15 @@
const auto* inst = b.Subtract(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kSubtract);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kSubtract);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -314,15 +313,15 @@
const auto* inst = b.Multiply(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kMultiply);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kMultiply);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -334,15 +333,15 @@
const auto* inst = b.Divide(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kDivide);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kDivide);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -354,15 +353,15 @@
const auto* inst = b.Modulo(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
ASSERT_TRUE(inst->Is<Binary>());
- EXPECT_EQ(inst->kind, Binary::Kind::kModulo);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kModulo);
ASSERT_TRUE(inst->LHS()->Is<Constant>());
- auto lhs = inst->LHS()->As<Constant>()->value;
+ auto lhs = inst->LHS()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
ASSERT_TRUE(inst->RHS()->Is<Constant>());
- auto rhs = inst->RHS()->As<Constant>()->value;
+ auto rhs = inst->RHS()->As<Constant>()->Value();
ASSERT_TRUE(rhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(2_i, rhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -372,7 +371,7 @@
Builder b{mod};
const auto* inst = b.And(b.ir.types.Get<type::I32>(), b.Constant(4_i), b.Constant(2_i));
- EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kAnd);
ASSERT_NE(inst->LHS(), nullptr);
ASSERT_EQ(inst->LHS()->Usage().Length(), 1u);
@@ -389,7 +388,7 @@
auto val = b.Constant(4_i);
const auto* inst = b.And(b.ir.types.Get<type::I32>(), val, val);
- EXPECT_EQ(inst->kind, Binary::Kind::kAnd);
+ EXPECT_EQ(inst->Kind(), Binary::Kind::kAnd);
ASSERT_EQ(inst->LHS(), inst->RHS());
ASSERT_NE(inst->LHS(), nullptr);
diff --git a/src/tint/ir/bitcast_test.cc b/src/tint/ir/bitcast_test.cc
index e78b734..6eda562 100644
--- a/src/tint/ir/bitcast_test.cc
+++ b/src/tint/ir/bitcast_test.cc
@@ -32,9 +32,10 @@
ASSERT_TRUE(inst->Is<ir::Bitcast>());
ASSERT_NE(inst->Type(), nullptr);
- ASSERT_EQ(inst->args.Length(), 1u);
- ASSERT_TRUE(inst->args[0]->Is<Constant>());
- auto val = inst->args[0]->As<Constant>()->value;
+ const auto args = inst->Args();
+ ASSERT_EQ(args.Length(), 1u);
+ ASSERT_TRUE(args[0]->Is<Constant>());
+ auto val = args[0]->As<Constant>()->Value();
ASSERT_TRUE(val->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, val->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -44,10 +45,11 @@
Builder b{mod};
const auto* inst = b.Bitcast(b.ir.types.Get<type::I32>(), b.Constant(4_i));
- ASSERT_EQ(inst->args.Length(), 1u);
- ASSERT_NE(inst->args[0], nullptr);
- ASSERT_EQ(inst->args[0]->Usage().Length(), 1u);
- EXPECT_EQ(inst->args[0]->Usage()[0], inst);
+ const auto args = inst->Args();
+ ASSERT_EQ(args.Length(), 1u);
+ ASSERT_NE(args[0], nullptr);
+ ASSERT_EQ(args[0]->Usage().Length(), 1u);
+ EXPECT_EQ(args[0]->Usage()[0], inst);
}
} // namespace
diff --git a/src/tint/ir/block.cc b/src/tint/ir/block.cc
index 2030c19..f5b5bfb 100644
--- a/src/tint/ir/block.cc
+++ b/src/tint/ir/block.cc
@@ -22,4 +22,13 @@
Block::~Block() = default;
+void Block::BranchTo(FlowNode* to, utils::VectorRef<Value*> args) {
+ TINT_ASSERT(IR, to);
+ branch_.target = to;
+ branch_.args = args;
+ if (to) {
+ to->AddInboundBranch(this);
+ }
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h
index 5af32fe..2b01fd0 100644
--- a/src/tint/ir/block.h
+++ b/src/tint/ir/block.h
@@ -15,6 +15,8 @@
#ifndef SRC_TINT_IR_BLOCK_H_
#define SRC_TINT_IR_BLOCK_H_
+#include <utility>
+
#include "src/tint/ir/block_param.h"
#include "src/tint/ir/branch.h"
#include "src/tint/ir/flow_node.h"
@@ -30,20 +32,48 @@
public:
/// Constructor
Block();
+ Block(const Block&) = delete;
+ Block(Block&&) = delete;
~Block() override;
- /// @returns true if this is a dead block. This can happen in the case like a loop merge block
- /// which is never reached.
- bool IsDead() const override { return branch.target == nullptr; }
+ Block& operator=(const Block&) = delete;
+ Block& operator=(Block&&) = delete;
- /// The node this block branches too.
- Branch branch = {};
+ /// Sets the blocks branch target to the given node.
+ /// @param to the node to branch too
+ /// @param args the branch arguments
+ void BranchTo(FlowNode* to, utils::VectorRef<Value*> args = {});
- /// The instructions in the block
- utils::Vector<const Instruction*, 16> instructions;
+ /// @returns true if this is block has a branch target set
+ bool HasBranchTarget() const override { return branch_.target != nullptr; }
- /// The parameters passed into the block
- utils::Vector<const BlockParam*, 0> params;
+ /// @return the node this block branches too.
+ const ir::Branch& Branch() const { return branch_; }
+
+ /// Sets the instructions in the block
+ /// @param instructions the instructions to set
+ void SetInstructions(utils::VectorRef<const Instruction*> instructions) {
+ instructions_ = std::move(instructions);
+ }
+
+ /// @returns the instructions in the block
+ utils::VectorRef<const Instruction*> Instructions() const { return instructions_; }
+ /// @returns the instructions in the block
+ utils::Vector<const Instruction*, 16>& Instructions() { return instructions_; }
+
+ /// Sets the params to the block
+ /// @param params the params for the block
+ void SetParams(utils::VectorRef<const BlockParam*> params) { params_ = std::move(params); }
+ /// @returns the params to the block
+ utils::Vector<const BlockParam*, 0>& Params() { return params_; }
+
+ /// @return the parameters passed into the block
+ utils::VectorRef<const BlockParam*> Params() const { return params_; }
+
+ private:
+ ir::Branch branch_ = {};
+ utils::Vector<const Instruction*, 16> instructions_;
+ utils::Vector<const BlockParam*, 0> params_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/block_param.cc b/src/tint/ir/block_param.cc
index f014d19..a3a0be5 100644
--- a/src/tint/ir/block_param.cc
+++ b/src/tint/ir/block_param.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-BlockParam::BlockParam(const type::Type* ty) : type(ty) {}
+BlockParam::BlockParam(const type::Type* ty) : type_(ty) {}
BlockParam::~BlockParam() = default;
diff --git a/src/tint/ir/block_param.h b/src/tint/ir/block_param.h
index 8ba68a7..036ddbf 100644
--- a/src/tint/ir/block_param.h
+++ b/src/tint/ir/block_param.h
@@ -34,10 +34,11 @@
BlockParam& operator=(BlockParam&& inst) = delete;
/// @returns the type of the var
- const type::Type* Type() const override { return type; }
+ const type::Type* Type() const override { return type_; }
+ private:
/// the result type of the instruction
- const type::Type* type = nullptr;
+ const type::Type* type_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc
index 983815e..b0b04d6 100644
--- a/src/tint/ir/builder.cc
+++ b/src/tint/ir/builder.cc
@@ -32,7 +32,7 @@
// Everything in the module scope must have been const-eval's, so everything will go into a
// single block. So, we can create the root terminator for the root-block now.
- ir.root_block->branch.target = CreateRootTerminator();
+ ir.root_block->BranchTo(CreateRootTerminator());
}
return ir.root_block;
}
@@ -56,11 +56,11 @@
TINT_ASSERT(IR, return_type);
auto* ir_func = ir.flow_nodes.Create<Function>(name, return_type, stage, wg_size);
- ir_func->start_target = CreateBlock();
- ir_func->end_target = CreateFunctionTerminator();
+ ir_func->SetStartTarget(CreateBlock());
+ ir_func->SetEndTarget(CreateFunctionTerminator());
- // Function is always branching into the start target
- ir_func->start_target->inbound_branches.Push(ir_func);
+ // Function is always branching into the Start().target
+ ir_func->StartTarget()->AddInboundBranch(ir_func);
return ir_func;
}
@@ -69,53 +69,48 @@
TINT_ASSERT(IR, condition);
auto* ir_if = ir.flow_nodes.Create<If>(condition);
- ir_if->true_.target = CreateBlock();
- ir_if->false_.target = CreateBlock();
- ir_if->merge.target = CreateBlock();
+ ir_if->True().target = CreateBlock();
+ ir_if->False().target = CreateBlock();
+ ir_if->Merge().target = CreateBlock();
// An if always branches to both the true and false block.
- ir_if->true_.target->inbound_branches.Push(ir_if);
- ir_if->false_.target->inbound_branches.Push(ir_if);
+ ir_if->True().target->AddInboundBranch(ir_if);
+ ir_if->False().target->AddInboundBranch(ir_if);
return ir_if;
}
Loop* Builder::CreateLoop() {
auto* ir_loop = ir.flow_nodes.Create<Loop>();
- ir_loop->start.target = CreateBlock();
- ir_loop->continuing.target = CreateBlock();
- ir_loop->merge.target = CreateBlock();
+ ir_loop->Start().target = CreateBlock();
+ ir_loop->Continuing().target = CreateBlock();
+ ir_loop->Merge().target = CreateBlock();
// A loop always branches to the start block.
- ir_loop->start.target->inbound_branches.Push(ir_loop);
+ ir_loop->Start().target->AddInboundBranch(ir_loop);
return ir_loop;
}
Switch* Builder::CreateSwitch(Value* condition) {
auto* ir_switch = ir.flow_nodes.Create<Switch>(condition);
- ir_switch->merge.target = CreateBlock();
+ ir_switch->Merge().target = CreateBlock();
return ir_switch;
}
Block* Builder::CreateCase(Switch* s, utils::VectorRef<Switch::CaseSelector> selectors) {
- s->cases.Push(Switch::Case{selectors, {CreateBlock(), utils::Empty}});
+ s->Cases().Push(Switch::Case{selectors, {CreateBlock(), utils::Empty}});
- Block* b = s->cases.Back().start.target->As<Block>();
+ Block* b = s->Cases().Back().Start().target->As<Block>();
// Switch branches into the case block
- b->inbound_branches.Push(s);
+ b->AddInboundBranch(s);
return b;
}
-void Builder::Branch(Block* from, FlowNode* to, utils::VectorRef<Value*> args) {
- TINT_ASSERT(IR, from);
- TINT_ASSERT(IR, to);
- from->branch.target = to;
- from->branch.args = args;
- to->inbound_branches.Push(from);
-}
-
-Binary* Builder::CreateBinary(Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs) {
+Binary* Builder::CreateBinary(enum Binary::Kind kind,
+ const type::Type* type,
+ Value* lhs,
+ Value* rhs) {
return ir.values.Create<ir::Binary>(kind, type, lhs, rhs);
}
@@ -183,7 +178,7 @@
return CreateBinary(Binary::Kind::kModulo, type, lhs, rhs);
}
-Unary* Builder::CreateUnary(Unary::Kind kind, const type::Type* type, Value* val) {
+Unary* Builder::CreateUnary(enum Unary::Kind kind, const type::Type* type, Value* val) {
return ir.values.Create<ir::Unary>(kind, type, val);
}
@@ -247,4 +242,8 @@
return ir.values.Create<ir::BlockParam>(type);
}
+ir::FunctionParam* Builder::FunctionParam(const type::Type* type) {
+ return ir.values.Create<ir::FunctionParam>(type);
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h
index 7c0cab6..ebd7a87 100644
--- a/src/tint/ir/builder.h
+++ b/src/tint/ir/builder.h
@@ -20,12 +20,14 @@
#include "src/tint/constant/scalar.h"
#include "src/tint/ir/binary.h"
#include "src/tint/ir/bitcast.h"
+#include "src/tint/ir/block_param.h"
#include "src/tint/ir/builtin.h"
#include "src/tint/ir/constant.h"
#include "src/tint/ir/construct.h"
#include "src/tint/ir/convert.h"
#include "src/tint/ir/discard.h"
#include "src/tint/ir/function.h"
+#include "src/tint/ir/function_param.h"
#include "src/tint/ir/function_terminator.h"
#include "src/tint/ir/if.h"
#include "src/tint/ir/load.h"
@@ -96,26 +98,20 @@
/// @returns the start block for the case flow node
Block* CreateCase(Switch* s, utils::VectorRef<Switch::CaseSelector> selectors);
- /// Branches the given block to the given flow node.
- /// @param from the block to branch from
- /// @param to the node to branch too
- /// @param args arguments to the branch
- void Branch(Block* from, FlowNode* to, utils::VectorRef<Value*> args = {});
-
/// Creates a constant::Value
/// @param args the arguments
/// @returns the new constant value
template <typename T, typename... ARGS>
utils::traits::EnableIf<utils::traits::IsTypeOrDerived<T, constant::Value>, const T>* create(
ARGS&&... args) {
- return ir.constants.Create<T>(std::forward<ARGS>(args)...);
+ return ir.constants_arena.Create<T>(std::forward<ARGS>(args)...);
}
/// Creates a new ir::Constant
/// @param val the constant value
/// @returns the new constant
ir::Constant* Constant(const constant::Value* val) {
- return ir.values.Create<ir::Constant>(val);
+ return ir.constants.GetOrCreate(val, [&]() { return ir.values.Create<ir::Constant>(val); });
}
/// Creates a ir::Constant for an i32 Scalar
@@ -159,7 +155,7 @@
/// @param lhs the left-hand-side of the operation
/// @param rhs the right-hand-side of the operation
/// @returns the operation
- Binary* CreateBinary(Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs);
+ Binary* CreateBinary(enum Binary::Kind kind, const type::Type* type, Value* lhs, Value* rhs);
/// Creates an And operation
/// @param type the result type of the expression
@@ -278,7 +274,7 @@
/// @param type the result type of the binary expression
/// @param val the value of the operation
/// @returns the operation
- Unary* CreateUnary(Unary::Kind kind, const type::Type* type, Value* val);
+ Unary* CreateUnary(enum Unary::Kind kind, const type::Type* type, Value* val);
/// Creates a Complement operation
/// @param type the result type of the expression
@@ -360,6 +356,11 @@
/// @returns the value
ir::BlockParam* BlockParam(const type::Type* type);
+ /// Creates a new `FunctionParam`
+ /// @param type the parameter type
+ /// @returns the value
+ ir::FunctionParam* FunctionParam(const type::Type* type);
+
/// Retrieves the root block for the module, creating if necessary
/// @returns the root block
ir::Block* CreateRootBlockIfNeeded();
diff --git a/src/tint/ir/call.cc b/src/tint/ir/call.cc
index dd503ea..909079f 100644
--- a/src/tint/ir/call.cc
+++ b/src/tint/ir/call.cc
@@ -21,8 +21,8 @@
namespace tint::ir {
Call::Call(const type::Type* res_ty, utils::VectorRef<Value*> arguments)
- : result_type(res_ty), args(std::move(arguments)) {
- for (auto* arg : args) {
+ : result_type_(res_ty), args_(std::move(arguments)) {
+ for (auto* arg : args_) {
arg->AddUsage(this);
}
}
diff --git a/src/tint/ir/call.h b/src/tint/ir/call.h
index 5393810..f4e12f9 100644
--- a/src/tint/ir/call.h
+++ b/src/tint/ir/call.h
@@ -31,13 +31,10 @@
Call& operator=(Call&& inst) = delete;
/// @returns the type of the value
- const type::Type* Type() const override { return result_type; }
+ const type::Type* Type() const override { return result_type_; }
- /// The instruction type
- const type::Type* result_type = nullptr;
-
- /// The constructor arguments
- utils::Vector<Value*, 1> args;
+ /// @returns the call arguments
+ utils::VectorRef<Value*> Args() const { return args_; }
protected:
/// Constructor
@@ -46,6 +43,10 @@
/// @param result_type the result type
/// @param args the constructor arguments
Call(const type::Type* result_type, utils::VectorRef<Value*> args);
+
+ private:
+ const type::Type* result_type_;
+ utils::Vector<Value*, 1> args_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/constant.cc b/src/tint/ir/constant.cc
index 8b5260c..a49d140 100644
--- a/src/tint/ir/constant.cc
+++ b/src/tint/ir/constant.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-Constant::Constant(const constant::Value* val) : value(val) {}
+Constant::Constant(const constant::Value* val) : value_(val) {}
Constant::~Constant() = default;
diff --git a/src/tint/ir/constant.h b/src/tint/ir/constant.h
index 68e0dc7..dcf3e5b 100644
--- a/src/tint/ir/constant.h
+++ b/src/tint/ir/constant.h
@@ -26,13 +26,21 @@
/// Constructor
/// @param val the value stored in the constant
explicit Constant(const constant::Value* val);
+ Constant(const Constant&) = delete;
+ Constant(Constant&&) = delete;
~Constant() override;
- /// @returns the type of the constant
- const type::Type* Type() const override { return value->Type(); }
+ Constant& operator=(const Constant&) = delete;
+ Constant& operator=(Constant&&) = delete;
- /// The constants value
- const constant::Value* const value;
+ /// @returns the constants value
+ const constant::Value* Value() const { return value_; }
+
+ /// @returns the type of the constant
+ const type::Type* Type() const override { return value_->Type(); }
+
+ private:
+ const constant::Value* const value_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/constant_test.cc b/src/tint/ir/constant_test.cc
index 7005751..7c3024a 100644
--- a/src/tint/ir/constant_test.cc
+++ b/src/tint/ir/constant_test.cc
@@ -30,13 +30,13 @@
utils::StringStream str;
auto* c = b.Constant(1.2_f);
- EXPECT_EQ(1.2_f, c->value->As<constant::Scalar<f32>>()->ValueAs<f32>());
+ EXPECT_EQ(1.2_f, c->Value()->As<constant::Scalar<f32>>()->ValueAs<f32>());
- EXPECT_TRUE(c->value->Is<constant::Scalar<f32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+ EXPECT_TRUE(c->Value()->Is<constant::Scalar<f32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
}
TEST_F(IR_ConstantTest, f16) {
@@ -46,13 +46,13 @@
utils::StringStream str;
auto* c = b.Constant(1.1_h);
- EXPECT_EQ(1.1_h, c->value->As<constant::Scalar<f16>>()->ValueAs<f16>());
+ EXPECT_EQ(1.1_h, c->Value()->As<constant::Scalar<f16>>()->ValueAs<f16>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
- EXPECT_TRUE(c->value->Is<constant::Scalar<f16>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+ EXPECT_TRUE(c->Value()->Is<constant::Scalar<f16>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
}
TEST_F(IR_ConstantTest, i32) {
@@ -62,13 +62,13 @@
utils::StringStream str;
auto* c = b.Constant(1_i);
- EXPECT_EQ(1_i, c->value->As<constant::Scalar<i32>>()->ValueAs<i32>());
+ EXPECT_EQ(1_i, c->Value()->As<constant::Scalar<i32>>()->ValueAs<i32>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
- EXPECT_TRUE(c->value->Is<constant::Scalar<i32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+ EXPECT_TRUE(c->Value()->Is<constant::Scalar<i32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
}
TEST_F(IR_ConstantTest, u32) {
@@ -78,13 +78,13 @@
utils::StringStream str;
auto* c = b.Constant(2_u);
- EXPECT_EQ(2_u, c->value->As<constant::Scalar<u32>>()->ValueAs<u32>());
+ EXPECT_EQ(2_u, c->Value()->As<constant::Scalar<u32>>()->ValueAs<u32>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
- EXPECT_TRUE(c->value->Is<constant::Scalar<u32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<bool>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+ EXPECT_TRUE(c->Value()->Is<constant::Scalar<u32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<bool>>());
}
TEST_F(IR_ConstantTest, bool) {
@@ -95,19 +95,19 @@
utils::StringStream str;
auto* c = b.Constant(false);
- EXPECT_FALSE(c->value->As<constant::Scalar<bool>>()->ValueAs<bool>());
+ EXPECT_FALSE(c->Value()->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
{
utils::StringStream str;
auto c = b.Constant(true);
- EXPECT_TRUE(c->value->As<constant::Scalar<bool>>()->ValueAs<bool>());
+ EXPECT_TRUE(c->Value()->As<constant::Scalar<bool>>()->ValueAs<bool>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<f16>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<i32>>());
- EXPECT_FALSE(c->value->Is<constant::Scalar<u32>>());
- EXPECT_TRUE(c->value->Is<constant::Scalar<bool>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<f16>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<i32>>());
+ EXPECT_FALSE(c->Value()->Is<constant::Scalar<u32>>());
+ EXPECT_TRUE(c->Value()->Is<constant::Scalar<bool>>());
}
}
diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc
index 655c455..85666ff 100644
--- a/src/tint/ir/debug.cc
+++ b/src/tint/ir/debug.cc
@@ -60,81 +60,81 @@
if (node_to_name.count(b) == 0) {
out << name_for(b) << R"( [label="block"])" << std::endl;
}
- out << name_for(b) << " -> " << name_for(b->branch.target);
+ out << name_for(b) << " -> " << name_for(b->Branch().target);
// Dashed lines to merge blocks
- if (merge_nodes.count(b->branch.target) != 0) {
+ if (merge_nodes.count(b->Branch().target) != 0) {
out << " [style=dashed]";
}
out << std::endl;
- Graph(b->branch.target);
+ Graph(b->Branch().target);
},
[&](const ir::Switch* s) {
out << name_for(s) << R"( [label="switch"])" << std::endl;
- out << name_for(s->merge.target) << R"( [label="switch merge"])" << std::endl;
- merge_nodes.insert(s->merge.target);
+ out << name_for(s->Merge().target) << R"( [label="switch merge"])" << std::endl;
+ merge_nodes.insert(s->Merge().target);
size_t i = 0;
- for (const auto& c : s->cases) {
- out << name_for(c.start.target)
+ for (const auto& c : s->Cases()) {
+ out << name_for(c.Start().target)
<< R"( [label="case )" + std::to_string(i++) + R"("])" << std::endl;
}
out << name_for(s) << " -> {";
- for (const auto& c : s->cases) {
- if (&c != &(s->cases[0])) {
+ for (const auto& c : s->Cases()) {
+ if (&c != &(s->Cases().Front())) {
out << ", ";
}
- out << name_for(c.start.target);
+ out << name_for(c.Start().target);
}
out << "}" << std::endl;
- for (const auto& c : s->cases) {
- Graph(c.start.target);
+ for (const auto& c : s->Cases()) {
+ Graph(c.Start().target);
}
- Graph(s->merge.target);
+ Graph(s->Merge().target);
},
[&](const ir::If* i) {
out << name_for(i) << R"( [label="if"])" << std::endl;
- out << name_for(i->true_.target) << R"( [label="true"])" << std::endl;
- out << name_for(i->false_.target) << R"( [label="false"])" << std::endl;
- out << name_for(i->merge.target) << R"( [label="if merge"])" << std::endl;
- merge_nodes.insert(i->merge.target);
+ out << name_for(i->True().target) << R"( [label="true"])" << std::endl;
+ out << name_for(i->False().target) << R"( [label="false"])" << std::endl;
+ out << name_for(i->Merge().target) << R"( [label="if merge"])" << std::endl;
+ merge_nodes.insert(i->Merge().target);
out << name_for(i) << " -> {";
- out << name_for(i->true_.target) << ", " << name_for(i->false_.target);
+ out << name_for(i->True().target) << ", " << name_for(i->False().target);
out << "}" << std::endl;
// Subgraph if true/false branches so they draw on the same line
out << "subgraph sub_" << name_for(i) << " {" << std::endl;
out << R"(rank="same")" << std::endl;
- out << name_for(i->true_.target) << std::endl;
- out << name_for(i->false_.target) << std::endl;
+ out << name_for(i->True().target) << std::endl;
+ out << name_for(i->False().target) << std::endl;
out << "}" << std::endl;
- Graph(i->true_.target);
- Graph(i->false_.target);
- Graph(i->merge.target);
+ Graph(i->True().target);
+ Graph(i->False().target);
+ Graph(i->Merge().target);
},
[&](const ir::Loop* l) {
out << name_for(l) << R"( [label="loop"])" << std::endl;
- out << name_for(l->start.target) << R"( [label="start"])" << std::endl;
- out << name_for(l->continuing.target) << R"( [label="continuing"])" << std::endl;
- out << name_for(l->merge.target) << R"( [label="loop merge"])" << std::endl;
- merge_nodes.insert(l->merge.target);
+ out << name_for(l->Start().target) << R"( [label="start"])" << std::endl;
+ out << name_for(l->Continuing().target) << R"( [label="continuing"])" << std::endl;
+ out << name_for(l->Merge().target) << R"( [label="loop merge"])" << std::endl;
+ merge_nodes.insert(l->Merge().target);
// Subgraph the continuing and merge so they get drawn on the same line
out << "subgraph sub_" << name_for(l) << " {" << std::endl;
out << R"(rank="same")" << std::endl;
- out << name_for(l->continuing.target) << std::endl;
- out << name_for(l->merge.target) << std::endl;
+ out << name_for(l->Continuing().target) << std::endl;
+ out << name_for(l->Merge().target) << std::endl;
out << "}" << std::endl;
- out << name_for(l) << " -> " << name_for(l->start.target) << std::endl;
+ out << name_for(l) << " -> " << name_for(l->Start().target) << std::endl;
- Graph(l->start.target);
- Graph(l->continuing.target);
- Graph(l->merge.target);
+ Graph(l->Start().target);
+ Graph(l->Continuing().target);
+ Graph(l->Merge().target);
},
[&](const ir::FunctionTerminator*) {
// Already done
@@ -145,10 +145,10 @@
for (const auto* func : mod->functions) {
// Cluster each function to label and draw a box around it.
out << "subgraph cluster_" << name_for(func) << " {" << std::endl;
- out << R"(label=")" << func->name.Name() << R"(")" << std::endl;
- out << name_for(func->start_target) << R"( [label="start"])" << std::endl;
- out << name_for(func->end_target) << R"( [label="end"])" << std::endl;
- Graph(func->start_target);
+ out << R"(label=")" << func->Name().Name() << R"(")" << std::endl;
+ out << name_for(func->StartTarget()) << R"( [label="start"])" << std::endl;
+ out << name_for(func->EndTarget()) << R"( [label="end"])" << std::endl;
+ Graph(func->StartTarget());
out << "}" << std::endl;
}
out << "}";
diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc
index f68b21a..76af740 100644
--- a/src/tint/ir/disassembler.cc
+++ b/src/tint/ir/disassembler.cc
@@ -81,7 +81,7 @@
}
void Disassembler::EmitBlockInstructions(const Block* b) {
- for (const auto* inst : b->instructions) {
+ for (const auto* inst : b->Instructions()) {
Indent();
EmitInstruction(inst);
out_ << std::endl;
@@ -114,25 +114,31 @@
[&](const ir::Function* f) {
TINT_SCOPED_ASSIGNMENT(in_function_, true);
- Indent() << "%fn" << IdOf(f) << " = func " << f->name.Name()
- << "():" << f->return_type->FriendlyName();
+ Indent() << "%fn" << IdOf(f) << " = func " << f->Name().Name() << "(";
+ for (auto* p : f->Params()) {
+ if (p != f->Params().Front()) {
+ out_ << ", ";
+ }
+ out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
+ }
+ out_ << "):" << f->ReturnType()->FriendlyName();
- if (f->pipeline_stage != Function::PipelineStage::kUndefined) {
- out_ << " [@" << f->pipeline_stage;
+ if (f->Stage() != Function::PipelineStage::kUndefined) {
+ out_ << " [@" << f->Stage();
- if (f->workgroup_size) {
- auto arr = f->workgroup_size.value();
+ if (f->WorkgroupSize()) {
+ auto arr = f->WorkgroupSize().value();
out_ << " @workgroup_size(" << arr[0] << ", " << arr[1] << ", " << arr[2]
<< ")";
}
- if (!f->return_attributes.IsEmpty()) {
+ if (!f->ReturnAttributes().IsEmpty()) {
out_ << " ra:";
- for (auto attr : f->return_attributes) {
+ for (auto attr : f->ReturnAttributes()) {
out_ << " @" << attr;
if (attr == Function::ReturnAttribute::kLocation) {
- out_ << "(" << f->return_location.value() << ")";
+ out_ << "(" << f->ReturnLocation().value() << ")";
}
}
}
@@ -143,23 +149,23 @@
{
ScopedIndent func_indent(indent_size_);
- ScopedStopNode scope(stop_nodes_, f->end_target);
- Walk(f->start_target);
+ ScopedStopNode scope(stop_nodes_, f->EndTarget());
+ Walk(f->StartTarget());
}
out_ << "} ";
- Walk(f->end_target);
+ Walk(f->EndTarget());
},
[&](const ir::Block* b) {
// If this block is dead, nothing to do
- if (b->IsDead()) {
+ if (!b->HasBranchTarget()) {
return;
}
Indent() << "%fn" << IdOf(b) << " = block";
- if (!b->params.IsEmpty()) {
+ if (!b->Params().IsEmpty()) {
out_ << " (";
- for (auto* p : b->params) {
- if (p != b->params.Front()) {
+ for (const auto* p : b->Params()) {
+ if (p != b->Params().Front()) {
out_ << ", ";
}
EmitValue(p);
@@ -175,20 +181,20 @@
Indent() << "}";
std::string suffix = "";
- if (b->branch.target->Is<FunctionTerminator>()) {
+ if (b->Branch().target->Is<FunctionTerminator>()) {
out_ << " -> %func_end";
suffix = "return";
- } else if (b->branch.target->Is<RootTerminator>()) {
+ } else if (b->Branch().target->Is<RootTerminator>()) {
// Nothing to do
} else {
out_ << " -> "
- << "%fn" << IdOf(b->branch.target);
+ << "%fn" << IdOf(b->Branch().target);
suffix = "branch";
}
- if (!b->branch.args.IsEmpty()) {
+ if (!b->Branch().args.IsEmpty()) {
out_ << " ";
- for (const auto* v : b->branch.args) {
- if (v != b->branch.args.Front()) {
+ for (const auto* v : b->Branch().args) {
+ if (v != b->Branch().args.Front()) {
out_ << ", ";
}
EmitValue(v);
@@ -199,18 +205,18 @@
}
out_ << std::endl;
- if (!b->branch.target->Is<FunctionTerminator>()) {
+ if (!b->Branch().target->Is<FunctionTerminator>()) {
out_ << std::endl;
}
- Walk(b->branch.target);
+ Walk(b->Branch().target);
},
[&](const ir::Switch* s) {
Indent() << "%fn" << IdOf(s) << " = switch ";
- EmitValue(s->condition);
+ EmitValue(s->Condition());
out_ << " [";
- for (const auto& c : s->cases) {
- if (&c != &s->cases.Front()) {
+ for (const auto& c : s->Cases()) {
+ if (&c != &s->Cases().Front()) {
out_ << ", ";
}
out_ << "c: (";
@@ -225,17 +231,17 @@
EmitValue(selector.val);
}
}
- out_ << ", %fn" << IdOf(c.start.target) << ")";
+ out_ << ", %fn" << IdOf(c.Start().target) << ")";
}
- if (s->merge.target->IsConnected()) {
- out_ << ", m: %fn" << IdOf(s->merge.target);
+ if (s->Merge().target->IsConnected()) {
+ out_ << ", m: %fn" << IdOf(s->Merge().target);
}
out_ << "]" << std::endl;
{
ScopedIndent switch_indent(indent_size_);
- ScopedStopNode scope(stop_nodes_, s->merge.target);
- for (const auto& c : s->cases) {
+ ScopedStopNode scope(stop_nodes_, s->Merge().target);
+ for (const auto& c : s->Cases()) {
Indent() << "# case ";
for (const auto& selector : c.selectors) {
if (&selector != &c.selectors.Front()) {
@@ -249,86 +255,86 @@
}
}
out_ << std::endl;
- Walk(c.start.target);
+ Walk(c.Start().target);
}
}
- if (s->merge.target->IsConnected()) {
+ if (s->Merge().target->IsConnected()) {
Indent() << "# switch merge" << std::endl;
- Walk(s->merge.target);
+ Walk(s->Merge().target);
}
},
[&](const ir::If* i) {
Indent() << "%fn" << IdOf(i) << " = if ";
- EmitValue(i->condition);
+ EmitValue(i->Condition());
- bool has_true = !i->true_.target->IsDead();
- bool has_false = !i->false_.target->IsDead();
+ bool has_true = i->True().target->HasBranchTarget();
+ bool has_false = i->False().target->HasBranchTarget();
out_ << " [";
if (has_true) {
- out_ << "t: %fn" << IdOf(i->true_.target);
+ out_ << "t: %fn" << IdOf(i->True().target);
}
if (has_false) {
if (has_true) {
out_ << ", ";
}
- out_ << "f: %fn" << IdOf(i->false_.target);
+ out_ << "f: %fn" << IdOf(i->False().target);
}
- if (i->merge.target->IsConnected()) {
- out_ << ", m: %fn" << IdOf(i->merge.target);
+ if (i->Merge().target->IsConnected()) {
+ out_ << ", m: %fn" << IdOf(i->Merge().target);
}
out_ << "]" << std::endl;
{
ScopedIndent if_indent(indent_size_);
- ScopedStopNode scope(stop_nodes_, i->merge.target);
+ ScopedStopNode scope(stop_nodes_, i->Merge().target);
if (has_true) {
Indent() << "# true branch" << std::endl;
- Walk(i->true_.target);
+ Walk(i->True().target);
}
if (has_false) {
Indent() << "# false branch" << std::endl;
- Walk(i->false_.target);
+ Walk(i->False().target);
}
}
- if (i->merge.target->IsConnected()) {
+ if (i->Merge().target->IsConnected()) {
Indent() << "# if merge" << std::endl;
- Walk(i->merge.target);
+ Walk(i->Merge().target);
}
},
[&](const ir::Loop* l) {
- Indent() << "%fn" << IdOf(l) << " = loop [s: %fn" << IdOf(l->start.target);
+ Indent() << "%fn" << IdOf(l) << " = loop [s: %fn" << IdOf(l->Start().target);
- if (l->continuing.target->IsConnected()) {
- out_ << ", c: %fn" << IdOf(l->continuing.target);
+ if (l->Continuing().target->IsConnected()) {
+ out_ << ", c: %fn" << IdOf(l->Continuing().target);
}
- if (l->merge.target->IsConnected()) {
- out_ << ", m: %fn" << IdOf(l->merge.target);
+ if (l->Merge().target->IsConnected()) {
+ out_ << ", m: %fn" << IdOf(l->Merge().target);
}
out_ << "]" << std::endl;
{
- ScopedStopNode loop_scope(stop_nodes_, l->merge.target);
+ ScopedStopNode loop_scope(stop_nodes_, l->Merge().target);
ScopedIndent loop_indent(indent_size_);
{
- ScopedStopNode inner_scope(stop_nodes_, l->continuing.target);
+ ScopedStopNode inner_scope(stop_nodes_, l->Continuing().target);
Indent() << "# loop start" << std::endl;
- Walk(l->start.target);
+ Walk(l->Start().target);
}
- if (l->continuing.target->IsConnected()) {
+ if (l->Continuing().target->IsConnected()) {
Indent() << "# loop continuing" << std::endl;
- Walk(l->continuing.target);
+ Walk(l->Continuing().target);
}
}
- if (l->merge.target->IsConnected()) {
+ if (l->Merge().target->IsConnected()) {
Indent() << "# loop merge" << std::endl;
- Walk(l->merge.target);
+ Walk(l->Merge().target);
}
},
[&](const ir::FunctionTerminator*) {
@@ -401,12 +407,13 @@
}
});
};
- emit(constant->value);
+ emit(constant->Value());
},
[&](const ir::Instruction* i) { out_ << "%" << IdOf(i); },
[&](const ir::BlockParam* p) {
out_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
},
+ [&](const ir::FunctionParam* p) { out_ << "%" << IdOf(p); },
[&](Default) { out_ << "Unknown value: " << val->TypeInfo().name; });
}
@@ -438,18 +445,18 @@
[&](const ir::Load* l) {
EmitValueWithType(l);
out_ << " = load ";
- EmitValue(l->from);
+ EmitValue(l->From());
},
[&](const ir::Store* s) {
out_ << "store ";
- EmitValue(s->to);
+ EmitValue(s->To());
out_ << ", ";
- EmitValue(s->from);
+ EmitValue(s->From());
},
[&](const ir::UserCall* uc) {
EmitValueWithType(uc);
- out_ << " = call " << uc->name.Name();
- if (uc->args.Length() > 0) {
+ out_ << " = call " << uc->Name().Name();
+ if (!uc->Args().IsEmpty()) {
out_ << ", ";
}
EmitArgs(uc);
@@ -457,16 +464,16 @@
[&](const ir::Var* v) {
EmitValueWithType(v);
out_ << " = var";
- if (v->initializer) {
+ if (v->Initializer()) {
out_ << ", ";
- EmitValue(v->initializer);
+ EmitValue(v->Initializer());
}
});
}
void Disassembler::EmitArgs(const Call* call) {
bool first = true;
- for (const auto* arg : call->args) {
+ for (const auto* arg : call->Args()) {
if (!first) {
out_ << ", ";
}
@@ -478,7 +485,7 @@
void Disassembler::EmitBinary(const Binary* b) {
EmitValueWithType(b);
out_ << " = ";
- switch (b->kind) {
+ switch (b->Kind()) {
case Binary::Kind::kAdd:
out_ << "add";
break;
@@ -537,7 +544,7 @@
void Disassembler::EmitUnary(const Unary* u) {
EmitValueWithType(u);
out_ << " = ";
- switch (u->kind) {
+ switch (u->Kind()) {
case Unary::Kind::kComplement:
out_ << "complement";
break;
diff --git a/src/tint/ir/flow_node.h b/src/tint/ir/flow_node.h
index 905f077..289873b 100644
--- a/src/tint/ir/flow_node.h
+++ b/src/tint/ir/flow_node.h
@@ -25,22 +25,30 @@
public:
~FlowNode() override;
- /// The list of flow nodes which branch into this node. This list maybe empty for several
- /// reasons:
- /// - Node is a start node
- /// - Node is a merge target outside control flow (if that returns in both branches)
- /// - Node is a continue target outside control flow (loop that returns)
- utils::Vector<FlowNode*, 2> inbound_branches;
-
/// @returns true if this node has inbound branches and branches out
- bool IsConnected() const { return !IsDead() && !inbound_branches.IsEmpty(); }
+ bool IsConnected() const { return HasBranchTarget() && !inbound_branches_.IsEmpty(); }
- /// @returns true if the node does not branch out
- virtual bool IsDead() const { return false; }
+ /// @returns true if the node has a branch target
+ virtual bool HasBranchTarget() const { return false; }
+
+ /// @returns the inbound branch list for the flow node
+ utils::VectorRef<FlowNode*> InboundBranches() const { return inbound_branches_; }
+
+ /// Adds the given node to the inbound branches
+ /// @param node the node to add
+ void AddInboundBranch(FlowNode* node) { inbound_branches_.Push(node); }
protected:
/// Constructor
FlowNode();
+
+ private:
+ /// The list of flow nodes which branch into this node. This list maybe empty for several
+ /// reasons:
+ /// - Node is a start node
+ /// - Node is a merge target outside control flow (e.g. an if that returns in both branches)
+ /// - Node is a continue target outside control flow (e.g. a loop that returns)
+ utils::Vector<FlowNode*, 2> inbound_branches_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc
index bb9c3f9..595db58 100644
--- a/src/tint/ir/from_program.cc
+++ b/src/tint/ir/from_program.cc
@@ -98,17 +98,13 @@
using ResultType = utils::Result<Module, diag::List>;
-bool IsBranched(const Block* b) {
- return b->branch.target != nullptr;
-}
-
bool IsConnected(const FlowNode* b) {
// Function is always connected as it's the start.
if (b->Is<ir::Function>()) {
return true;
}
- for (auto* parent : b->inbound_branches) {
+ for (auto* parent : b->InboundBranches()) {
if (IsConnected(parent)) {
return true;
}
@@ -146,7 +142,7 @@
/* src */ {&program_->Symbols()},
/* dst */ {&builder_.ir.symbols, &builder_.ir.types},
},
- /* dst */ {&builder_.ir.constants},
+ /* dst */ {&builder_.ir.constants_arena},
};
/// The stack of flow control blocks.
@@ -184,14 +180,14 @@
void BranchTo(FlowNode* node, utils::VectorRef<Value*> args = {}) {
TINT_ASSERT(IR, current_flow_block_);
- TINT_ASSERT(IR, !IsBranched(current_flow_block_));
+ TINT_ASSERT(IR, !current_flow_block_->HasBranchTarget());
- builder_.Branch(current_flow_block_, node, args);
+ current_flow_block_->BranchTo(node, args);
current_flow_block_ = nullptr;
}
void BranchToIfNeeded(FlowNode* node) {
- if (!current_flow_block_ || IsBranched(current_flow_block_)) {
+ if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) {
return;
}
BranchTo(node);
@@ -271,20 +267,17 @@
if (ast_func->IsEntryPoint()) {
switch (ast_func->PipelineStage()) {
case ast::PipelineStage::kVertex:
- ir_func->pipeline_stage = Function::PipelineStage::kVertex;
+ ir_func->SetStage(Function::PipelineStage::kVertex);
break;
case ast::PipelineStage::kFragment:
- ir_func->pipeline_stage = Function::PipelineStage::kFragment;
+ ir_func->SetStage(Function::PipelineStage::kFragment);
break;
case ast::PipelineStage::kCompute: {
- ir_func->pipeline_stage = Function::PipelineStage::kCompute;
+ ir_func->SetStage(Function::PipelineStage::kCompute);
auto wg_size = sem->WorkgroupSize();
- ir_func->workgroup_size = {
- wg_size[0].value(),
- wg_size[1].value_or(1),
- wg_size[2].value_or(1),
- };
+ ir_func->SetWorkgroupSize(wg_size[0].value(), wg_size[1].value_or(1),
+ wg_size[2].value_or(1));
break;
}
default: {
@@ -293,14 +286,15 @@
}
}
+ utils::Vector<Function::ReturnAttribute, 1> return_attributes;
for (auto* attr : ast_func->return_type_attributes) {
tint::Switch(
attr, //
[&](const ast::LocationAttribute*) {
- ir_func->return_attributes.Push(Function::ReturnAttribute::kLocation);
+ return_attributes.Push(Function::ReturnAttribute::kLocation);
},
[&](const ast::InvariantAttribute*) {
- ir_func->return_attributes.Push(Function::ReturnAttribute::kInvariant);
+ return_attributes.Push(Function::ReturnAttribute::kInvariant);
},
[&](const ast::BuiltinAttribute* b) {
if (auto* ident_sem =
@@ -309,16 +303,13 @@
->As<sem::BuiltinEnumExpression<builtin::BuiltinValue>>()) {
switch (ident_sem->Value()) {
case builtin::BuiltinValue::kPosition:
- ir_func->return_attributes.Push(
- Function::ReturnAttribute::kPosition);
+ return_attributes.Push(Function::ReturnAttribute::kPosition);
break;
case builtin::BuiltinValue::kFragDepth:
- ir_func->return_attributes.Push(
- Function::ReturnAttribute::kFragDepth);
+ return_attributes.Push(Function::ReturnAttribute::kFragDepth);
break;
case builtin::BuiltinValue::kSampleMask:
- ir_func->return_attributes.Push(
- Function::ReturnAttribute::kSampleMask);
+ return_attributes.Push(Function::ReturnAttribute::kSampleMask);
break;
default:
TINT_ICE(IR, diagnostics_)
@@ -332,20 +323,34 @@
}
});
}
+ ir_func->SetReturnAttributes(return_attributes);
}
- ir_func->return_location = sem->ReturnLocation();
+ ir_func->SetReturnLocation(sem->ReturnLocation());
+
+ scopes_.Push();
+ TINT_DEFER(scopes_.Pop());
+
+ utils::Vector<FunctionParam*, 1> params;
+ for (auto* p : ast_func->params) {
+ const auto* param_sem = program_->Sem().Get(p);
+ auto* ty = param_sem->Type()->Clone(clone_ctx_.type_ctx);
+ auto* param = builder_.FunctionParam(ty);
+
+ scopes_.Set(p->name->symbol, param);
+ builder_.ir.SetName(param, p->name->symbol.NameView());
+ params.Push(param);
+ }
+ ir_func->SetParams(params);
{
FlowStackScope scope(this, ir_func);
- current_flow_block_ = ir_func->start_target;
+ current_flow_block_ = ir_func->StartTarget();
EmitBlock(ast_func->body);
- // TODO(dsinclair): Store parameters
-
// If the branch target has already been set then a `return` was called. Only set in the
// case where `return` wasn't called.
- BranchToIfNeeded(current_function_->end_target);
+ BranchToIfNeeded(current_function_->EndTarget());
}
TINT_ASSERT(IR, flow_stack_.IsEmpty());
@@ -359,7 +364,7 @@
// If the current flow block has a branch target then the rest of the statements in this
// block are dead code. Skip them.
- if (!current_flow_block_ || IsBranched(current_flow_block_)) {
+ if (!current_flow_block_ || current_flow_block_->HasBranchTarget()) {
break;
}
}
@@ -414,7 +419,7 @@
return;
}
auto store = builder_.Store(lhs.Get(), rhs.Get());
- current_flow_block_->instructions.Push(store);
+ current_flow_block_->Instructions().Push(store);
}
void EmitIncrementDecrement(const ast::IncrementDecrementStatement* stmt) {
@@ -425,7 +430,7 @@
// Load from the LHS.
auto* lhs_value = builder_.Load(lhs.Get());
- current_flow_block_->instructions.Push(lhs_value);
+ current_flow_block_->Instructions().Push(lhs_value);
auto* ty = lhs_value->Type();
@@ -438,10 +443,10 @@
} else {
inst = builder_.Subtract(ty, lhs_value, rhs);
}
- current_flow_block_->instructions.Push(inst);
+ current_flow_block_->Instructions().Push(inst);
auto store = builder_.Store(lhs.Get(), inst);
- current_flow_block_->instructions.Push(store);
+ current_flow_block_->Instructions().Push(store);
}
void EmitCompoundAssignment(const ast::CompoundAssignmentStatement* stmt) {
@@ -457,7 +462,7 @@
// Load from the LHS.
auto* lhs_value = builder_.Load(lhs.Get());
- current_flow_block_->instructions.Push(lhs_value);
+ current_flow_block_->Instructions().Push(lhs_value);
auto* ty = lhs_value->Type();
@@ -507,10 +512,10 @@
TINT_ICE(IR, diagnostics_) << "missing binary operand type";
return;
}
- current_flow_block_->instructions.Push(inst);
+ current_flow_block_->Instructions().Push(inst);
auto store = builder_.Store(lhs.Get(), inst);
- current_flow_block_->instructions.Push(store);
+ current_flow_block_->Instructions().Push(store);
}
void EmitBlock(const ast::BlockStatement* block) {
@@ -538,27 +543,27 @@
{
FlowStackScope scope(this, if_node);
- current_flow_block_ = if_node->true_.target->As<Block>();
+ current_flow_block_ = if_node->True().target->As<Block>();
EmitBlock(stmt->body);
- // If the true branch did not execute control flow, then go to the merge target
- BranchToIfNeeded(if_node->merge.target);
+ // If the true branch did not execute control flow, then go to the Merge().target
+ BranchToIfNeeded(if_node->Merge().target);
- current_flow_block_ = if_node->false_.target->As<Block>();
+ current_flow_block_ = if_node->False().target->As<Block>();
if (stmt->else_statement) {
EmitStatement(stmt->else_statement);
}
- // If the false branch did not execute control flow, then go to the merge target
- BranchToIfNeeded(if_node->merge.target);
+ // If the false branch did not execute control flow, then go to the Merge().target
+ BranchToIfNeeded(if_node->Merge().target);
}
current_flow_block_ = nullptr;
// If both branches went somewhere, then they both returned, continued or broke. So, there
// is no need for the if merge-block and there is nothing to branch to the merge block
// anyway.
- if (IsConnected(if_node->merge.target)) {
- current_flow_block_ = if_node->merge.target->As<Block>();
+ if (IsConnected(if_node->Merge().target)) {
+ current_flow_block_ = if_node->Merge().target->As<Block>();
}
}
@@ -572,7 +577,7 @@
{
FlowStackScope scope(this, loop_node);
- current_flow_block_ = loop_node->start.target->As<Block>();
+ current_flow_block_ = loop_node->Start().target->As<Block>();
// The loop doesn't use EmitBlock because it needs the scope stack to not get popped
// until after the continuing block.
@@ -581,21 +586,22 @@
EmitStatements(stmt->body->statements);
// The current block didn't `break`, `return` or `continue`, go to the continuing block.
- BranchToIfNeeded(loop_node->continuing.target);
+ BranchToIfNeeded(loop_node->Continuing().target);
- current_flow_block_ = loop_node->continuing.target->As<Block>();
+ current_flow_block_ = loop_node->Continuing().target->As<Block>();
if (stmt->continuing) {
EmitBlock(stmt->continuing);
}
// Branch back to the start node if the continue target didn't branch out already
- BranchToIfNeeded(loop_node->start.target);
+ BranchToIfNeeded(loop_node->Start().target);
}
// The loop merge can get disconnected if the loop returns directly, or the continuing
- // target branches, eventually, to the merge, but nothing branched to the continuing target.
- current_flow_block_ = loop_node->merge.target->As<Block>();
- if (!IsConnected(loop_node->merge.target)) {
+ // target branches, eventually, to the merge, but nothing branched to the
+ // Continuing().target.
+ current_flow_block_ = loop_node->Merge().target->As<Block>();
+ if (!IsConnected(loop_node->Merge().target)) {
current_flow_block_ = nullptr;
}
}
@@ -603,9 +609,8 @@
void EmitWhile(const ast::WhileStatement* stmt) {
auto* loop_node = builder_.CreateLoop();
// Continue is always empty, just go back to the start
- TINT_ASSERT(IR, loop_node->continuing.target->Is<Block>());
- builder_.Branch(loop_node->continuing.target->As<Block>(), loop_node->start.target,
- utils::Empty);
+ TINT_ASSERT(IR, loop_node->Continuing().target->Is<Block>());
+ loop_node->Continuing().target->As<Block>()->BranchTo(loop_node->Start().target);
BranchTo(loop_node);
@@ -614,9 +619,9 @@
{
FlowStackScope scope(this, loop_node);
- current_flow_block_ = loop_node->start.target->As<Block>();
+ current_flow_block_ = loop_node->Start().target->As<Block>();
- // Emit the while condition into the start target of the loop
+ // Emit the while condition into the Start().target of the loop
auto reg = EmitExpression(stmt->condition);
if (!reg) {
return;
@@ -624,31 +629,24 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_node = builder_.CreateIf(reg.Get());
- TINT_ASSERT(IR, if_node->true_.target->Is<Block>());
- builder_.Branch(if_node->true_.target->As<Block>(), if_node->merge.target,
- utils::Empty);
-
- TINT_ASSERT(IR, if_node->false_.target->Is<Block>());
- builder_.Branch(if_node->false_.target->As<Block>(), loop_node->merge.target,
- utils::Empty);
+ if_node->True().target->As<Block>()->BranchTo(if_node->Merge().target);
+ if_node->False().target->As<Block>()->BranchTo(loop_node->Merge().target);
BranchTo(if_node);
- current_flow_block_ = if_node->merge.target->As<Block>();
+ current_flow_block_ = if_node->Merge().target->As<Block>();
EmitBlock(stmt->body);
- BranchToIfNeeded(loop_node->continuing.target);
+ BranchToIfNeeded(loop_node->Continuing().target);
}
- // The while loop always has a path to the merge target as the break statement comes before
- // anything inside the loop.
- current_flow_block_ = loop_node->merge.target->As<Block>();
+ // The while loop always has a path to the Merge().target as the break statement comes
+ // before anything inside the loop.
+ current_flow_block_ = loop_node->Merge().target->As<Block>();
}
void EmitForLoop(const ast::ForLoopStatement* stmt) {
auto* loop_node = builder_.CreateLoop();
- TINT_ASSERT(IR, loop_node->continuing.target->Is<Block>());
- builder_.Branch(loop_node->continuing.target->As<Block>(), loop_node->start.target,
- utils::Empty);
+ loop_node->Continuing().target->As<Block>()->BranchTo(loop_node->Start().target);
// Make sure the initializer ends up in a contained scope
scopes_.Push();
@@ -666,7 +664,7 @@
{
FlowStackScope scope(this, loop_node);
- current_flow_block_ = loop_node->start.target->As<Block>();
+ current_flow_block_ = loop_node->Start().target->As<Block>();
if (stmt->condition) {
// Emit the condition into the target target of the loop
@@ -677,30 +675,25 @@
// Create an `if (cond) {} else {break;}` control flow
auto* if_node = builder_.CreateIf(reg.Get());
- TINT_ASSERT(IR, if_node->true_.target->Is<Block>());
- builder_.Branch(if_node->true_.target->As<Block>(), if_node->merge.target,
- utils::Empty);
-
- TINT_ASSERT(IR, if_node->false_.target->Is<Block>());
- builder_.Branch(if_node->false_.target->As<Block>(), loop_node->merge.target,
- utils::Empty);
+ if_node->True().target->As<Block>()->BranchTo(if_node->Merge().target);
+ if_node->False().target->As<Block>()->BranchTo(loop_node->Merge().target);
BranchTo(if_node);
- current_flow_block_ = if_node->merge.target->As<Block>();
+ current_flow_block_ = if_node->Merge().target->As<Block>();
}
EmitBlock(stmt->body);
- BranchToIfNeeded(loop_node->continuing.target);
+ BranchToIfNeeded(loop_node->Continuing().target);
if (stmt->continuing) {
- current_flow_block_ = loop_node->continuing.target->As<Block>();
+ current_flow_block_ = loop_node->Continuing().target->As<Block>();
EmitStatement(stmt->continuing);
}
}
- // The while loop always has a path to the merge target as the break statement comes before
- // anything inside the loop.
- current_flow_block_ = loop_node->merge.target->As<Block>();
+ // The while loop always has a path to the Merge().target as the break statement comes
+ // before anything inside the loop.
+ current_flow_block_ = loop_node->Merge().target->As<Block>();
}
void EmitSwitch(const ast::SwitchStatement* stmt) {
@@ -732,13 +725,13 @@
current_flow_block_ = builder_.CreateCase(switch_node, selectors);
EmitBlock(c->Body()->Declaration());
- BranchToIfNeeded(switch_node->merge.target);
+ BranchToIfNeeded(switch_node->Merge().target);
}
}
current_flow_block_ = nullptr;
- if (IsConnected(switch_node->merge.target)) {
- current_flow_block_ = switch_node->merge.target->As<Block>();
+ if (IsConnected(switch_node->Merge().target)) {
+ current_flow_block_ = switch_node->Merge().target->As<Block>();
}
}
@@ -752,7 +745,7 @@
ret_value.Push(ret.Get());
}
- BranchTo(current_function_->end_target, std::move(ret_value));
+ BranchTo(current_function_->EndTarget(), std::move(ret_value));
}
void EmitBreak(const ast::BreakStatement*) {
@@ -760,9 +753,9 @@
TINT_ASSERT(IR, current_control);
if (auto* c = current_control->As<Loop>()) {
- BranchTo(c->merge.target);
+ BranchTo(c->Merge().target);
} else if (auto* s = current_control->As<Switch>()) {
- BranchTo(s->merge.target);
+ BranchTo(s->Merge().target);
} else {
TINT_UNREACHABLE(IR, diagnostics_);
}
@@ -773,7 +766,7 @@
TINT_ASSERT(IR, current_control);
if (auto* c = current_control->As<Loop>()) {
- BranchTo(c->continuing.target);
+ BranchTo(c->Continuing().target);
} else {
TINT_UNREACHABLE(IR, diagnostics_);
}
@@ -785,7 +778,7 @@
// figuring out the multi-level exit that is triggered.
void EmitDiscard(const ast::DiscardStatement*) {
auto* inst = builder_.Discard();
- current_flow_block_->instructions.Push(inst);
+ current_flow_block_->Instructions().Push(inst);
}
void EmitBreakIf(const ast::BreakIfStatement* stmt) {
@@ -806,17 +799,17 @@
auto* loop = current_control->As<Loop>();
- current_flow_block_ = if_node->true_.target->As<Block>();
- BranchTo(loop->merge.target);
+ current_flow_block_ = if_node->True().target->As<Block>();
+ BranchTo(loop->Merge().target);
- current_flow_block_ = if_node->false_.target->As<Block>();
- BranchTo(if_node->merge.target);
+ current_flow_block_ = if_node->False().target->As<Block>();
+ BranchTo(if_node->Merge().target);
- current_flow_block_ = if_node->merge.target->As<Block>();
+ current_flow_block_ = if_node->Merge().target->As<Block>();
// The `break-if` has to be the last item in the continuing block. The false branch of the
// `break-if` will always take us back to the start of the loop.
- BranchTo(loop->start.target);
+ BranchTo(loop->Start().target);
}
utils::Result<Value*> EmitExpression(const ast::Expression* expr) {
@@ -863,7 +856,7 @@
// If this expression maps to sem::Load, insert a load instruction to get the result.
if (result && sem->Is<sem::Load>()) {
auto* load = builder_.Load(result.Get());
- current_flow_block_->instructions.Push(load);
+ current_flow_block_->Instructions().Push(load);
return load;
}
@@ -882,14 +875,14 @@
ref->Access());
auto* val = builder_.Declare(ty);
- current_flow_block_->instructions.Push(val);
+ current_flow_block_->Instructions().Push(val);
if (v->initializer) {
auto init = EmitExpression(v->initializer);
if (!init) {
return;
}
- val->initializer = init.Get();
+ val->SetInitializer(init.Get());
}
// Store the declaration so we can get the instruction to store too
scopes_.Set(v->name->symbol, val);
@@ -956,7 +949,7 @@
break;
}
- current_flow_block_->instructions.Push(inst);
+ current_flow_block_->Instructions().Push(inst);
return inst;
}
@@ -983,7 +976,7 @@
BranchTo(if_node);
auto* result = builder_.BlockParam(builder_.ir.types.Get<type::Bool>());
- if_node->merge.target->As<Block>()->params.Push(result);
+ if_node->Merge().target->As<Block>()->SetParams(utils::Vector{result});
utils::Result<Value*> rhs;
{
@@ -997,17 +990,17 @@
if (expr->op == ast::BinaryOp::kLogicalAnd) {
// If the lhs is false, then that is the result we want to pass to the merge block
// as our argument
- current_flow_block_ = if_node->false_.target->As<Block>();
- BranchTo(if_node->merge.target, std::move(alt_args));
+ current_flow_block_ = if_node->False().target->As<Block>();
+ BranchTo(if_node->Merge().target, std::move(alt_args));
- current_flow_block_ = if_node->true_.target->As<Block>();
+ current_flow_block_ = if_node->True().target->As<Block>();
} else {
// If the lhs is true, then that is the result we want to pass to the merge block
// as our argument
- current_flow_block_ = if_node->true_.target->As<Block>();
- BranchTo(if_node->merge.target, std::move(alt_args));
+ current_flow_block_ = if_node->True().target->As<Block>();
+ BranchTo(if_node->Merge().target, std::move(alt_args));
- current_flow_block_ = if_node->false_.target->As<Block>();
+ current_flow_block_ = if_node->False().target->As<Block>();
}
rhs = EmitExpression(expr->rhs);
@@ -1017,9 +1010,9 @@
utils::Vector<Value*, 1> args;
args.Push(rhs.Get());
- BranchTo(if_node->merge.target, std::move(args));
+ BranchTo(if_node->Merge().target, std::move(args));
}
- current_flow_block_ = if_node->merge.target->As<Block>();
+ current_flow_block_ = if_node->Merge().target->As<Block>();
return result;
}
@@ -1101,7 +1094,7 @@
return utils::Failure;
}
- current_flow_block_->instructions.Push(inst);
+ current_flow_block_->Instructions().Push(inst);
return inst;
}
@@ -1115,7 +1108,7 @@
auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx);
auto* inst = builder_.Bitcast(ty, val.Get());
- current_flow_block_->instructions.Push(inst);
+ current_flow_block_->Instructions().Push(inst);
return inst;
}
@@ -1178,7 +1171,7 @@
if (inst == nullptr) {
return utils::Failure;
}
- current_flow_block_->instructions.Push(inst);
+ current_flow_block_->Instructions().Push(inst);
return inst;
}
diff --git a/src/tint/ir/from_program_binary_test.cc b/src/tint/ir/from_program_binary_test.cc
index 85377ff..b74a8c8 100644
--- a/src/tint/ir/from_program_binary_test.cc
+++ b/src/tint/ir/from_program_binary_test.cc
@@ -830,7 +830,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():bool {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:bool):bool {
%fn2 = block {
} -> %func_end true # return
} %func_end
diff --git a/src/tint/ir/from_program_call_test.cc b/src/tint/ir/from_program_call_test.cc
index 038f4dd..6339c22 100644
--- a/src/tint/ir/from_program_call_test.cc
+++ b/src/tint/ir/from_program_call_test.cc
@@ -77,14 +77,14 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func():void {
+ EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func my_func(%p:f32):void {
%fn2 = block {
} -> %func_end # return
} %func_end
%fn3 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
%fn4 = block {
- %1:void = call my_func, 6.0f
+ %2:void = call my_func, 6.0f
} -> %func_end # return
} %func_end
diff --git a/src/tint/ir/from_program_literal_test.cc b/src/tint/ir/from_program_literal_test.cc
index de5e81e..f746bc0 100644
--- a/src/tint/ir/from_program_literal_test.cc
+++ b/src/tint/ir/from_program_literal_test.cc
@@ -25,17 +25,19 @@
namespace tint::ir {
namespace {
-Value* GlobalVarInitializer(const Module& m) {
- if (m.root_block->instructions.Length() == 0u) {
+const Value* GlobalVarInitializer(const Module& m) {
+ const auto instr = m.root_block->Instructions();
+
+ if (instr.Length() == 0u) {
ADD_FAILURE() << "m.root_block has no instruction";
return nullptr;
}
- auto* var = m.root_block->instructions[0]->As<ir::Var>();
+ auto* var = instr[0]->As<ir::Var>();
if (!var) {
ADD_FAILURE() << "m.root_block.instructions[0] was not a var";
return nullptr;
}
- return var->initializer;
+ return var->Initializer();
}
using namespace tint::number_suffixes; // NOLINT
@@ -51,7 +53,7 @@
auto* init = GlobalVarInitializer(m.Get());
ASSERT_TRUE(Is<Constant>(init));
- auto* val = init->As<Constant>()->value;
+ auto* val = init->As<Constant>()->Value();
EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
EXPECT_TRUE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
@@ -65,11 +67,35 @@
auto* init = GlobalVarInitializer(m.Get());
ASSERT_TRUE(Is<Constant>(init));
- auto* val = init->As<Constant>()->value;
+ auto* val = init->As<Constant>()->Value();
EXPECT_TRUE(val->Is<constant::Scalar<bool>>());
EXPECT_FALSE(val->As<constant::Scalar<bool>>()->ValueAs<bool>());
}
+TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_Deduped) {
+ GlobalVar("a", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(true));
+ GlobalVar("b", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(false));
+ GlobalVar("c", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(true));
+ GlobalVar("d", ty.bool_(), builtin::AddressSpace::kPrivate, Expr(false));
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ auto instr = m.Get().root_block->Instructions();
+ auto* var_a = instr[0]->As<ir::Var>();
+ ASSERT_NE(var_a, nullptr);
+ auto* var_b = instr[1]->As<ir::Var>();
+ ASSERT_NE(var_b, nullptr);
+ auto* var_c = instr[2]->As<ir::Var>();
+ ASSERT_NE(var_c, nullptr);
+ auto* var_d = instr[3]->As<ir::Var>();
+ ASSERT_NE(var_d, nullptr);
+
+ ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+ ASSERT_EQ(var_b->Initializer(), var_d->Initializer());
+ ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
TEST_F(IR_BuilderImplTest, EmitLiteral_F32) {
auto* expr = Expr(1.2_f);
GlobalVar("a", ty.f32(), builtin::AddressSpace::kPrivate, expr);
@@ -79,11 +105,31 @@
auto* init = GlobalVarInitializer(m.Get());
ASSERT_TRUE(Is<Constant>(init));
- auto* val = init->As<Constant>()->value;
+ auto* val = init->As<Constant>()->Value();
EXPECT_TRUE(val->Is<constant::Scalar<f32>>());
EXPECT_EQ(1.2_f, val->As<constant::Scalar<f32>>()->ValueAs<f32>());
}
+TEST_F(IR_BuilderImplTest, EmitLiteral_F32_Deduped) {
+ GlobalVar("a", ty.f32(), builtin::AddressSpace::kPrivate, Expr(1.2_f));
+ GlobalVar("b", ty.f32(), builtin::AddressSpace::kPrivate, Expr(1.25_f));
+ GlobalVar("c", ty.f32(), builtin::AddressSpace::kPrivate, Expr(1.2_f));
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ auto instr = m.Get().root_block->Instructions();
+ auto* var_a = instr[0]->As<ir::Var>();
+ ASSERT_NE(var_a, nullptr);
+ auto* var_b = instr[1]->As<ir::Var>();
+ ASSERT_NE(var_b, nullptr);
+ auto* var_c = instr[2]->As<ir::Var>();
+ ASSERT_NE(var_c, nullptr);
+
+ ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+ ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
TEST_F(IR_BuilderImplTest, EmitLiteral_F16) {
Enable(builtin::Extension::kF16);
auto* expr = Expr(1.2_h);
@@ -94,11 +140,32 @@
auto* init = GlobalVarInitializer(m.Get());
ASSERT_TRUE(Is<Constant>(init));
- auto* val = init->As<Constant>()->value;
+ auto* val = init->As<Constant>()->Value();
EXPECT_TRUE(val->Is<constant::Scalar<f16>>());
EXPECT_EQ(1.2_h, val->As<constant::Scalar<f16>>()->ValueAs<f32>());
}
+TEST_F(IR_BuilderImplTest, EmitLiteral_F16_Deduped) {
+ Enable(builtin::Extension::kF16);
+ GlobalVar("a", ty.f16(), builtin::AddressSpace::kPrivate, Expr(1.2_h));
+ GlobalVar("b", ty.f16(), builtin::AddressSpace::kPrivate, Expr(1.25_h));
+ GlobalVar("c", ty.f16(), builtin::AddressSpace::kPrivate, Expr(1.2_h));
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ auto instr = m.Get().root_block->Instructions();
+ auto* var_a = instr[0]->As<ir::Var>();
+ ASSERT_NE(var_a, nullptr);
+ auto* var_b = instr[1]->As<ir::Var>();
+ ASSERT_NE(var_b, nullptr);
+ auto* var_c = instr[2]->As<ir::Var>();
+ ASSERT_NE(var_c, nullptr);
+
+ ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+ ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
TEST_F(IR_BuilderImplTest, EmitLiteral_I32) {
auto* expr = Expr(-2_i);
GlobalVar("a", ty.i32(), builtin::AddressSpace::kPrivate, expr);
@@ -108,11 +175,31 @@
auto* init = GlobalVarInitializer(m.Get());
ASSERT_TRUE(Is<Constant>(init));
- auto* val = init->As<Constant>()->value;
+ auto* val = init->As<Constant>()->Value();
EXPECT_TRUE(val->Is<constant::Scalar<i32>>());
EXPECT_EQ(-2_i, val->As<constant::Scalar<i32>>()->ValueAs<f32>());
}
+TEST_F(IR_BuilderImplTest, EmitLiteral_I32_Deduped) {
+ GlobalVar("a", ty.i32(), builtin::AddressSpace::kPrivate, Expr(-2_i));
+ GlobalVar("b", ty.i32(), builtin::AddressSpace::kPrivate, Expr(2_i));
+ GlobalVar("c", ty.i32(), builtin::AddressSpace::kPrivate, Expr(-2_i));
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ auto instr = m.Get().root_block->Instructions();
+ auto* var_a = instr[0]->As<ir::Var>();
+ ASSERT_NE(var_a, nullptr);
+ auto* var_b = instr[1]->As<ir::Var>();
+ ASSERT_NE(var_b, nullptr);
+ auto* var_c = instr[2]->As<ir::Var>();
+ ASSERT_NE(var_c, nullptr);
+
+ ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+ ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
TEST_F(IR_BuilderImplTest, EmitLiteral_U32) {
auto* expr = Expr(2_u);
GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, expr);
@@ -122,10 +209,30 @@
auto* init = GlobalVarInitializer(m.Get());
ASSERT_TRUE(Is<Constant>(init));
- auto* val = init->As<Constant>()->value;
+ auto* val = init->As<Constant>()->Value();
EXPECT_TRUE(val->Is<constant::Scalar<u32>>());
EXPECT_EQ(2_u, val->As<constant::Scalar<u32>>()->ValueAs<f32>());
}
+TEST_F(IR_BuilderImplTest, EmitLiteral_U32_Deduped) {
+ GlobalVar("a", ty.u32(), builtin::AddressSpace::kPrivate, Expr(2_u));
+ GlobalVar("b", ty.u32(), builtin::AddressSpace::kPrivate, Expr(3_u));
+ GlobalVar("c", ty.u32(), builtin::AddressSpace::kPrivate, Expr(2_u));
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ auto instr = m.Get().root_block->Instructions();
+ auto* var_a = instr[0]->As<ir::Var>();
+ ASSERT_NE(var_a, nullptr);
+ auto* var_b = instr[1]->As<ir::Var>();
+ ASSERT_NE(var_b, nullptr);
+ auto* var_c = instr[2]->As<ir::Var>();
+ ASSERT_NE(var_c, nullptr);
+
+ ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
+ ASSERT_NE(var_a->Initializer(), var_b->Initializer());
+}
+
} // namespace
} // namespace tint::ir
diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc
index 47dc911..1049386 100644
--- a/src/tint/ir/from_program_test.cc
+++ b/src/tint/ir/from_program_test.cc
@@ -63,13 +63,13 @@
ASSERT_EQ(1u, m->functions.Length());
auto* f = m->functions[0];
- ASSERT_NE(f->start_target, nullptr);
- ASSERT_NE(f->end_target, nullptr);
+ ASSERT_NE(f->StartTarget(), nullptr);
+ ASSERT_NE(f->EndTarget(), nullptr);
- EXPECT_EQ(1u, f->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, f->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
- EXPECT_EQ(m->functions[0]->pipeline_stage, Function::PipelineStage::kUndefined);
+ EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f():void {
%fn2 = block {
@@ -79,6 +79,57 @@
)");
}
+TEST_F(IR_BuilderImplTest, Func_WithParam) {
+ Func("f", utils::Vector{Param("a", ty.u32())}, ty.u32(), utils::Vector{Return("a")});
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ ASSERT_EQ(1u, m->functions.Length());
+
+ auto* f = m->functions[0];
+ ASSERT_NE(f->StartTarget(), nullptr);
+ ASSERT_NE(f->EndTarget(), nullptr);
+
+ EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
+
+ EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
+
+ EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32):u32 {
+ %fn2 = block {
+ } -> %func_end %a # return
+} %func_end
+
+)");
+}
+
+TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) {
+ Func("f", utils::Vector{Param("a", ty.u32()), Param("b", ty.i32()), Param("c", ty.bool_())},
+ ty.void_(), utils::Empty);
+
+ auto m = Build();
+ ASSERT_TRUE(m) << (!m ? m.Failure() : "");
+
+ ASSERT_EQ(1u, m->functions.Length());
+
+ auto* f = m->functions[0];
+ ASSERT_NE(f->StartTarget(), nullptr);
+ ASSERT_NE(f->EndTarget(), nullptr);
+
+ EXPECT_EQ(1u, f->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, f->EndTarget()->InboundBranches().Length());
+
+ EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kUndefined);
+
+ EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func f(%a:u32, %b:i32, %c:bool):void {
+ %fn2 = block {
+ } -> %func_end # return
+} %func_end
+
+)");
+}
+
TEST_F(IR_BuilderImplTest, EntryPoint) {
Func("f", utils::Empty, ty.void_(), utils::Empty,
utils::Vector{Stage(ast::PipelineStage::kFragment)});
@@ -86,7 +137,7 @@
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
- EXPECT_EQ(m->functions[0]->pipeline_stage, Function::PipelineStage::kFragment);
+ EXPECT_EQ(m->functions[0]->Stage(), Function::PipelineStage::kFragment);
}
TEST_F(IR_BuilderImplTest, IfStatement) {
@@ -97,19 +148,19 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(flow->true_.target, nullptr);
- ASSERT_NE(flow->false_.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->True().target, nullptr);
+ ASSERT_NE(flow->False().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(2u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -141,19 +192,19 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(flow->true_.target, nullptr);
- ASSERT_NE(flow->false_.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->True().target, nullptr);
+ ASSERT_NE(flow->False().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -184,19 +235,19 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(flow->true_.target, nullptr);
- ASSERT_NE(flow->false_.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->True().target, nullptr);
+ ASSERT_NE(flow->False().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -227,19 +278,19 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(flow->true_.target, nullptr);
- ASSERT_NE(flow->false_.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->True().target, nullptr);
+ ASSERT_NE(flow->False().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -267,15 +318,15 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(if_flow->true_.target, nullptr);
- ASSERT_NE(if_flow->false_.target, nullptr);
- ASSERT_NE(if_flow->merge.target, nullptr);
+ ASSERT_NE(if_flow->True().target, nullptr);
+ ASSERT_NE(if_flow->False().target, nullptr);
+ ASSERT_NE(if_flow->Merge().target, nullptr);
auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
ASSERT_NE(loop_flow, nullptr);
- ASSERT_NE(loop_flow->start.target, nullptr);
- ASSERT_NE(loop_flow->continuing.target, nullptr);
- ASSERT_NE(loop_flow->merge.target, nullptr);
+ ASSERT_NE(loop_flow->Start().target, nullptr);
+ ASSERT_NE(loop_flow->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow->Merge().target, nullptr);
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -316,19 +367,19 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(flow->start.target, nullptr);
- ASSERT_NE(flow->continuing.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target, nullptr);
+ ASSERT_NE(flow->Continuing().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
- EXPECT_EQ(0u, flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -357,28 +408,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(loop_flow->start.target, nullptr);
- ASSERT_NE(loop_flow->continuing.target, nullptr);
- ASSERT_NE(loop_flow->merge.target, nullptr);
+ ASSERT_NE(loop_flow->Start().target, nullptr);
+ ASSERT_NE(loop_flow->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow->Merge().target, nullptr);
auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(if_flow->true_.target, nullptr);
- ASSERT_NE(if_flow->false_.target, nullptr);
- ASSERT_NE(if_flow->merge.target, nullptr);
+ ASSERT_NE(if_flow->True().target, nullptr);
+ ASSERT_NE(if_flow->False().target, nullptr);
+ ASSERT_NE(if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -424,28 +475,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(loop_flow->start.target, nullptr);
- ASSERT_NE(loop_flow->continuing.target, nullptr);
- ASSERT_NE(loop_flow->merge.target, nullptr);
+ ASSERT_NE(loop_flow->Start().target, nullptr);
+ ASSERT_NE(loop_flow->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow->Merge().target, nullptr);
auto* break_if_flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(break_if_flow->true_.target, nullptr);
- ASSERT_NE(break_if_flow->false_.target, nullptr);
- ASSERT_NE(break_if_flow->merge.target, nullptr);
+ ASSERT_NE(break_if_flow->True().target, nullptr);
+ ASSERT_NE(break_if_flow->False().target, nullptr);
+ ASSERT_NE(break_if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, break_if_flow->inbound_branches.Length());
- EXPECT_EQ(1u, break_if_flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, break_if_flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, break_if_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, break_if_flow->InboundBranches().Length());
+ EXPECT_EQ(1u, break_if_flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, break_if_flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, break_if_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -535,28 +586,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(loop_flow->start.target, nullptr);
- ASSERT_NE(loop_flow->continuing.target, nullptr);
- ASSERT_NE(loop_flow->merge.target, nullptr);
+ ASSERT_NE(loop_flow->Start().target, nullptr);
+ ASSERT_NE(loop_flow->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow->Merge().target, nullptr);
auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(if_flow->true_.target, nullptr);
- ASSERT_NE(if_flow->false_.target, nullptr);
- ASSERT_NE(if_flow->merge.target, nullptr);
+ ASSERT_NE(if_flow->True().target, nullptr);
+ ASSERT_NE(if_flow->False().target, nullptr);
+ ASSERT_NE(if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(0u, loop_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -597,19 +648,19 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(loop_flow->start.target, nullptr);
- ASSERT_NE(loop_flow->continuing.target, nullptr);
- ASSERT_NE(loop_flow->merge.target, nullptr);
+ ASSERT_NE(loop_flow->Start().target, nullptr);
+ ASSERT_NE(loop_flow->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
- EXPECT_EQ(0u, loop_flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(0u, loop_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -642,25 +693,25 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(loop_flow->start.target, nullptr);
- ASSERT_NE(loop_flow->continuing.target, nullptr);
- ASSERT_NE(loop_flow->merge.target, nullptr);
+ ASSERT_NE(loop_flow->Start().target, nullptr);
+ ASSERT_NE(loop_flow->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow->Merge().target, nullptr);
auto* break_if_flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(break_if_flow->true_.target, nullptr);
- ASSERT_NE(break_if_flow->false_.target, nullptr);
- ASSERT_NE(break_if_flow->merge.target, nullptr);
+ ASSERT_NE(break_if_flow->True().target, nullptr);
+ ASSERT_NE(break_if_flow->False().target, nullptr);
+ ASSERT_NE(break_if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
- EXPECT_EQ(0u, loop_flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
+ EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
// This is 1 because only the loop branch happens. The subsequent if return is dead code.
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -685,28 +736,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* loop_flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(loop_flow->start.target, nullptr);
- ASSERT_NE(loop_flow->continuing.target, nullptr);
- ASSERT_NE(loop_flow->merge.target, nullptr);
+ ASSERT_NE(loop_flow->Start().target, nullptr);
+ ASSERT_NE(loop_flow->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow->Merge().target, nullptr);
auto* if_flow = FindSingleFlowNode<ir::If>(m.Get());
- ASSERT_NE(if_flow->true_.target, nullptr);
- ASSERT_NE(if_flow->false_.target, nullptr);
- ASSERT_NE(if_flow->merge.target, nullptr);
+ ASSERT_NE(if_flow->True().target, nullptr);
+ ASSERT_NE(if_flow->False().target, nullptr);
+ ASSERT_NE(if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, loop_flow->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow->start.target->inbound_branches.Length());
- EXPECT_EQ(0u, loop_flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(0u, if_flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, loop_flow->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(0u, if_flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -756,96 +807,96 @@
auto block_exit = [&](const ir::FlowNode* node) -> const ir::FlowNode* {
if (auto* block = As<ir::Block>(node)) {
- return block->branch.target;
+ return block->Branch().target;
}
return nullptr;
};
- auto* loop_flow_a = As<ir::Loop>(m->functions[0]->start_target->branch.target);
+ auto* loop_flow_a = As<ir::Loop>(m->functions[0]->StartTarget()->Branch().target);
ASSERT_NE(loop_flow_a, nullptr);
- ASSERT_NE(loop_flow_a->start.target, nullptr);
- ASSERT_NE(loop_flow_a->continuing.target, nullptr);
- ASSERT_NE(loop_flow_a->merge.target, nullptr);
+ ASSERT_NE(loop_flow_a->Start().target, nullptr);
+ ASSERT_NE(loop_flow_a->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow_a->Merge().target, nullptr);
- auto* loop_flow_b = As<ir::Loop>(block_exit(loop_flow_a->start.target));
+ auto* loop_flow_b = As<ir::Loop>(block_exit(loop_flow_a->Start().target));
ASSERT_NE(loop_flow_b, nullptr);
- ASSERT_NE(loop_flow_b->start.target, nullptr);
- ASSERT_NE(loop_flow_b->continuing.target, nullptr);
- ASSERT_NE(loop_flow_b->merge.target, nullptr);
+ ASSERT_NE(loop_flow_b->Start().target, nullptr);
+ ASSERT_NE(loop_flow_b->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow_b->Merge().target, nullptr);
- auto* if_flow_a = As<ir::If>(block_exit(loop_flow_b->start.target));
+ auto* if_flow_a = As<ir::If>(block_exit(loop_flow_b->Start().target));
ASSERT_NE(if_flow_a, nullptr);
- ASSERT_NE(if_flow_a->true_.target, nullptr);
- ASSERT_NE(if_flow_a->false_.target, nullptr);
- ASSERT_NE(if_flow_a->merge.target, nullptr);
+ ASSERT_NE(if_flow_a->True().target, nullptr);
+ ASSERT_NE(if_flow_a->False().target, nullptr);
+ ASSERT_NE(if_flow_a->Merge().target, nullptr);
- auto* if_flow_b = As<ir::If>(block_exit(if_flow_a->merge.target));
+ auto* if_flow_b = As<ir::If>(block_exit(if_flow_a->Merge().target));
ASSERT_NE(if_flow_b, nullptr);
- ASSERT_NE(if_flow_b->true_.target, nullptr);
- ASSERT_NE(if_flow_b->false_.target, nullptr);
- ASSERT_NE(if_flow_b->merge.target, nullptr);
+ ASSERT_NE(if_flow_b->True().target, nullptr);
+ ASSERT_NE(if_flow_b->False().target, nullptr);
+ ASSERT_NE(if_flow_b->Merge().target, nullptr);
- auto* loop_flow_c = As<ir::Loop>(block_exit(loop_flow_b->continuing.target));
+ auto* loop_flow_c = As<ir::Loop>(block_exit(loop_flow_b->Continuing().target));
ASSERT_NE(loop_flow_c, nullptr);
- ASSERT_NE(loop_flow_c->start.target, nullptr);
- ASSERT_NE(loop_flow_c->continuing.target, nullptr);
- ASSERT_NE(loop_flow_c->merge.target, nullptr);
+ ASSERT_NE(loop_flow_c->Start().target, nullptr);
+ ASSERT_NE(loop_flow_c->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow_c->Merge().target, nullptr);
- auto* loop_flow_d = As<ir::Loop>(block_exit(loop_flow_c->merge.target));
+ auto* loop_flow_d = As<ir::Loop>(block_exit(loop_flow_c->Merge().target));
ASSERT_NE(loop_flow_d, nullptr);
- ASSERT_NE(loop_flow_d->start.target, nullptr);
- ASSERT_NE(loop_flow_d->continuing.target, nullptr);
- ASSERT_NE(loop_flow_d->merge.target, nullptr);
+ ASSERT_NE(loop_flow_d->Start().target, nullptr);
+ ASSERT_NE(loop_flow_d->Continuing().target, nullptr);
+ ASSERT_NE(loop_flow_d->Merge().target, nullptr);
- auto* if_flow_c = As<ir::If>(block_exit(loop_flow_d->continuing.target));
+ auto* if_flow_c = As<ir::If>(block_exit(loop_flow_d->Continuing().target));
ASSERT_NE(if_flow_c, nullptr);
- ASSERT_NE(if_flow_c->true_.target, nullptr);
- ASSERT_NE(if_flow_c->false_.target, nullptr);
- ASSERT_NE(if_flow_c->merge.target, nullptr);
+ ASSERT_NE(if_flow_c->True().target, nullptr);
+ ASSERT_NE(if_flow_c->False().target, nullptr);
+ ASSERT_NE(if_flow_c->Merge().target, nullptr);
- auto* if_flow_d = As<ir::If>(block_exit(loop_flow_b->merge.target));
+ auto* if_flow_d = As<ir::If>(block_exit(loop_flow_b->Merge().target));
ASSERT_NE(if_flow_d, nullptr);
- ASSERT_NE(if_flow_d->true_.target, nullptr);
- ASSERT_NE(if_flow_d->false_.target, nullptr);
- ASSERT_NE(if_flow_d->merge.target, nullptr);
+ ASSERT_NE(if_flow_d->True().target, nullptr);
+ ASSERT_NE(if_flow_d->False().target, nullptr);
+ ASSERT_NE(if_flow_d->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, loop_flow_a->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow_a->start.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_a->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_a->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_b->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow_b->start.target->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow_b->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_b->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_c->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow_c->start.target->inbound_branches.Length());
- EXPECT_EQ(0u, loop_flow_c->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_c->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_d->inbound_branches.Length());
- EXPECT_EQ(2u, loop_flow_d->start.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_d->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, loop_flow_d->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_a->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_a->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_a->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_a->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_b->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_b->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_b->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_b->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_c->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_c->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_c->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_c->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_d->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_d->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_d->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow_d->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->start_target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, loop_flow_a->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow_a->Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_a->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_a->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_b->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow_b->Start().target->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow_b->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_b->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_c->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow_c->Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, loop_flow_c->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_c->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_d->InboundBranches().Length());
+ EXPECT_EQ(2u, loop_flow_d->Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_d->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, loop_flow_d->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_a->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_a->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_a->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_a->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_b->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_b->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_b->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_b->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_c->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_c->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_c->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_c->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_d->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_d->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_d->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow_d->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->StartTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -964,28 +1015,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(flow->start.target, nullptr);
- ASSERT_NE(flow->continuing.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target, nullptr);
+ ASSERT_NE(flow->Continuing().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
- ASSERT_NE(flow->start.target->As<ir::Block>()->branch.target, nullptr);
- ASSERT_TRUE(flow->start.target->As<ir::Block>()->branch.target->Is<ir::If>());
- auto* if_flow = flow->start.target->As<ir::Block>()->branch.target->As<ir::If>();
- ASSERT_NE(if_flow->true_.target, nullptr);
- ASSERT_NE(if_flow->false_.target, nullptr);
- ASSERT_NE(if_flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target->As<ir::Block>()->Branch().target, nullptr);
+ ASSERT_TRUE(flow->Start().target->As<ir::Block>()->Branch().target->Is<ir::If>());
+ auto* if_flow = flow->Start().target->As<ir::Block>()->Branch().target->As<ir::If>();
+ ASSERT_NE(if_flow->True().target, nullptr);
+ ASSERT_NE(if_flow->False().target, nullptr);
+ ASSERT_NE(if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1030,28 +1081,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(flow->start.target, nullptr);
- ASSERT_NE(flow->continuing.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target, nullptr);
+ ASSERT_NE(flow->Continuing().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
- ASSERT_NE(flow->start.target->As<ir::Block>()->branch.target, nullptr);
- ASSERT_TRUE(flow->start.target->As<ir::Block>()->branch.target->Is<ir::If>());
- auto* if_flow = flow->start.target->As<ir::Block>()->branch.target->As<ir::If>();
- ASSERT_NE(if_flow->true_.target, nullptr);
- ASSERT_NE(if_flow->false_.target, nullptr);
- ASSERT_NE(if_flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target->As<ir::Block>()->Branch().target, nullptr);
+ ASSERT_TRUE(flow->Start().target->As<ir::Block>()->Branch().target->Is<ir::If>());
+ auto* if_flow = flow->Start().target->As<ir::Block>()->Branch().target->As<ir::If>();
+ ASSERT_NE(if_flow->True().target, nullptr);
+ ASSERT_NE(if_flow->False().target, nullptr);
+ ASSERT_NE(if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
- EXPECT_EQ(0u, flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
+ EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1104,28 +1155,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(flow->start.target, nullptr);
- ASSERT_NE(flow->continuing.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target, nullptr);
+ ASSERT_NE(flow->Continuing().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
- ASSERT_NE(flow->start.target->As<ir::Block>()->branch.target, nullptr);
- ASSERT_TRUE(flow->start.target->As<ir::Block>()->branch.target->Is<ir::If>());
- auto* if_flow = flow->start.target->As<ir::Block>()->branch.target->As<ir::If>();
- ASSERT_NE(if_flow->true_.target, nullptr);
- ASSERT_NE(if_flow->false_.target, nullptr);
- ASSERT_NE(if_flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target->As<ir::Block>()->Branch().target, nullptr);
+ ASSERT_TRUE(flow->Start().target->As<ir::Block>()->Branch().target->Is<ir::If>());
+ auto* if_flow = flow->Start().target->As<ir::Block>()->Branch().target->As<ir::If>();
+ ASSERT_NE(if_flow->True().target, nullptr);
+ ASSERT_NE(if_flow->False().target, nullptr);
+ ASSERT_NE(if_flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->true_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->false_.target->inbound_branches.Length());
- EXPECT_EQ(1u, if_flow->merge.target->inbound_branches.Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->True().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->False().target->InboundBranches().Length());
+ EXPECT_EQ(1u, if_flow->Merge().target->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()), R"()");
}
@@ -1138,18 +1189,18 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Loop>(m.Get());
- ASSERT_NE(flow->start.target, nullptr);
- ASSERT_NE(flow->continuing.target, nullptr);
- ASSERT_NE(flow->merge.target, nullptr);
+ ASSERT_NE(flow->Start().target, nullptr);
+ ASSERT_NE(flow->Continuing().target, nullptr);
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(2u, flow->start.target->inbound_branches.Length());
- EXPECT_EQ(0u, flow->continuing.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Continuing().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1180,31 +1231,33 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
- ASSERT_NE(flow->merge.target, nullptr);
- ASSERT_EQ(3u, flow->cases.Length());
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- ASSERT_EQ(1u, flow->cases[0].selectors.Length());
- ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+ auto cases = flow->Cases();
+ ASSERT_EQ(3u, cases.Length());
+
+ ASSERT_EQ(1u, cases[0].selectors.Length());
+ ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
EXPECT_EQ(0_i,
- flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+ cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
- ASSERT_EQ(1u, flow->cases[1].selectors.Length());
- ASSERT_TRUE(flow->cases[1].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+ ASSERT_EQ(1u, cases[1].selectors.Length());
+ ASSERT_TRUE(cases[1].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
EXPECT_EQ(1_i,
- flow->cases[1].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+ cases[1].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
- ASSERT_EQ(1u, flow->cases[2].selectors.Length());
- EXPECT_TRUE(flow->cases[2].selectors[0].IsDefault());
+ ASSERT_EQ(1u, cases[2].selectors.Length());
+ EXPECT_TRUE(cases[2].selectors[0].IsDefault());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[1].start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[2].start.target->inbound_branches.Length());
- EXPECT_EQ(3u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[2].Start().target->InboundBranches().Length());
+ EXPECT_EQ(3u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1244,27 +1297,28 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
- ASSERT_NE(flow->merge.target, nullptr);
- ASSERT_EQ(1u, flow->cases.Length());
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- ASSERT_EQ(3u, flow->cases[0].selectors.Length());
- ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+ auto cases = flow->Cases();
+ ASSERT_EQ(1u, cases.Length());
+ ASSERT_EQ(3u, cases[0].selectors.Length());
+ ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
EXPECT_EQ(0_i,
- flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+ cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
- ASSERT_TRUE(flow->cases[0].selectors[1].val->value->Is<constant::Scalar<tint::i32>>());
+ ASSERT_TRUE(cases[0].selectors[1].val->Value()->Is<constant::Scalar<tint::i32>>());
EXPECT_EQ(1_i,
- flow->cases[0].selectors[1].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+ cases[0].selectors[1].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
- EXPECT_TRUE(flow->cases[0].selectors[2].IsDefault());
+ EXPECT_TRUE(cases[0].selectors[2].IsDefault());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1292,19 +1346,20 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
- ASSERT_NE(flow->merge.target, nullptr);
- ASSERT_EQ(1u, flow->cases.Length());
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- ASSERT_EQ(1u, flow->cases[0].selectors.Length());
- EXPECT_TRUE(flow->cases[0].selectors[0].IsDefault());
+ auto cases = flow->Cases();
+ ASSERT_EQ(1u, cases.Length());
+ ASSERT_EQ(1u, cases[0].selectors.Length());
+ EXPECT_TRUE(cases[0].selectors[0].IsDefault());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1334,26 +1389,27 @@
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
- ASSERT_NE(flow->merge.target, nullptr);
- ASSERT_EQ(2u, flow->cases.Length());
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- ASSERT_EQ(1u, flow->cases[0].selectors.Length());
- ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+ auto cases = flow->Cases();
+ ASSERT_EQ(2u, cases.Length());
+ ASSERT_EQ(1u, cases[0].selectors.Length());
+ ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
EXPECT_EQ(0_i,
- flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+ cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
- ASSERT_EQ(1u, flow->cases[1].selectors.Length());
- EXPECT_TRUE(flow->cases[1].selectors[0].IsDefault());
+ ASSERT_EQ(1u, cases[1].selectors.Length());
+ EXPECT_TRUE(cases[1].selectors[0].IsDefault());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[1].start.target->inbound_branches.Length());
- EXPECT_EQ(2u, flow->merge.target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length());
+ EXPECT_EQ(2u, flow->Merge().target->InboundBranches().Length());
// This is 1 because the if is dead-code eliminated and the return doesn't happen.
- EXPECT_EQ(1u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
@@ -1390,25 +1446,26 @@
ASSERT_EQ(FindSingleFlowNode<ir::If>(m.Get()), nullptr);
auto* flow = FindSingleFlowNode<ir::Switch>(m.Get());
- ASSERT_NE(flow->merge.target, nullptr);
- ASSERT_EQ(2u, flow->cases.Length());
+ ASSERT_NE(flow->Merge().target, nullptr);
ASSERT_EQ(1u, m->functions.Length());
auto* func = m->functions[0];
- ASSERT_EQ(1u, flow->cases[0].selectors.Length());
- ASSERT_TRUE(flow->cases[0].selectors[0].val->value->Is<constant::Scalar<tint::i32>>());
+ auto cases = flow->Cases();
+ ASSERT_EQ(2u, cases.Length());
+ ASSERT_EQ(1u, cases[0].selectors.Length());
+ ASSERT_TRUE(cases[0].selectors[0].val->Value()->Is<constant::Scalar<tint::i32>>());
EXPECT_EQ(0_i,
- flow->cases[0].selectors[0].val->value->As<constant::Scalar<tint::i32>>()->ValueOf());
+ cases[0].selectors[0].val->Value()->As<constant::Scalar<tint::i32>>()->ValueOf());
- ASSERT_EQ(1u, flow->cases[1].selectors.Length());
- EXPECT_TRUE(flow->cases[1].selectors[0].IsDefault());
+ ASSERT_EQ(1u, cases[1].selectors.Length());
+ EXPECT_TRUE(cases[1].selectors[0].IsDefault());
- EXPECT_EQ(1u, flow->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length());
- EXPECT_EQ(1u, flow->cases[1].start.target->inbound_branches.Length());
- EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length());
- EXPECT_EQ(2u, func->end_target->inbound_branches.Length());
+ EXPECT_EQ(1u, flow->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length());
+ EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length());
+ EXPECT_EQ(0u, flow->Merge().target->InboundBranches().Length());
+ EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length());
EXPECT_EQ(Disassemble(m.Get()),
R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] {
diff --git a/src/tint/ir/function.cc b/src/tint/ir/function.cc
index 2a5f8c3..d6cd6fc 100644
--- a/src/tint/ir/function.cc
+++ b/src/tint/ir/function.cc
@@ -18,11 +18,11 @@
namespace tint::ir {
-Function::Function(Symbol n,
+Function::Function(Symbol name,
type::Type* rt,
PipelineStage stage,
std::optional<std::array<uint32_t, 3>> wg_size)
- : Base(), name(n), pipeline_stage(stage), workgroup_size(wg_size), return_type(rt) {}
+ : Base(), name_(name), return_type_(rt), pipeline_stage_(stage), workgroup_size_(wg_size) {}
Function::~Function() = default;
diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h
index 95487b9..86465aa 100644
--- a/src/tint/ir/function.h
+++ b/src/tint/ir/function.h
@@ -17,8 +17,10 @@
#include <array>
#include <optional>
+#include <utility>
#include "src/tint/ir/flow_node.h"
+#include "src/tint/ir/function_param.h"
#include "src/tint/symbol.h"
#include "src/tint/type/type.h"
@@ -70,29 +72,81 @@
type::Type* rt,
PipelineStage stage = PipelineStage::kUndefined,
std::optional<std::array<uint32_t, 3>> wg_size = {});
+ Function(Function&&) = delete;
+ Function(const Function&) = delete;
~Function() override;
- /// The function name
- Symbol name;
+ Function& operator=(Function&&) = delete;
+ Function& operator=(const Function&) = delete;
- /// The pipeline stage for the function, `kUndefined` if the function is not an entry point
- PipelineStage pipeline_stage = PipelineStage::kUndefined;
+ /// @returns the function name
+ Symbol Name() const { return name_; }
- /// If this is a `compute` entry point, holds the workgroup size information
- std::optional<std::array<uint32_t, 3>> workgroup_size;
+ /// Sets the function stage
+ /// @param stage the stage to set
+ void SetStage(PipelineStage stage) { pipeline_stage_ = stage; }
- /// The function return type
- const type::Type* return_type = nullptr;
- /// The function return attributes if any
- utils::Vector<ReturnAttribute, 1> return_attributes;
- /// If the return attribute is `kLocation` this stores the location value.
- std::optional<uint32_t> return_location;
+ /// @returns the function pipeline stage
+ PipelineStage Stage() const { return pipeline_stage_; }
- /// The start target is the first block in a function.
- Block* start_target = nullptr;
- /// The end target is the end of the function. It is used as the branch target if a return is
- /// encountered in the function.
- FunctionTerminator* end_target = nullptr;
+ /// Sets the workgroup size
+ /// @param x the x size
+ /// @param y the y size
+ /// @param z the z size
+ void SetWorkgroupSize(uint32_t x, uint32_t y, uint32_t z) { workgroup_size_ = {x, y, z}; }
+
+ /// @returns the workgroup size information
+ std::optional<std::array<uint32_t, 3>> WorkgroupSize() const { return workgroup_size_; }
+
+ /// @returns the return type for the function
+ const type::Type* ReturnType() const { return return_type_; }
+
+ /// Sets the return attributes
+ /// @param attrs the attributes to set
+ void SetReturnAttributes(utils::VectorRef<ReturnAttribute> attrs) {
+ return_attributes_ = std::move(attrs);
+ }
+ /// @returns the return attributes
+ utils::VectorRef<ReturnAttribute> ReturnAttributes() const { return return_attributes_; }
+
+ /// Sets the return location
+ /// @param loc the location to set
+ void SetReturnLocation(std::optional<uint32_t> loc) { return_location_ = loc; }
+ /// @returns the return location
+ std::optional<uint32_t> ReturnLocation() const { return return_location_; }
+
+ /// Sets the function parameters
+ /// @param params the function paramters
+ void SetParams(utils::VectorRef<FunctionParam*> params) { params_ = std::move(params); }
+
+ /// @returns the function parameters
+ utils::VectorRef<FunctionParam*> Params() const { return params_; }
+
+ /// Sets the start target for the function
+ /// @param target the start target
+ void SetStartTarget(Block* target) { start_target_ = target; }
+ /// @returns the function start target
+ Block* StartTarget() const { return start_target_; }
+
+ /// Sets the end target for the function
+ /// @param target the end target
+ void SetEndTarget(FunctionTerminator* target) { end_target_ = target; }
+ /// @returns the function end target
+ FunctionTerminator* EndTarget() const { return end_target_; }
+
+ private:
+ Symbol name_;
+ const type::Type* return_type_;
+ PipelineStage pipeline_stage_;
+ std::optional<std::array<uint32_t, 3>> workgroup_size_;
+
+ utils::Vector<ReturnAttribute, 1> return_attributes_;
+ std::optional<uint32_t> return_location_;
+
+ utils::Vector<FunctionParam*, 1> params_;
+
+ Block* start_target_ = nullptr;
+ FunctionTerminator* end_target_ = nullptr;
};
utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value);
diff --git a/src/tint/ir/function_param.cc b/src/tint/ir/function_param.cc
new file mode 100644
index 0000000..a9570fa
--- /dev/null
+++ b/src/tint/ir/function_param.cc
@@ -0,0 +1,25 @@
+// 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/ir/function_param.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::FunctionParam);
+
+namespace tint::ir {
+
+FunctionParam::FunctionParam(const type::Type* ty) : type_(ty) {}
+
+FunctionParam::~FunctionParam() = default;
+
+} // namespace tint::ir
diff --git a/src/tint/ir/function_param.h b/src/tint/ir/function_param.h
new file mode 100644
index 0000000..1bbb812
--- /dev/null
+++ b/src/tint/ir/function_param.h
@@ -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.
+
+#ifndef SRC_TINT_IR_FUNCTION_PARAM_H_
+#define SRC_TINT_IR_FUNCTION_PARAM_H_
+
+#include "src/tint/ir/value.h"
+#include "src/tint/utils/castable.h"
+
+namespace tint::ir {
+
+/// A function parameter in the IR.
+class FunctionParam : public utils::Castable<FunctionParam, Value> {
+ public:
+ /// Constructor
+ /// @param type the type of the var
+ explicit FunctionParam(const type::Type* type);
+ FunctionParam(const FunctionParam& inst) = delete;
+ FunctionParam(FunctionParam&& inst) = delete;
+ ~FunctionParam() override;
+
+ FunctionParam& operator=(const FunctionParam& inst) = delete;
+ FunctionParam& operator=(FunctionParam&& inst) = delete;
+
+ /// @returns the type of the var
+ const type::Type* Type() const override { return type_; }
+
+ private:
+ /// The type of the parameter
+ const type::Type* type_;
+};
+
+} // namespace tint::ir
+
+#endif // SRC_TINT_IR_FUNCTION_PARAM_H_
diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc
index b59d87f..d235b5e 100644
--- a/src/tint/ir/if.cc
+++ b/src/tint/ir/if.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-If::If(Value* cond) : Base(), condition(cond) {}
+If::If(Value* cond) : Base(), condition_(cond) {}
If::~If() = default;
diff --git a/src/tint/ir/if.h b/src/tint/ir/if.h
index aadc5c9..d02c32a 100644
--- a/src/tint/ir/if.h
+++ b/src/tint/ir/if.h
@@ -32,17 +32,36 @@
/// Constructor
/// @param cond the if condition
explicit If(Value* cond);
+ If(const If&) = delete;
+ If(If&&) = delete;
~If() override;
- /// The true branch block
+ If& operator=(const If&) = delete;
+ If& operator=(If&&) = delete;
+
+ /// @returns the if condition
+ const Value* Condition() const { return condition_; }
+
+ /// @returns the true branch block
+ const Branch& True() const { return true_; }
+ /// @returns the true branch block
+ Branch& True() { return true_; }
+
+ /// @returns the false branch block
+ const Branch& False() const { return false_; }
+ /// @returns the false branch block
+ Branch& False() { return false_; }
+
+ /// @returns the merge branch block
+ const Branch& Merge() const { return merge_; }
+ /// @returns the merge branch block
+ Branch& Merge() { return merge_; }
+
+ private:
Branch true_ = {};
- /// The false branch block
Branch false_ = {};
- /// An block to converge the true/false branches. The block always exists, but there maybe no
- /// branches into it. (e.g. if both branches `return`)
- Branch merge = {};
- /// Value holding the condition result
- const Value* condition = nullptr;
+ Branch merge_ = {};
+ Value* condition_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/load.cc b/src/tint/ir/load.cc
index 1fe55c0..3b64cdd 100644
--- a/src/tint/ir/load.cc
+++ b/src/tint/ir/load.cc
@@ -19,10 +19,10 @@
namespace tint::ir {
-Load::Load(const type::Type* type, Value* f) : Base(), result_type(type), from(f) {
- TINT_ASSERT(IR, result_type);
- TINT_ASSERT(IR, from);
- from->AddUsage(this);
+Load::Load(const type::Type* type, Value* f) : Base(), result_type_(type), from_(f) {
+ TINT_ASSERT(IR, result_type_);
+ TINT_ASSERT(IR, from_);
+ from_->AddUsage(this);
}
Load::~Load() = default;
diff --git a/src/tint/ir/load.h b/src/tint/ir/load.h
index b15eced..e1a365f 100644
--- a/src/tint/ir/load.h
+++ b/src/tint/ir/load.h
@@ -35,13 +35,14 @@
Load& operator=(Load&& inst) = delete;
/// @returns the type of the value
- const type::Type* Type() const override { return result_type; }
+ const type::Type* Type() const override { return result_type_; }
- /// the result type of the instruction
- const type::Type* result_type = nullptr;
+ /// @returns the avlue being loaded from
+ Value* From() const { return from_; }
- /// the value being loaded
- Value* from = nullptr;
+ private:
+ const type::Type* result_type_;
+ Value* from_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/load_test.cc b/src/tint/ir/load_test.cc
index 2c6e5c1..9881e98 100644
--- a/src/tint/ir/load_test.cc
+++ b/src/tint/ir/load_test.cc
@@ -33,12 +33,12 @@
const auto* inst = b.Load(var);
ASSERT_TRUE(inst->Is<Load>());
- ASSERT_EQ(inst->from, var);
+ ASSERT_EQ(inst->From(), var);
EXPECT_EQ(inst->Type(), store_type);
- ASSERT_TRUE(inst->from->Is<ir::Var>());
- EXPECT_EQ(inst->from, var);
+ ASSERT_TRUE(inst->From()->Is<ir::Var>());
+ EXPECT_EQ(inst->From(), var);
}
TEST_F(IR_InstructionTest, Load_Usage) {
@@ -50,9 +50,9 @@
store_type, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite));
const auto* inst = b.Load(var);
- ASSERT_NE(inst->from, nullptr);
- ASSERT_EQ(inst->from->Usage().Length(), 1u);
- EXPECT_EQ(inst->from->Usage()[0], inst);
+ ASSERT_NE(inst->From(), nullptr);
+ ASSERT_EQ(inst->From()->Usage().Length(), 1u);
+ EXPECT_EQ(inst->From()->Usage()[0], inst);
}
} // namespace
diff --git a/src/tint/ir/loop.h b/src/tint/ir/loop.h
index e0066f4..590f49c 100644
--- a/src/tint/ir/loop.h
+++ b/src/tint/ir/loop.h
@@ -26,16 +26,32 @@
public:
/// Constructor
Loop();
+ Loop(const Loop&) = delete;
+ Loop(Loop&&) = delete;
~Loop() override;
- /// The start block is the first block in a loop.
- Branch start = {};
- /// The continue target of the block.
- Branch continuing = {};
- /// The loop merge target. If the `loop` does a `return` then this block may not actually
- /// end up in the control flow. We need it if the loop does a `break` we know where to break
- /// too.
- Branch merge = {};
+ Loop& operator=(const Loop&) = delete;
+ Loop& operator=(Loop&&) = delete;
+
+ /// @returns the switch start branch
+ const Branch& Start() const { return start_; }
+ /// @returns the switch start branch
+ Branch& Start() { return start_; }
+
+ /// @returns the switch continuing branch
+ const Branch& Continuing() const { return continuing_; }
+ /// @returns the switch continuing branch
+ Branch& Continuing() { return continuing_; }
+
+ /// @returns the switch merge branch
+ const Branch& Merge() const { return merge_; }
+ /// @returns the switch merge branch
+ Branch& Merge() { return merge_; }
+
+ private:
+ Branch start_ = {};
+ Branch continuing_ = {};
+ Branch merge_ = {};
};
} // namespace tint::ir
diff --git a/src/tint/ir/module.h b/src/tint/ir/module.h
index 244bc0a..c6cf74c 100644
--- a/src/tint/ir/module.h
+++ b/src/tint/ir/module.h
@@ -18,6 +18,7 @@
#include <string>
#include "src/tint/constant/value.h"
+#include "src/tint/ir/constant.h"
#include "src/tint/ir/function.h"
#include "src/tint/ir/instruction.h"
#include "src/tint/ir/value.h"
@@ -67,7 +68,7 @@
/// The flow node allocator
utils::BlockAllocator<FlowNode> flow_nodes;
/// The constant allocator
- utils::BlockAllocator<constant::Value> constants;
+ utils::BlockAllocator<constant::Value> constants_arena;
/// The value allocator
utils::BlockAllocator<Value> values;
@@ -82,6 +83,29 @@
/// The symbol table for the module
SymbolTable symbols{prog_id_};
+
+ /// ConstantHasher provides a hash function for a constant::Value pointer, hashing the value
+ /// instead of the pointer itself.
+ struct ConstantHasher {
+ /// @param c the constant pointer to create a hash for
+ /// @return the hash value
+ inline std::size_t operator()(const constant::Value* c) const { return c->Hash(); }
+ };
+
+ /// ConstantEquals provides an equality function for two constant::Value pointers, comparing
+ /// their values instead of the pointers.
+ struct ConstantEquals {
+ /// @param a the first constant pointer to compare
+ /// @param b the second constant pointer to compare
+ /// @return the hash value
+ inline bool operator()(const constant::Value* a, const constant::Value* b) const {
+ return a->Equal(b);
+ }
+ };
+
+ /// The map of constant::Value to their ir::Constant.
+ utils::Hashmap<const constant::Value*, ir::Constant*, 16, ConstantHasher, ConstantEquals>
+ constants;
};
} // namespace tint::ir
diff --git a/src/tint/ir/root_terminator.h b/src/tint/ir/root_terminator.h
index 361aa6d..c6338ba 100644
--- a/src/tint/ir/root_terminator.h
+++ b/src/tint/ir/root_terminator.h
@@ -25,7 +25,12 @@
public:
/// Constructor
RootTerminator();
+ RootTerminator(const RootTerminator&) = delete;
+ RootTerminator(RootTerminator&&) = delete;
~RootTerminator() override;
+
+ RootTerminator& operator=(const RootTerminator&) = delete;
+ RootTerminator& operator=(RootTerminator&&) = delete;
};
} // namespace tint::ir
diff --git a/src/tint/ir/store.cc b/src/tint/ir/store.cc
index 87f3620..8b8de46 100644
--- a/src/tint/ir/store.cc
+++ b/src/tint/ir/store.cc
@@ -19,11 +19,11 @@
namespace tint::ir {
-Store::Store(Value* t, Value* f) : Base(), to(t), from(f) {
- TINT_ASSERT(IR, to);
- TINT_ASSERT(IR, from);
- to->AddUsage(this);
- from->AddUsage(this);
+Store::Store(Value* to, Value* from) : Base(), to_(to), from_(from) {
+ TINT_ASSERT(IR, to_);
+ TINT_ASSERT(IR, from_);
+ to_->AddUsage(this);
+ from_->AddUsage(this);
}
Store::~Store() = default;
diff --git a/src/tint/ir/store.h b/src/tint/ir/store.h
index 9095c41..af5377d 100644
--- a/src/tint/ir/store.h
+++ b/src/tint/ir/store.h
@@ -34,10 +34,15 @@
Store& operator=(const Store& inst) = delete;
Store& operator=(Store&& inst) = delete;
- /// the value being stored to
- Value* to = nullptr;
- /// the value being stored
- Value* from = nullptr;
+ /// @returns the value being stored too
+ Value* To() const { return to_; }
+
+ /// @returns the value being stored
+ Value* From() const { return from_; }
+
+ private:
+ Value* to_;
+ Value* from_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/store_test.cc b/src/tint/ir/store_test.cc
index 902ca95..1906cc6 100644
--- a/src/tint/ir/store_test.cc
+++ b/src/tint/ir/store_test.cc
@@ -33,10 +33,10 @@
const auto* inst = b.Store(to, b.Constant(4_i));
ASSERT_TRUE(inst->Is<Store>());
- ASSERT_EQ(inst->to, to);
+ ASSERT_EQ(inst->To(), to);
- ASSERT_TRUE(inst->from->Is<Constant>());
- auto lhs = inst->from->As<Constant>()->value;
+ ASSERT_TRUE(inst->From()->Is<Constant>());
+ auto lhs = inst->From()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -48,13 +48,13 @@
auto* to = b.Discard();
const auto* inst = b.Store(to, b.Constant(4_i));
- ASSERT_NE(inst->to, nullptr);
- ASSERT_EQ(inst->to->Usage().Length(), 1u);
- EXPECT_EQ(inst->to->Usage()[0], inst);
+ ASSERT_NE(inst->To(), nullptr);
+ ASSERT_EQ(inst->To()->Usage().Length(), 1u);
+ EXPECT_EQ(inst->To()->Usage()[0], inst);
- ASSERT_NE(inst->from, nullptr);
- ASSERT_EQ(inst->from->Usage().Length(), 1u);
- EXPECT_EQ(inst->from->Usage()[0], inst);
+ ASSERT_NE(inst->From(), nullptr);
+ ASSERT_EQ(inst->From()->Usage().Length(), 1u);
+ EXPECT_EQ(inst->From()->Usage()[0], inst);
}
} // namespace
diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc
index ad6a145..3bccc83 100644
--- a/src/tint/ir/switch.cc
+++ b/src/tint/ir/switch.cc
@@ -18,7 +18,7 @@
namespace tint::ir {
-Switch::Switch(Value* cond) : Base(), condition(cond) {}
+Switch::Switch(Value* cond) : Base(), condition_(cond) {}
Switch::~Switch() = default;
diff --git a/src/tint/ir/switch.h b/src/tint/ir/switch.h
index 6be7b62..2e977ec 100644
--- a/src/tint/ir/switch.h
+++ b/src/tint/ir/switch.h
@@ -41,21 +41,40 @@
utils::Vector<CaseSelector, 4> selectors;
/// The start block for the case block.
Branch start = {};
+
+ /// @returns the case start target
+ const Branch& Start() const { return start; }
+ /// @returns the case start target
+ Branch& Start() { return start; }
};
/// Constructor
/// @param cond the condition
explicit Switch(Value* cond);
+ Switch(const Switch&) = delete;
+ Switch(Switch&&) = delete;
~Switch() override;
- /// The switch merge target
- Branch merge = {};
+ Switch& operator=(const Switch&) = delete;
+ Switch& operator=(Switch&&) = delete;
- /// The switch case statements
- utils::Vector<Case, 4> cases;
+ /// @returns the switch merge branch
+ const Branch& Merge() const { return merge_; }
+ /// @returns the switch merge branch
+ Branch& Merge() { return merge_; }
- /// Value holding the condition result
- const Value* condition = nullptr;
+ /// @returns the switch cases
+ utils::VectorRef<Case> Cases() const { return cases_; }
+ /// @returns the switch cases
+ utils::Vector<Case, 4>& Cases() { return cases_; }
+
+ /// @returns the condition
+ const Value* Condition() const { return condition_; }
+
+ private:
+ Branch merge_ = {};
+ utils::Vector<Case, 4> cases_;
+ Value* condition_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/test_helper.h b/src/tint/ir/test_helper.h
index d9055ce..b7279f3 100644
--- a/src/tint/ir/test_helper.h
+++ b/src/tint/ir/test_helper.h
@@ -33,7 +33,6 @@
class TestHelperBase : public BASE, public ProgramBuilder {
public:
TestHelperBase() = default;
-
~TestHelperBase() override = default;
/// Build the module, cleaning up the program before returning.
diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc
index b1aa67f..9be189c 100644
--- a/src/tint/ir/to_program.cc
+++ b/src/tint/ir/to_program.cc
@@ -91,14 +91,14 @@
const ast::Function* Fn(const Function* fn) {
SCOPED_NESTING();
- auto name = Sym(fn->name);
+ auto name = Sym(fn->Name());
// TODO(crbug.com/tint/1915): Properly implement this when we've fleshed out Function
utils::Vector<const ast::Parameter*, 1> params{};
- auto ret_ty = Type(fn->return_type);
+ auto ret_ty = Type(fn->ReturnType());
if (!ret_ty) {
return nullptr;
}
- auto* body = FlowNodeGraph(fn->start_target);
+ auto* body = FlowNodeGraph(fn->StartTarget());
if (!body) {
return nullptr;
}
@@ -126,7 +126,7 @@
branch->target,
[&](const ir::Block* block) {
- for (auto* inst : block->instructions) {
+ for (const auto* inst : block->Instructions()) {
auto stmt = Stmt(inst);
if (TINT_UNLIKELY(!stmt)) {
return kError;
@@ -135,7 +135,7 @@
stmts.Push(s);
}
}
- branch = &block->branch;
+ branch = &block->Branch();
return kContinue;
},
@@ -145,8 +145,8 @@
return kError;
}
stmts.Push(stmt);
- branch = &if_->merge;
- return branch->target->inbound_branches.IsEmpty() ? kStop : kContinue;
+ branch = &if_->Merge();
+ return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue;
},
[&](const ir::Switch* switch_) {
@@ -155,8 +155,8 @@
return kError;
}
stmts.Push(stmt);
- branch = &switch_->merge;
- return branch->target->inbound_branches.IsEmpty() ? kStop : kContinue;
+ branch = &switch_->Merge();
+ return branch->target->InboundBranches().IsEmpty() ? kStop : kContinue;
},
[&](const ir::FunctionTerminator*) {
@@ -189,25 +189,25 @@
const ast::IfStatement* If(const ir::If* i) {
SCOPED_NESTING();
- auto* cond = Expr(i->condition);
- auto* t = FlowNodeGraph(i->true_.target, i->merge.target);
+ auto* cond = Expr(i->Condition());
+ auto* t = FlowNodeGraph(i->True().target, i->Merge().target);
if (TINT_UNLIKELY(!t)) {
return nullptr;
}
- if (!IsEmpty(i->false_.target, i->merge.target)) {
- // If the else target is an if flow node with the same merge target as this if, then
+ if (!IsEmpty(i->False().target, i->Merge().target)) {
+ // If the else target is an if flow node with the same Merge().target as this if, then
// emit an 'else if' instead of a block statement for the else.
- if (auto* else_if = As<ir::If>(NextNonEmptyNode(i->false_.target));
+ if (auto* else_if = As<ir::If>(NextNonEmptyNode(i->False().target));
else_if &&
- NextNonEmptyNode(i->merge.target) == NextNonEmptyNode(else_if->merge.target)) {
+ NextNonEmptyNode(i->Merge().target) == NextNonEmptyNode(else_if->Merge().target)) {
auto* f = If(else_if);
if (!f) {
return nullptr;
}
return b.If(cond, t, b.Else(f));
} else {
- auto* f = FlowNodeGraph(i->false_.target, i->merge.target);
+ auto* f = FlowNodeGraph(i->False().target, i->Merge().target);
if (!f) {
return nullptr;
}
@@ -221,16 +221,16 @@
const ast::SwitchStatement* Switch(const ir::Switch* s) {
SCOPED_NESTING();
- auto* cond = Expr(s->condition);
+ auto* cond = Expr(s->Condition());
if (!cond) {
return nullptr;
}
- auto cases = utils::Transform(
- s->cases, //
+ auto cases = utils::Transform<1>(
+ s->Cases(), //
[&](const ir::Switch::Case& c) -> const tint::ast::CaseStatement* {
SCOPED_NESTING();
- auto* body = FlowNodeGraph(c.start.target, s->merge.target);
+ auto* body = FlowNodeGraph(c.start.target, s->Merge().target);
if (!body) {
return nullptr;
}
@@ -292,10 +292,10 @@
bool IsEmpty(const ir::FlowNode* node, const ir::FlowNode* stop_at) {
while (node != stop_at) {
if (auto* block = node->As<ir::Block>()) {
- if (block->instructions.Length() > 0) {
+ if (!block->Instructions().IsEmpty()) {
return false;
}
- node = block->branch.target;
+ node = block->Branch().target;
} else {
return false;
}
@@ -307,13 +307,13 @@
const ir::FlowNode* NextNonEmptyNode(const ir::FlowNode* node) {
while (node) {
if (auto* block = node->As<ir::Block>()) {
- for (auto* inst : block->instructions) {
+ for (const auto* inst : block->Instructions()) {
// Load instructions will be inlined, so ignore them.
if (!inst->Is<ir::Load>()) {
return node;
}
}
- node = block->branch.target;
+ node = block->Branch().target;
} else {
return node;
}
@@ -351,8 +351,8 @@
}
auto ty = Type(ptr->StoreType());
const ast::Expression* init = nullptr;
- if (var->initializer) {
- init = Expr(var->initializer);
+ if (var->Initializer()) {
+ init = Expr(var->Initializer());
if (!init) {
return nullptr;
}
@@ -368,18 +368,19 @@
}
const ast::AssignmentStatement* Store(const ir::Store* store) {
- auto* expr = Expr(store->from);
- return b.Assign(NameOf(store->to), expr);
+ auto* expr = Expr(store->From());
+ return b.Assign(NameOf(store->To()), expr);
}
const ast::CallExpression* Call(const ir::Call* call) {
- auto args = utils::Transform(call->args, [&](const ir::Value* arg) { return Expr(arg); });
+ auto args =
+ utils::Transform<2>(call->Args(), [&](const ir::Value* arg) { return Expr(arg); });
if (args.Any(utils::IsNull)) {
return nullptr;
}
return tint::Switch(
call, //
- [&](const ir::UserCall* c) { return b.Call(Sym(c->name), std::move(args)); },
+ [&](const ir::UserCall* c) { return b.Call(Sym(c->Name()), std::move(args)); },
[&](Default) {
UNHANDLED_CASE(call);
return nullptr;
@@ -401,18 +402,18 @@
const ast::Expression* ConstExpr(const ir::Constant* c) {
return tint::Switch(
c->Type(), //
- [&](const type::I32*) { return b.Expr(c->value->ValueAs<i32>()); },
- [&](const type::U32*) { return b.Expr(c->value->ValueAs<u32>()); },
- [&](const type::F32*) { return b.Expr(c->value->ValueAs<f32>()); },
- [&](const type::F16*) { return b.Expr(c->value->ValueAs<f16>()); },
- [&](const type::Bool*) { return b.Expr(c->value->ValueAs<bool>()); },
+ [&](const type::I32*) { return b.Expr(c->Value()->ValueAs<i32>()); },
+ [&](const type::U32*) { return b.Expr(c->Value()->ValueAs<u32>()); },
+ [&](const type::F32*) { return b.Expr(c->Value()->ValueAs<f32>()); },
+ [&](const type::F16*) { return b.Expr(c->Value()->ValueAs<f16>()); },
+ [&](const type::Bool*) { return b.Expr(c->Value()->ValueAs<bool>()); },
[&](Default) {
UNHANDLED_CASE(c);
return nullptr;
});
}
- const ast::Expression* LoadExpr(const ir::Load* l) { return Expr(l->from); }
+ const ast::Expression* LoadExpr(const ir::Load* l) { return Expr(l->From()); }
const ast::Expression* VarExpr(const ir::Var* v) { return b.Expr(NameOf(v)); }
diff --git a/src/tint/ir/transform/add_empty_entry_point.cc b/src/tint/ir/transform/add_empty_entry_point.cc
index 809a6ad..6788d7c 100644
--- a/src/tint/ir/transform/add_empty_entry_point.cc
+++ b/src/tint/ir/transform/add_empty_entry_point.cc
@@ -29,7 +29,7 @@
void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
for (auto* func : ir->functions) {
- if (func->pipeline_stage != Function::PipelineStage::kUndefined) {
+ if (func->Stage() != Function::PipelineStage::kUndefined) {
return;
}
}
@@ -38,7 +38,7 @@
auto* ep =
builder.CreateFunction(ir->symbols.New("unused_entry_point"), ir->types.Get<type::Void>(),
Function::PipelineStage::kCompute, std::array{1u, 1u, 1u});
- builder.Branch(ep->start_target, ep->end_target);
+ ep->StartTarget()->BranchTo(ep->EndTarget());
ir->functions.Push(ep);
}
diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc
index baba8e0..a363d83 100644
--- a/src/tint/ir/transform/add_empty_entry_point_test.cc
+++ b/src/tint/ir/transform/add_empty_entry_point_test.cc
@@ -40,7 +40,7 @@
TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) {
auto* ep = b.CreateFunction(mod.symbols.New("main"), mod.types.Get<type::Void>(),
Function::PipelineStage::kFragment);
- b.Branch(ep->start_target, ep->end_target);
+ ep->StartTarget()->BranchTo(ep->EndTarget());
mod.functions.Push(ep);
auto* expect = R"(
diff --git a/src/tint/ir/unary.cc b/src/tint/ir/unary.cc
index db13d5d..2a0ee94 100644
--- a/src/tint/ir/unary.cc
+++ b/src/tint/ir/unary.cc
@@ -19,8 +19,8 @@
namespace tint::ir {
-Unary::Unary(Kind k, const type::Type* res_ty, Value* val)
- : kind(k), result_type(res_ty), val_(val) {
+Unary::Unary(enum Kind k, const type::Type* res_ty, Value* val)
+ : kind_(k), result_type_(res_ty), val_(val) {
TINT_ASSERT(IR, val_);
val_->AddUsage(this);
}
diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h
index 46edd45..98ad9b8 100644
--- a/src/tint/ir/unary.h
+++ b/src/tint/ir/unary.h
@@ -33,7 +33,7 @@
/// @param kind the kind of unary instruction
/// @param result_type the result type
/// @param val the input value for the instruction
- Unary(Kind kind, const type::Type* result_type, Value* val);
+ Unary(enum Kind kind, const type::Type* result_type, Value* val);
Unary(const Unary& inst) = delete;
Unary(Unary&& inst) = delete;
~Unary() override;
@@ -42,19 +42,18 @@
Unary& operator=(Unary&& inst) = delete;
/// @returns the type of the value
- const type::Type* Type() const override { return result_type; }
+ const type::Type* Type() const override { return result_type_; }
/// @returns the value for the instruction
const Value* Val() const { return val_; }
- /// the kind of unary instruction
- Kind kind = Kind::kNegation;
-
- /// the result type of the instruction
- const type::Type* result_type = nullptr;
+ /// @returns the kind of unary instruction
+ enum Kind Kind() const { return kind_; }
private:
- Value* val_ = nullptr;
+ enum Kind kind_;
+ const type::Type* result_type_;
+ Value* val_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc
index 280c7de..6cefed3 100644
--- a/src/tint/ir/unary_test.cc
+++ b/src/tint/ir/unary_test.cc
@@ -29,10 +29,10 @@
const auto* inst = b.Complement(b.ir.types.Get<type::I32>(), b.Constant(4_i));
ASSERT_TRUE(inst->Is<Unary>());
- EXPECT_EQ(inst->kind, Unary::Kind::kComplement);
+ EXPECT_EQ(inst->Kind(), Unary::Kind::kComplement);
ASSERT_TRUE(inst->Val()->Is<Constant>());
- auto lhs = inst->Val()->As<Constant>()->value;
+ auto lhs = inst->Val()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -43,10 +43,10 @@
const auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
ASSERT_TRUE(inst->Is<Unary>());
- EXPECT_EQ(inst->kind, Unary::Kind::kNegation);
+ EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation);
ASSERT_TRUE(inst->Val()->Is<Constant>());
- auto lhs = inst->Val()->As<Constant>()->value;
+ auto lhs = inst->Val()->As<Constant>()->Value();
ASSERT_TRUE(lhs->Is<constant::Scalar<i32>>());
EXPECT_EQ(4_i, lhs->As<constant::Scalar<i32>>()->ValueAs<i32>());
}
@@ -56,7 +56,7 @@
Builder b{mod};
const auto* inst = b.Negation(b.ir.types.Get<type::I32>(), b.Constant(4_i));
- EXPECT_EQ(inst->kind, Unary::Kind::kNegation);
+ EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation);
ASSERT_NE(inst->Val(), nullptr);
ASSERT_EQ(inst->Val()->Usage().Length(), 1u);
diff --git a/src/tint/ir/user_call.cc b/src/tint/ir/user_call.cc
index f718284..e44e285 100644
--- a/src/tint/ir/user_call.cc
+++ b/src/tint/ir/user_call.cc
@@ -23,7 +23,7 @@
namespace tint::ir {
UserCall::UserCall(const type::Type* ty, Symbol n, utils::VectorRef<Value*> arguments)
- : Base(ty, std::move(arguments)), name(n) {}
+ : Base(ty, std::move(arguments)), name_(n) {}
UserCall::~UserCall() = default;
diff --git a/src/tint/ir/user_call.h b/src/tint/ir/user_call.h
index 5ada8f8..ba52e20 100644
--- a/src/tint/ir/user_call.h
+++ b/src/tint/ir/user_call.h
@@ -36,8 +36,11 @@
UserCall& operator=(const UserCall& inst) = delete;
UserCall& operator=(UserCall&& inst) = delete;
- /// The function name
- Symbol name;
+ /// @returns the called function name
+ Symbol Name() const { return name_; }
+
+ private:
+ Symbol name_;
};
} // namespace tint::ir
diff --git a/src/tint/ir/var.cc b/src/tint/ir/var.cc
index 9bf4b8d..9b43329 100644
--- a/src/tint/ir/var.cc
+++ b/src/tint/ir/var.cc
@@ -19,8 +19,14 @@
namespace tint::ir {
-Var::Var(const type::Type* ty) : type(ty) {}
+Var::Var(const type::Type* ty) : type_(ty) {}
Var::~Var() = default;
+void Var::SetInitializer(Value* initializer) {
+ initializer_ = initializer;
+ initializer_->AddUsage(this);
+ // TODO(dsinclair): Probably should do a RemoveUsage on an existing initializer if set
+}
+
} // namespace tint::ir
diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h
index c874a62..67ccf0d 100644
--- a/src/tint/ir/var.h
+++ b/src/tint/ir/var.h
@@ -36,13 +36,17 @@
Var& operator=(Var&& inst) = delete;
/// @returns the type of the var
- const type::Type* Type() const override { return type; }
+ const type::Type* Type() const override { return type_; }
- /// the result type of the instruction
- const type::Type* type = nullptr;
+ /// Sets the var initializer
+ /// @param initializer the initializer
+ void SetInitializer(Value* initializer);
+ /// @returns the initializer
+ const Value* Initializer() const { return initializer_; }
- /// The optional initializer
- Value* initializer = nullptr;
+ private:
+ const type::Type* type_;
+ Value* initializer_ = nullptr;
};
} // namespace tint::ir
diff --git a/src/tint/transform/manager_test.cc b/src/tint/transform/manager_test.cc
index 3b0ee25..a81f7bf 100644
--- a/src/tint/transform/manager_test.cc
+++ b/src/tint/transform/manager_test.cc
@@ -52,7 +52,7 @@
ir::Builder builder(*mod);
auto* func =
builder.CreateFunction(mod->symbols.New("ir_func"), mod->types.Get<type::Void>());
- builder.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
mod->functions.Push(func);
}
};
@@ -70,7 +70,7 @@
ir::Builder builder(mod);
auto* func =
builder.CreateFunction(builder.ir.symbols.New("main"), builder.ir.types.Get<type::Void>());
- builder.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
builder.ir.functions.Push(func);
return mod;
}
@@ -102,9 +102,10 @@
manager.Add<IR_AddFunction>();
manager.Run(&ir, {}, outputs);
+
ASSERT_EQ(ir.functions.Length(), 2u);
- EXPECT_EQ(ir.functions[0]->name.Name(), "main");
- EXPECT_EQ(ir.functions[1]->name.Name(), "ir_func");
+ EXPECT_EQ(ir.functions[0]->Name().Name(), "main");
+ EXPECT_EQ(ir.functions[1]->Name().Name(), "ir_func");
}
TEST_F(TransformManagerTest, AST_MixedTransforms_AST_Before_IR) {
@@ -149,9 +150,9 @@
manager.Run(&ir, {}, outputs);
ASSERT_EQ(ir.functions.Length(), 3u);
- EXPECT_EQ(ir.functions[0]->name.Name(), "ast_func");
- EXPECT_EQ(ir.functions[1]->name.Name(), "main");
- EXPECT_EQ(ir.functions[2]->name.Name(), "ir_func");
+ EXPECT_EQ(ir.functions[0]->Name().Name(), "ast_func");
+ EXPECT_EQ(ir.functions[1]->Name().Name(), "main");
+ EXPECT_EQ(ir.functions[2]->Name().Name(), "ir_func");
}
TEST_F(TransformManagerTest, IR_MixedTransforms_IR_Before_AST) {
@@ -164,9 +165,9 @@
manager.Run(&ir, {}, outputs);
ASSERT_EQ(ir.functions.Length(), 3u);
- EXPECT_EQ(ir.functions[0]->name.Name(), "ast_func");
- EXPECT_EQ(ir.functions[1]->name.Name(), "main");
- EXPECT_EQ(ir.functions[2]->name.Name(), "ir_func");
+ EXPECT_EQ(ir.functions[0]->Name().Name(), "ast_func");
+ EXPECT_EQ(ir.functions[1]->Name().Name(), "main");
+ EXPECT_EQ(ir.functions[2]->Name().Name(), "ir_func");
}
#endif // TINT_BUILD_IR
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc
index 4ec4d9d..bd9a735 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc
@@ -19,6 +19,7 @@
#include "src/tint/ir/block.h"
#include "src/tint/ir/function_terminator.h"
#include "src/tint/ir/if.h"
+#include "src/tint/ir/load.h"
#include "src/tint/ir/module.h"
#include "src/tint/ir/store.h"
#include "src/tint/ir/transform/add_empty_entry_point.h"
@@ -109,7 +110,7 @@
}
uint32_t GeneratorImplIr::Constant(const ir::Constant* constant) {
- return Constant(constant->value);
+ return Constant(constant->Value());
}
uint32_t GeneratorImplIr::Constant(const constant::Value* constant) {
@@ -213,15 +214,15 @@
auto id = module_.NextId();
// Emit the function name.
- module_.PushDebug(spv::Op::OpName, {id, Operand(func->name.Name())});
+ module_.PushDebug(spv::Op::OpName, {id, Operand(func->Name().Name())});
// Emit OpEntryPoint and OpExecutionMode declarations if needed.
- if (func->pipeline_stage != ir::Function::PipelineStage::kUndefined) {
+ if (func->Stage() != ir::Function::PipelineStage::kUndefined) {
EmitEntryPoint(func, id);
}
// Get the ID for the return type.
- auto return_type_id = Type(func->return_type);
+ auto return_type_id = Type(func->ReturnType());
// Get the ID for the function type (creating it if needed).
// TODO(jrprice): Add the parameter types when they are supported in the IR.
@@ -247,7 +248,7 @@
TINT_DEFER(current_function_ = Function());
// Emit the body of the function.
- EmitBlock(func->start_target);
+ EmitBlock(func->StartTarget());
// Add the function to the module.
module_.PushFunction(current_function_);
@@ -255,13 +256,13 @@
void GeneratorImplIr::EmitEntryPoint(const ir::Function* func, uint32_t id) {
SpvExecutionModel stage = SpvExecutionModelMax;
- switch (func->pipeline_stage) {
+ switch (func->Stage()) {
case ir::Function::PipelineStage::kCompute: {
stage = SpvExecutionModelGLCompute;
module_.PushExecutionMode(
spv::Op::OpExecutionMode,
- {id, U32Operand(SpvExecutionModeLocalSize), func->workgroup_size->at(0),
- func->workgroup_size->at(1), func->workgroup_size->at(2)});
+ {id, U32Operand(SpvExecutionModeLocalSize), func->WorkgroupSize()->at(0),
+ func->WorkgroupSize()->at(1), func->WorkgroupSize()->at(2)});
break;
}
case ir::Function::PipelineStage::kFragment: {
@@ -281,7 +282,7 @@
}
// TODO(jrprice): Add the interface list of all referenced global variables.
- module_.PushEntryPoint(spv::Op::OpEntryPoint, {U32Operand(stage), id, func->name.Name()});
+ module_.PushEntryPoint(spv::Op::OpEntryPoint, {U32Operand(stage), id, func->Name().Name()});
}
void GeneratorImplIr::EmitBlock(const ir::Block* block) {
@@ -292,10 +293,11 @@
}
// Emit the instructions.
- for (auto* inst : block->instructions) {
+ for (const auto* inst : block->Instructions()) {
auto result = Switch(
inst, //
[&](const ir::Binary* b) { return EmitBinary(b); },
+ [&](const ir::Load* l) { return EmitLoad(l); },
[&](const ir::Store* s) {
EmitStore(s);
return 0u;
@@ -311,43 +313,43 @@
// Handle the branch at the end of the block.
Switch(
- block->branch.target,
+ block->Branch().target,
[&](const ir::Block* b) { current_function_.push_inst(spv::Op::OpBranch, {Label(b)}); },
[&](const ir::If* i) { EmitIf(i); },
[&](const ir::FunctionTerminator*) {
// TODO(jrprice): Handle the return value, which will be a branch argument.
- if (!block->branch.args.IsEmpty()) {
+ if (!block->Branch().args.IsEmpty()) {
TINT_ICE(Writer, diagnostics_) << "unimplemented return value";
}
current_function_.push_inst(spv::Op::OpReturn, {});
},
[&](Default) {
- if (!block->branch.target) {
+ if (!block->Branch().target) {
// A block may not have an outward branch (e.g. an unreachable merge block).
current_function_.push_inst(spv::Op::OpUnreachable, {});
} else {
TINT_ICE(Writer, diagnostics_)
- << "unimplemented branch target: " << block->branch.target->TypeInfo().name;
+ << "unimplemented branch target: " << block->Branch().target->TypeInfo().name;
}
});
}
void GeneratorImplIr::EmitIf(const ir::If* i) {
- auto* merge_block = i->merge.target->As<ir::Block>();
- auto* true_block = i->true_.target->As<ir::Block>();
- auto* false_block = i->false_.target->As<ir::Block>();
+ auto* merge_block = i->Merge().target->As<ir::Block>();
+ auto* true_block = i->True().target->As<ir::Block>();
+ auto* false_block = i->False().target->As<ir::Block>();
// Generate labels for the blocks. We emit the true or false block if it:
// 1. contains instructions, or
- // 2. branches somewhere other then the merge target.
+ // 2. branches somewhere other then the Merge().target.
// Otherwise we skip them and branch straight to the merge block.
uint32_t merge_label = Label(merge_block);
uint32_t true_label = merge_label;
uint32_t false_label = merge_label;
- if (!true_block->instructions.IsEmpty() || true_block->branch.target != merge_block) {
+ if (!true_block->Instructions().IsEmpty() || true_block->Branch().target != merge_block) {
true_label = Label(true_block);
}
- if (!false_block->instructions.IsEmpty() || false_block->branch.target != merge_block) {
+ if (!false_block->Instructions().IsEmpty() || false_block->Branch().target != merge_block) {
false_label = Label(false_block);
}
@@ -355,7 +357,7 @@
current_function_.push_inst(spv::Op::OpSelectionMerge,
{merge_label, U32Operand(SpvSelectionControlMaskNone)});
current_function_.push_inst(spv::Op::OpBranchConditional,
- {Value(i->condition), true_label, false_label});
+ {Value(i->Condition()), true_label, false_label});
// Emit the `true` and `false` blocks, if they're not being skipped.
if (true_label != merge_label) {
@@ -374,7 +376,7 @@
// Determine the opcode.
spv::Op op = spv::Op::Max;
- switch (binary->kind) {
+ switch (binary->Kind()) {
case ir::Binary::Kind::kAdd: {
op = binary->Type()->is_integer_scalar_or_vector() ? spv::Op::OpIAdd : spv::Op::OpFAdd;
break;
@@ -385,7 +387,7 @@
}
default: {
TINT_ICE(Writer, diagnostics_)
- << "unimplemented binary instruction: " << static_cast<uint32_t>(binary->kind);
+ << "unimplemented binary instruction: " << static_cast<uint32_t>(binary->Kind());
}
}
@@ -396,8 +398,14 @@
return id;
}
+uint32_t GeneratorImplIr::EmitLoad(const ir::Load* load) {
+ auto id = module_.NextId();
+ current_function_.push_inst(spv::Op::OpLoad, {Type(load->Type()), id, Value(load->From())});
+ return id;
+}
+
void GeneratorImplIr::EmitStore(const ir::Store* store) {
- current_function_.push_inst(spv::Op::OpStore, {Value(store->to), Value(store->from)});
+ current_function_.push_inst(spv::Op::OpStore, {Value(store->To()), Value(store->From())});
}
uint32_t GeneratorImplIr::EmitVar(const ir::Var* var) {
@@ -409,8 +417,8 @@
if (ptr->AddressSpace() == builtin::AddressSpace::kFunction) {
TINT_ASSERT(Writer, current_function_);
current_function_.push_var({ty, id, U32Operand(SpvStorageClassFunction)});
- if (var->initializer) {
- current_function_.push_inst(spv::Op::OpStore, {id, Value(var->initializer)});
+ if (var->Initializer()) {
+ current_function_.push_inst(spv::Op::OpStore, {id, Value(var->Initializer())});
}
} else {
TINT_ICE(Writer, diagnostics_)
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.h b/src/tint/writer/spirv/ir/generator_impl_ir.h
index 66ffe48..b72201f 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir.h
+++ b/src/tint/writer/spirv/ir/generator_impl_ir.h
@@ -32,6 +32,7 @@
class Block;
class If;
class Function;
+class Load;
class Module;
class Store;
class Value;
@@ -106,6 +107,11 @@
/// @returns the result ID of the instruction
uint32_t EmitBinary(const ir::Binary* binary);
+ /// Emit a load instruction.
+ /// @param load the load instruction to emit
+ /// @returns the result ID of the instruction
+ uint32_t EmitLoad(const ir::Load* load);
+
/// Emit a store instruction.
/// @param store the store instruction to emit
void EmitStore(const ir::Store* store);
@@ -151,25 +157,6 @@
}
};
- /// ConstantHasher provides a hash function for a constant::Value pointer, hashing the value
- /// instead of the pointer itself.
- struct ConstantHasher {
- /// @param c the constant::Value pointer to create a hash for
- /// @return the hash value
- inline std::size_t operator()(const constant::Value* c) const { return c->Hash(); }
- };
-
- /// ConstantEquals provides an equality function for two constant::Value pointers, comparing
- /// their values instead of the pointers.
- struct ConstantEquals {
- /// @param a the first constant::Value pointer to compare
- /// @param b the second constant::Value pointer to compare
- /// @return the hash value
- inline bool operator()(const constant::Value* a, const constant::Value* b) const {
- return a->Equal(b);
- }
- };
-
/// The map of types to their result IDs.
utils::Hashmap<const type::Type*, uint32_t, 8> types_;
@@ -177,7 +164,7 @@
utils::Hashmap<FunctionType, uint32_t, 8, FunctionType::Hasher> function_types_;
/// The map of constants to their result IDs.
- utils::Hashmap<const constant::Value*, uint32_t, 16, ConstantHasher, ConstantEquals> constants_;
+ utils::Hashmap<const constant::Value*, uint32_t, 16> constants_;
/// The map of instructions to their result IDs.
utils::Hashmap<const ir::Instruction*, uint32_t, 8> instructions_;
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
index 30dc059..e9231f3 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
@@ -21,10 +21,10 @@
TEST_F(SpvGeneratorImplTest, Binary_Add_I32) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- func->start_target->instructions.Push(
- b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i)));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -43,10 +43,10 @@
TEST_F(SpvGeneratorImplTest, Binary_Add_U32) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- func->start_target->instructions.Push(
- b.Add(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u)));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Add(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -65,10 +65,10 @@
TEST_F(SpvGeneratorImplTest, Binary_Add_F32) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- func->start_target->instructions.Push(
- b.Add(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f)));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Add(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -87,10 +87,10 @@
TEST_F(SpvGeneratorImplTest, Binary_Sub_I32) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- func->start_target->instructions.Push(
- b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i)));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -109,10 +109,10 @@
TEST_F(SpvGeneratorImplTest, Binary_Sub_U32) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- func->start_target->instructions.Push(
- b.Subtract(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u)));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Subtract(mod.types.Get<type::U32>(), b.Constant(1_u), b.Constant(2_u))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -131,10 +131,10 @@
TEST_F(SpvGeneratorImplTest, Binary_Sub_F32) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- func->start_target->instructions.Push(
- b.Subtract(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f)));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Subtract(mod.types.Get<type::F32>(), b.Constant(1_f), b.Constant(2_f))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -153,17 +153,17 @@
TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec2i) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- auto* lhs = mod.constants.Create<constant::Composite>(
+ auto* lhs = mod.constants_arena.Create<constant::Composite>(
mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
- utils::Vector{b.Constant(42_i)->value, b.Constant(-1_i)->value}, false, false);
- auto* rhs = mod.constants.Create<constant::Composite>(
+ utils::Vector{b.Constant(42_i)->Value(), b.Constant(-1_i)->Value()}, false, false);
+ auto* rhs = mod.constants_arena.Create<constant::Composite>(
mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
- utils::Vector{b.Constant(0_i)->value, b.Constant(-43_i)->value}, false, false);
- func->start_target->instructions.Push(
- b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u), b.Constant(lhs),
- b.Constant(rhs)));
+ utils::Vector{b.Constant(0_i)->Value(), b.Constant(-43_i)->Value()}, false, false);
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::I32>(), 2u),
+ b.Constant(lhs), b.Constant(rhs))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -187,21 +187,21 @@
TEST_F(SpvGeneratorImplTest, Binary_Sub_Vec4f) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
- auto* lhs = mod.constants.Create<constant::Composite>(
+ auto* lhs = mod.constants_arena.Create<constant::Composite>(
mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
- utils::Vector{b.Constant(42_f)->value, b.Constant(-1_f)->value, b.Constant(0_f)->value,
- b.Constant(1.25_f)->value},
+ utils::Vector{b.Constant(42_f)->Value(), b.Constant(-1_f)->Value(),
+ b.Constant(0_f)->Value(), b.Constant(1.25_f)->Value()},
false, false);
- auto* rhs = mod.constants.Create<constant::Composite>(
+ auto* rhs = mod.constants_arena.Create<constant::Composite>(
mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
- utils::Vector{b.Constant(0_f)->value, b.Constant(1.25_f)->value, b.Constant(-42_f)->value,
- b.Constant(1_f)->value},
+ utils::Vector{b.Constant(0_f)->Value(), b.Constant(1.25_f)->Value(),
+ b.Constant(-42_f)->Value(), b.Constant(1_f)->Value()},
false, false);
- func->start_target->instructions.Push(
- b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u), b.Constant(lhs),
- b.Constant(rhs)));
+ func->StartTarget()->SetInstructions(
+ utils::Vector{b.Subtract(mod.types.Get<type::Vector>(mod.types.Get<type::F32>(), 4u),
+ b.Constant(lhs), b.Constant(rhs))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -227,11 +227,10 @@
TEST_F(SpvGeneratorImplTest, Binary_Chain) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
auto* a = b.Subtract(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(2_i));
- func->start_target->instructions.Push(a);
- func->start_target->instructions.Push(b.Add(mod.types.Get<type::I32>(), a, a));
+ func->StartTarget()->SetInstructions(utils::Vector{a, b.Add(mod.types.Get<type::I32>(), a, a)});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
index 95fce03..6ab48aa 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
@@ -65,9 +65,9 @@
TEST_F(SpvGeneratorImplTest, Constant_Vec4Bool) {
auto* t = b.Constant(true);
auto* f = b.Constant(false);
- auto* v = mod.constants.Create<constant::Composite>(
+ auto* v = mod.constants_arena.Create<constant::Composite>(
mod.types.Get<type::Vector>(mod.types.Get<type::Bool>(), 4u),
- utils::Vector{t->value, f->value, f->value, t->value}, false, true);
+ utils::Vector{t->Value(), f->Value(), f->Value(), t->Value()}, false, true);
generator_.Constant(b.Constant(v));
EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeBool
%2 = OpTypeVector %3 4
@@ -81,8 +81,9 @@
auto* i = mod.types.Get<type::I32>();
auto* i_42 = b.Constant(i32(42));
auto* i_n1 = b.Constant(i32(-1));
- auto* v = mod.constants.Create<constant::Composite>(
- mod.types.Get<type::Vector>(i, 2u), utils::Vector{i_42->value, i_n1->value}, false, false);
+ auto* v = mod.constants_arena.Create<constant::Composite>(
+ mod.types.Get<type::Vector>(i, 2u), utils::Vector{i_42->Value(), i_n1->Value()}, false,
+ false);
generator_.Constant(b.Constant(v));
EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeInt 32 1
%2 = OpTypeVector %3 2
@@ -97,9 +98,9 @@
auto* u_42 = b.Constant(u32(42));
auto* u_0 = b.Constant(u32(0));
auto* u_4b = b.Constant(u32(4000000000));
- auto* v = mod.constants.Create<constant::Composite>(
- mod.types.Get<type::Vector>(u, 3u), utils::Vector{u_42->value, u_0->value, u_4b->value},
- false, true);
+ auto* v = mod.constants_arena.Create<constant::Composite>(
+ mod.types.Get<type::Vector>(u, 3u),
+ utils::Vector{u_42->Value(), u_0->Value(), u_4b->Value()}, false, true);
generator_.Constant(b.Constant(v));
EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeInt 32 0
%2 = OpTypeVector %3 3
@@ -116,9 +117,9 @@
auto* f_0 = b.Constant(f32(0));
auto* f_q = b.Constant(f32(0.25));
auto* f_n1 = b.Constant(f32(-1));
- auto* v = mod.constants.Create<constant::Composite>(
+ auto* v = mod.constants_arena.Create<constant::Composite>(
mod.types.Get<type::Vector>(f, 4u),
- utils::Vector{f_42->value, f_0->value, f_q->value, f_n1->value}, false, true);
+ utils::Vector{f_42->Value(), f_0->Value(), f_q->Value(), f_n1->Value()}, false, true);
generator_.Constant(b.Constant(v));
EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeFloat 32
%2 = OpTypeVector %3 4
@@ -134,8 +135,9 @@
auto* h = mod.types.Get<type::F16>();
auto* h_42 = b.Constant(f16(42));
auto* h_q = b.Constant(f16(0.25));
- auto* v = mod.constants.Create<constant::Composite>(
- mod.types.Get<type::Vector>(h, 2u), utils::Vector{h_42->value, h_q->value}, false, false);
+ auto* v = mod.constants_arena.Create<constant::Composite>(
+ mod.types.Get<type::Vector>(h, 2u), utils::Vector{h_42->Value(), h_q->Value()}, false,
+ false);
generator_.Constant(b.Constant(v));
EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeFloat 16
%2 = OpTypeVector %3 2
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
index 77b4a62..d2af246 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
@@ -19,7 +19,7 @@
TEST_F(SpvGeneratorImplTest, Function_Empty) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -35,7 +35,7 @@
// Test that we do not emit the same function type more than once.
TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
generator_.EmitFunction(func);
generator_.EmitFunction(func);
@@ -48,7 +48,7 @@
TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Compute) {
auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
ir::Function::PipelineStage::kCompute, {{32, 4, 1}});
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main"
@@ -66,7 +66,7 @@
TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Fragment) {
auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
ir::Function::PipelineStage::kFragment);
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Fragment %1 "main"
@@ -84,7 +84,7 @@
TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Vertex) {
auto* func = b.CreateFunction(mod.symbols.Register("main"), mod.types.Get<type::Void>(),
ir::Function::PipelineStage::kVertex);
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Vertex %1 "main"
@@ -101,15 +101,15 @@
TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Multiple) {
auto* f1 = b.CreateFunction(mod.symbols.Register("main1"), mod.types.Get<type::Void>(),
ir::Function::PipelineStage::kCompute, {{32, 4, 1}});
- b.Branch(f1->start_target, f1->end_target);
+ f1->StartTarget()->BranchTo(f1->EndTarget());
auto* f2 = b.CreateFunction(mod.symbols.Register("main2"), mod.types.Get<type::Void>(),
ir::Function::PipelineStage::kCompute, {{8, 2, 16}});
- b.Branch(f2->start_target, f2->end_target);
+ f2->StartTarget()->BranchTo(f2->EndTarget());
auto* f3 = b.CreateFunction(mod.symbols.Register("main3"), mod.types.Get<type::Void>(),
ir::Function::PipelineStage::kFragment);
- b.Branch(f3->start_target, f3->end_target);
+ f3->StartTarget()->BranchTo(f3->EndTarget());
generator_.EmitFunction(f1);
generator_.EmitFunction(f2);
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
index 7b41184..fa8cc6b 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
@@ -23,11 +23,11 @@
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
auto* i = b.CreateIf(b.Constant(true));
- b.Branch(i->true_.target->As<ir::Block>(), i->merge.target);
- b.Branch(i->false_.target->As<ir::Block>(), i->merge.target);
- b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+ i->True().target->As<ir::Block>()->BranchTo(i->Merge().target);
+ i->False().target->As<ir::Block>()->BranchTo(i->Merge().target);
+ i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
- b.Branch(func->start_target, i);
+ func->StartTarget()->BranchTo(i);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -49,15 +49,15 @@
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
auto* i = b.CreateIf(b.Constant(true));
- b.Branch(i->false_.target->As<ir::Block>(), i->merge.target);
- b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+ i->False().target->As<ir::Block>()->BranchTo(i->Merge().target);
+ i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
- auto* true_block = i->true_.target->As<ir::Block>();
- true_block->instructions.Push(
- b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i)));
- b.Branch(true_block, i->merge.target);
+ auto* true_block = i->True().target->As<ir::Block>();
+ true_block->SetInstructions(
+ utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i))});
+ true_block->BranchTo(i->Merge().target);
- b.Branch(func->start_target, i);
+ func->StartTarget()->BranchTo(i);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -84,15 +84,15 @@
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
auto* i = b.CreateIf(b.Constant(true));
- b.Branch(i->true_.target->As<ir::Block>(), i->merge.target);
- b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+ i->True().target->As<ir::Block>()->BranchTo(i->Merge().target);
+ i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
- auto* false_block = i->false_.target->As<ir::Block>();
- false_block->instructions.Push(
- b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i)));
- b.Branch(false_block, i->merge.target);
+ auto* false_block = i->False().target->As<ir::Block>();
+ false_block->SetInstructions(
+ utils::Vector{b.Add(mod.types.Get<type::I32>(), b.Constant(1_i), b.Constant(1_i))});
+ false_block->BranchTo(i->Merge().target);
- b.Branch(func->start_target, i);
+ func->StartTarget()->BranchTo(i);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -119,11 +119,11 @@
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
auto* i = b.CreateIf(b.Constant(true));
- b.Branch(i->true_.target->As<ir::Block>(), func->end_target);
- b.Branch(i->false_.target->As<ir::Block>(), func->end_target);
- i->merge.target->As<ir::Block>()->branch.target = nullptr;
+ i->True().target->As<ir::Block>()->BranchTo(func->EndTarget());
+ i->False().target->As<ir::Block>()->BranchTo(func->EndTarget());
+ i->Merge().target->As<ir::Block>()->BranchTo(nullptr);
- b.Branch(func->start_target, i);
+ func->StartTarget()->BranchTo(i);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
index ff94862..20e5167 100644
--- a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
+++ b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
@@ -22,12 +22,11 @@
TEST_F(SpvGeneratorImplTest, FunctionVar_NoInit) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
auto* ty = mod.types.Get<type::Pointer>(
mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
- auto* v = b.Declare(ty);
- func->start_target->instructions.Push(v);
+ func->StartTarget()->SetInstructions(utils::Vector{b.Declare(ty)});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -45,13 +44,14 @@
TEST_F(SpvGeneratorImplTest, FunctionVar_WithInit) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
auto* ty = mod.types.Get<type::Pointer>(
mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
auto* v = b.Declare(ty);
- func->start_target->instructions.Push(v);
- v->initializer = b.Constant(42_i);
+ v->SetInitializer(b.Constant(42_i));
+
+ func->StartTarget()->SetInstructions(utils::Vector{v});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -71,12 +71,12 @@
TEST_F(SpvGeneratorImplTest, FunctionVar_Name) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
auto* ty = mod.types.Get<type::Pointer>(
mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
auto* v = b.Declare(ty);
- func->start_target->instructions.Push(v);
+ func->StartTarget()->SetInstructions(utils::Vector{v});
mod.SetName(v, "myvar");
generator_.EmitFunction(func);
@@ -96,22 +96,22 @@
TEST_F(SpvGeneratorImplTest, FunctionVar_DeclInsideBlock) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
auto* ty = mod.types.Get<type::Pointer>(
mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
auto* v = b.Declare(ty);
- v->initializer = b.Constant(42_i);
+ v->SetInitializer(b.Constant(42_i));
auto* i = b.CreateIf(b.Constant(true));
- b.Branch(i->false_.target->As<ir::Block>(), func->end_target);
- b.Branch(i->merge.target->As<ir::Block>(), func->end_target);
+ i->False().target->As<ir::Block>()->BranchTo(func->EndTarget());
+ i->Merge().target->As<ir::Block>()->BranchTo(func->EndTarget());
- auto* true_block = i->true_.target->As<ir::Block>();
- true_block->instructions.Push(v);
- b.Branch(true_block, i->merge.target);
+ auto* true_block = i->True().target->As<ir::Block>();
+ true_block->SetInstructions(utils::Vector{v});
+ true_block->BranchTo(i->Merge().target);
- b.Branch(func->start_target, i);
+ func->StartTarget()->BranchTo(i);
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@@ -138,15 +138,39 @@
)");
}
+TEST_F(SpvGeneratorImplTest, FunctionVar_Load) {
+ auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
+ func->StartTarget()->BranchTo(func->EndTarget());
+
+ auto* store_ty = mod.types.Get<type::I32>();
+ auto* ty = mod.types.Get<type::Pointer>(store_ty, builtin::AddressSpace::kFunction,
+ builtin::Access::kReadWrite);
+ auto* v = b.Declare(ty);
+ func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v)});
+
+ generator_.EmitFunction(func);
+ EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%7 = OpTypeInt 32 1
+%6 = OpTypePointer Function %7
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+%5 = OpVariable %6 Function
+%8 = OpLoad %7 %5
+OpReturn
+OpFunctionEnd
+)");
+}
+
TEST_F(SpvGeneratorImplTest, FunctionVar_Store) {
auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get<type::Void>());
- b.Branch(func->start_target, func->end_target);
+ func->StartTarget()->BranchTo(func->EndTarget());
auto* ty = mod.types.Get<type::Pointer>(
mod.types.Get<type::I32>(), builtin::AddressSpace::kFunction, builtin::Access::kReadWrite);
auto* v = b.Declare(ty);
- func->start_target->instructions.Push(v);
- func->start_target->instructions.Push(b.Store(v, b.Constant(42_i)));
+ func->StartTarget()->SetInstructions(utils::Vector{v, b.Store(v, b.Constant(42_i))});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
