encounter/dawn-cmake
1
0
Fork 0
mirror of https://github.com/encounter/dawn-cmake.git synced 2025-05-13 10:51:35 +00:00
Code Issues Packages Projects Releases Wiki Activity

[ir] Shift back to accessors.

Browse Source 
Some of the IR classes require setters in order to update dependant
information. In order to keep the IR access symmetrical this CL moves
the IR back to using accessors and private fields.

Bug: tint:1718
Change-Id: I101edda004671e07c4594bdcae4b1576e5771782
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/133640
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
Commit-Queue: Dan Sinclair <dsinclair@chromium.org>
This commit is contained in:
dan sinclair 2023-05-18 22:16:08 +00:00 committed by Dawn LUCI CQ
parent 3cd9b53e88
commit 24cb81116d
57 changed files with 1219 additions and 1050 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

1
src/tint/BUILD.gn
View File

@ -2265,6 +2265,7 @@ if (tint_build_unittests) {
"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",
]

1
src/tint/CMakeLists.txt
View File

@ -1476,6 +1476,7 @@ if(TINT_BUILD_TESTS)
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()

4
src/tint/ir/binary.cc
View File

@ -19,8 +19,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Binary);
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);

19
src/tint/ir/binary.h
View File

@ -51,7 +51,7 @@ class Binary : public utils::Castable<Binary, Instruction> {
/// @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 @@ class Binary : public utils::Castable<Binary, Instruction> {
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 @@ class Binary : public utils::Castable<Binary, Instruction> {
/// @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

107
src/tint/ir/binary_test.cc
View File

@ -30,17 +30,16 @@ TEST_F(IR_InstructionTest, CreateAnd) {
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 @@ TEST_F(IR_InstructionTest, CreateOr) {
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 @@ TEST_F(IR_InstructionTest, CreateXor) {
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 @@ TEST_F(IR_InstructionTest, CreateEqual) {
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 @@ TEST_F(IR_InstructionTest, CreateNotEqual) {
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 @@ TEST_F(IR_InstructionTest, CreateLessThan) {
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 @@ TEST_F(IR_InstructionTest, CreateGreaterThan) {
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 @@ TEST_F(IR_InstructionTest, CreateLessThanEqual) {
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 @@ TEST_F(IR_InstructionTest, CreateGreaterThanEqual) {
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 @@ TEST_F(IR_InstructionTest, CreateNot) {
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 @@ TEST_F(IR_InstructionTest, CreateShiftLeft) {
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 @@ TEST_F(IR_InstructionTest, CreateShiftRight) {
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 @@ TEST_F(IR_InstructionTest, CreateAdd) {
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 @@ TEST_F(IR_InstructionTest, CreateSubtract) {
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 @@ TEST_F(IR_InstructionTest, CreateMultiply) {
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 @@ TEST_F(IR_InstructionTest, CreateDivide) {
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 @@ TEST_F(IR_InstructionTest, CreateModulo) {
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 @@ TEST_F(IR_InstructionTest, Binary_Usage) {
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 @@ TEST_F(IR_InstructionTest, Binary_Usage_DuplicateValue) {
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);

16
src/tint/ir/bitcast_test.cc
View File

@ -32,9 +32,10 @@ TEST_F(IR_InstructionTest, Bitcast) {
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 @@ TEST_F(IR_InstructionTest, Bitcast_Usage) {
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

9
src/tint/ir/block.cc
View File

@ -22,4 +22,13 @@ Block::Block() : Base() {}
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

48
src/tint/ir/block.h
View File

@ -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 @@ class Block : public utils::Castable<Block, FlowNode> {
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

2
src/tint/ir/block_param.cc
View File

@ -18,7 +18,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::BlockParam);
namespace tint::ir {
BlockParam::BlockParam(const type::Type* ty) : type(ty) {}
BlockParam::BlockParam(const type::Type* ty) : type_(ty) {}
BlockParam::~BlockParam() = default;

5
src/tint/ir/block_param.h
View File

@ -34,10 +34,11 @@ class BlockParam : public utils::Castable<BlockParam, Value> {
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

51
src/tint/ir/builder.cc
View File

@ -32,7 +32,7 @@ ir::Block* Builder::CreateRootBlockIfNeeded() {
// 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 @@ Function* Builder::CreateFunction(Symbol name,
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 @@ If* Builder::CreateIf(Value* condition) {
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 @@ Binary* Builder::Modulo(const type::Type* type, Value* lhs, Value* rhs) {
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);
}

10
src/tint/ir/builder.h
View File

@ -98,12 +98,6 @@ class Builder {
/// @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
@ -161,7 +155,7 @@ class Builder {
/// @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
@ -280,7 +274,7 @@ class Builder {
/// @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

4
src/tint/ir/call.cc
View File

@ -21,8 +21,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Call);
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);
}
}

13
src/tint/ir/call.h
View File

@ -31,13 +31,10 @@ class Call : public utils::Castable<Call, Instruction> {
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 @@ class Call : public utils::Castable<Call, Instruction> {
/// @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

2
src/tint/ir/constant.cc
View File

@ -18,7 +18,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Constant);
namespace tint::ir {
Constant::Constant(const constant::Value* val) : value(val) {}
Constant::Constant(const constant::Value* val) : value_(val) {}
Constant::~Constant() = default;

16
src/tint/ir/constant.h
View File

@ -26,13 +26,21 @@ class Constant : public utils::Castable<Constant, Value> {
/// 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

62
src/tint/ir/constant_test.cc
View File

@ -30,13 +30,13 @@ TEST_F(IR_ConstantTest, f32) {
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 @@ TEST_F(IR_ConstantTest, f16) {
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 @@ TEST_F(IR_ConstantTest, i32) {
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 @@ TEST_F(IR_ConstantTest, u32) {
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 @@ TEST_F(IR_ConstantTest, bool) {
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>>());
}
}

74
src/tint/ir/debug.cc
View File

@ -60,81 +60,81 @@ std::string Debug::AsDotGraph(const Module* mod) {
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 @@ std::string Debug::AsDotGraph(const Module* mod) {
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 << "}";

148
src/tint/ir/disassembler.cc
View File

@ -81,7 +81,7 @@ utils::StringStream& Disassembler::Indent() {
}
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,31 +114,31 @@ void Disassembler::Walk(const FlowNode* node) {
[&](const ir::Function* f) {
TINT_SCOPED_ASSIGNMENT(in_function_, true);
Indent() << "%fn" << IdOf(f) << " = func " << f->name.Name() << "(";
for (auto* p : f->params) {
if (p != f->params.Front()) {
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_ << "%" << IdOf(p) << ":" << p->Type()->FriendlyName();
}
out_ << "):" << f->return_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() << ")";
}
}
}
@ -149,23 +149,23 @@ void Disassembler::Walk(const FlowNode* node) {
{
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);
@ -181,20 +181,20 @@ void Disassembler::Walk(const FlowNode* node) {
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);
@ -205,18 +205,18 @@ void Disassembler::Walk(const FlowNode* node) {
}
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: (";
@ -231,17 +231,17 @@ void Disassembler::Walk(const FlowNode* node) {
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()) {
@ -255,86 +255,86 @@ void Disassembler::Walk(const FlowNode* node) {
}
}
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*) {
@ -407,7 +407,7 @@ void Disassembler::EmitValue(const Value* val) {
}
});
};
emit(constant->value);
emit(constant->Value());
},
[&](const ir::Instruction* i) { out_ << "%" << IdOf(i); },
[&](const ir::BlockParam* p) {
@ -445,18 +445,18 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
[&](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);
@ -464,16 +464,16 @@ void Disassembler::EmitInstruction(const Instruction* inst) {
[&](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_ << ", ";
}
@ -485,7 +485,7 @@ void Disassembler::EmitArgs(const Call* call) {
void Disassembler::EmitBinary(const Binary* b) {
EmitValueWithType(b);
out_ << " = ";
switch (b->kind) {
switch (b->Kind()) {
case Binary::Kind::kAdd:
out_ << "add";
break;
@ -544,7 +544,7 @@ void Disassembler::EmitBinary(const Binary* b) {
void Disassembler::EmitUnary(const Unary* u) {
EmitValueWithType(u);
out_ << " = ";
switch (u->kind) {
switch (u->Kind()) {
case Unary::Kind::kComplement:
out_ << "complement";
break;

28
src/tint/ir/flow_node.h
View File

@ -25,22 +25,30 @@ class FlowNode : public utils::Castable<FlowNode> {
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

210
src/tint/ir/from_program.cc
View File

@ -98,17 +98,13 @@ namespace {
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;
}
@ -184,14 +180,14 @@ class Impl {
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 @@ class Impl {
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 @@ class Impl {
}
}
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 @@ class Impl {
->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,8 +323,9 @@ class Impl {
}
});
}
ir_func->SetReturnAttributes(return_attributes);
}
ir_func->return_location = sem->ReturnLocation();
ir_func->SetReturnLocation(sem->ReturnLocation());
scopes_.Push();
TINT_DEFER(scopes_.Pop());
@ -348,17 +340,17 @@ class Impl {
builder_.ir.SetName(param, p->name->symbol.NameView());
params.Push(param);
}
ir_func->params = std::move(params);
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);
// 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());
@ -372,7 +364,7 @@ class Impl {
// 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;
}
}
@ -427,7 +419,7 @@ class Impl {
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) {
@ -438,7 +430,7 @@ class Impl {
// 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();
@ -451,10 +443,10 @@ class Impl {
} 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) {
@ -470,7 +462,7 @@ class Impl {
// 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();
@ -520,10 +512,10 @@ class Impl {
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) {
@ -551,27 +543,27 @@ class Impl {
{
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>();
}
}
@ -585,7 +577,7 @@ class Impl {
{
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.
@ -594,21 +586,22 @@ class Impl {
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;
}
}
@ -616,9 +609,8 @@ class Impl {
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);
@ -627,9 +619,9 @@ class Impl {
{
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;
@ -637,31 +629,24 @@ class Impl {
// 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();
@ -679,7 +664,7 @@ class Impl {
{
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
@ -690,30 +675,25 @@ class Impl {
// 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) {
@ -745,13 +725,13 @@ class Impl {
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>();
}
}
@ -765,7 +745,7 @@ class Impl {
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*) {
@ -773,9 +753,9 @@ class Impl {
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_);
}
@ -786,7 +766,7 @@ class Impl {
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_);
}
@ -798,7 +778,7 @@ class Impl {
// 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) {
@ -819,17 +799,17 @@ class Impl {
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) {
@ -876,7 +856,7 @@ class Impl {
// 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;
}
@ -895,14 +875,14 @@ class Impl {
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);
@ -969,7 +949,7 @@ class Impl {
break;
}
current_flow_block_->instructions.Push(inst);
current_flow_block_->Instructions().Push(inst);
return inst;
}
@ -996,7 +976,7 @@ class Impl {
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;
{
@ -1010,17 +990,17 @@ class Impl {
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);
@ -1030,9 +1010,9 @@ class Impl {
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;
}
@ -1114,7 +1094,7 @@ class Impl {
return utils::Failure;
}
current_flow_block_->instructions.Push(inst);
current_flow_block_->Instructions().Push(inst);
return inst;
}
@ -1128,7 +1108,7 @@ class Impl {
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;
}
@ -1191,7 +1171,7 @@ class Impl {
if (inst == nullptr) {
return utils::Failure;
}
current_flow_block_->instructions.Push(inst);
current_flow_block_->Instructions().Push(inst);
return inst;
}

81
src/tint/ir/from_program_literal_test.cc
View File

@ -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 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_True) {
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,7 +67,7 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_False) {
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>());
}
@ -79,18 +81,19 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_Deduped) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* var_a = m.Get().root_block->instructions[0]->As<ir::Var>();
auto instr = m.Get().root_block->Instructions();
auto* var_a = instr[0]->As<ir::Var>();
ASSERT_NE(var_a, nullptr);
auto* var_b = m.Get().root_block->instructions[1]->As<ir::Var>();
auto* var_b = instr[1]->As<ir::Var>();
ASSERT_NE(var_b, nullptr);
auto* var_c = m.Get().root_block->instructions[2]->As<ir::Var>();
auto* var_c = instr[2]->As<ir::Var>();
ASSERT_NE(var_c, nullptr);
auto* var_d = m.Get().root_block->instructions[3]->As<ir::Var>();
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);
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) {
@ -102,7 +105,7 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_F32) {
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>());
}
@ -115,15 +118,16 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_F32_Deduped) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* var_a = m.Get().root_block->instructions[0]->As<ir::Var>();
auto instr = m.Get().root_block->Instructions();
auto* var_a = instr[0]->As<ir::Var>();
ASSERT_NE(var_a, nullptr);
auto* var_b = m.Get().root_block->instructions[1]->As<ir::Var>();
auto* var_b = instr[1]->As<ir::Var>();
ASSERT_NE(var_b, nullptr);
auto* var_c = m.Get().root_block->instructions[2]->As<ir::Var>();
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);
ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
ASSERT_NE(var_a->Initializer(), var_b->Initializer());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_F16) {
@ -136,7 +140,7 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_F16) {
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>());
}
@ -150,15 +154,16 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_F16_Deduped) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* var_a = m.Get().root_block->instructions[0]->As<ir::Var>();
auto instr = m.Get().root_block->Instructions();
auto* var_a = instr[0]->As<ir::Var>();
ASSERT_NE(var_a, nullptr);
auto* var_b = m.Get().root_block->instructions[1]->As<ir::Var>();
auto* var_b = instr[1]->As<ir::Var>();
ASSERT_NE(var_b, nullptr);
auto* var_c = m.Get().root_block->instructions[2]->As<ir::Var>();
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);
ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
ASSERT_NE(var_a->Initializer(), var_b->Initializer());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_I32) {
@ -170,7 +175,7 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_I32) {
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>());
}
@ -183,15 +188,16 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_I32_Deduped) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* var_a = m.Get().root_block->instructions[0]->As<ir::Var>();
auto instr = m.Get().root_block->Instructions();
auto* var_a = instr[0]->As<ir::Var>();
ASSERT_NE(var_a, nullptr);
auto* var_b = m.Get().root_block->instructions[1]->As<ir::Var>();
auto* var_b = instr[1]->As<ir::Var>();
ASSERT_NE(var_b, nullptr);
auto* var_c = m.Get().root_block->instructions[2]->As<ir::Var>();
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);
ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
ASSERT_NE(var_a->Initializer(), var_b->Initializer());
}
TEST_F(IR_BuilderImplTest, EmitLiteral_U32) {
@ -203,7 +209,7 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_U32) {
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>());
}
@ -216,15 +222,16 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_U32_Deduped) {
auto m = Build();
ASSERT_TRUE(m) << (!m ? m.Failure() : "");
auto* var_a = m.Get().root_block->instructions[0]->As<ir::Var>();
auto instr = m.Get().root_block->Instructions();
auto* var_a = instr[0]->As<ir::Var>();
ASSERT_NE(var_a, nullptr);
auto* var_b = m.Get().root_block->instructions[1]->As<ir::Var>();
auto* var_b = instr[1]->As<ir::Var>();
ASSERT_NE(var_b, nullptr);
auto* var_c = m.Get().root_block->instructions[2]->As<ir::Var>();
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);
ASSERT_EQ(var_a->Initializer(), var_c->Initializer());
ASSERT_NE(var_a->Initializer(), var_b->Initializer());
}
} // namespace

682
src/tint/ir/from_program_test.cc
View File

@ -63,13 +63,13 @@ TEST_F(IR_BuilderImplTest, Func) {
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 {
@ -88,13 +88,13 @@ TEST_F(IR_BuilderImplTest, Func_WithParam) {
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(%a:u32):u32 {
%fn2 = block {
@ -114,13 +114,13 @@ TEST_F(IR_BuilderImplTest, Func_WithMultipleParam) {
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(%a:u32, %b:i32, %c:bool):void {
%fn2 = block {
@ -137,7 +137,7 @@ TEST_F(IR_BuilderImplTest, EntryPoint) {
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) {
@ -148,19 +148,19 @@ TEST_F(IR_BuilderImplTest, IfStatement) {
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)] {
@ -192,19 +192,19 @@ TEST_F(IR_BuilderImplTest, IfStatement_TrueReturns) {
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)] {
@ -235,19 +235,19 @@ TEST_F(IR_BuilderImplTest, IfStatement_FalseReturns) {
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)] {
@ -278,19 +278,19 @@ TEST_F(IR_BuilderImplTest, IfStatement_BothReturn) {
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)] {
@ -318,15 +318,15 @@ TEST_F(IR_BuilderImplTest, IfStatement_JumpChainToMerge) {
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)] {
@ -367,19 +367,19 @@ TEST_F(IR_BuilderImplTest, Loop_WithBreak) {
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)] {
@ -408,28 +408,28 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinue) {
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)] {
@ -475,28 +475,28 @@ TEST_F(IR_BuilderImplTest, Loop_WithContinuing_BreakIf) {
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)] {
@ -586,28 +586,28 @@ TEST_F(IR_BuilderImplTest, Loop_WithReturn) {
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)] {
@ -648,19 +648,19 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn) {
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)] {
@ -693,25 +693,25 @@ TEST_F(IR_BuilderImplTest, Loop_WithOnlyReturn_ContinuingBreakIf) {
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)] {
@ -736,28 +736,28 @@ TEST_F(IR_BuilderImplTest, Loop_WithIf_BothBranchesBreak) {
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)] {
@ -807,96 +807,96 @@ TEST_F(IR_BuilderImplTest, Loop_Nested) {
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)] {
@ -1015,28 +1015,28 @@ TEST_F(IR_BuilderImplTest, While) {
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)] {
@ -1081,28 +1081,28 @@ TEST_F(IR_BuilderImplTest, While_Return) {
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)] {
@ -1155,28 +1155,28 @@ TEST_F(IR_BuilderImplTest, DISABLED_For) {
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"()");
}
@ -1189,18 +1189,18 @@ TEST_F(IR_BuilderImplTest, For_NoInitCondOrContinuing) {
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)] {
@ -1231,31 +1231,33 @@ TEST_F(IR_BuilderImplTest, Switch) {
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)] {
@ -1295,27 +1297,28 @@ TEST_F(IR_BuilderImplTest, Switch_MultiSelector) {
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)] {
@ -1343,19 +1346,20 @@ TEST_F(IR_BuilderImplTest, Switch_OnlyDefault) {
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)] {
@ -1385,26 +1389,27 @@ TEST_F(IR_BuilderImplTest, Switch_WithBreak) {
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)] {
@ -1441,25 +1446,26 @@ TEST_F(IR_BuilderImplTest, Switch_AllReturn) {
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)] {

4
src/tint/ir/function.cc
View File

@ -18,11 +18,11 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Function);
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;

88
src/tint/ir/function.h
View File

@ -17,6 +17,7 @@
#include <array>
#include <optional>
#include <utility>
#include "src/tint/ir/flow_node.h"
#include "src/tint/ir/function_param.h"
@ -71,32 +72,81 @@ class Function : public utils::Castable<Function, FlowNode> {
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 parameters to the function
utils::Vector<FunctionParam*, 1> params;
/// 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}; }
/// 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;
/// @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);

2
src/tint/ir/function_param.cc
View File

@ -18,7 +18,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::FunctionParam);
namespace tint::ir {
FunctionParam::FunctionParam(const type::Type* ty) : type(ty) {}
FunctionParam::FunctionParam(const type::Type* ty) : type_(ty) {}
FunctionParam::~FunctionParam() = default;

5
src/tint/ir/function_param.h
View File

@ -34,10 +34,11 @@ class FunctionParam : public utils::Castable<FunctionParam, Value> {
FunctionParam& operator=(FunctionParam&& 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 type of the parameter
const type::Type* type;
const type::Type* type_;
};
} // namespace tint::ir

2
src/tint/ir/if.cc
View File

@ -18,7 +18,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::If);
namespace tint::ir {
If::If(Value* cond) : Base(), condition(cond) {}
If::If(Value* cond) : Base(), condition_(cond) {}
If::~If() = default;

33
src/tint/ir/if.h
View File

@ -32,17 +32,36 @@ class If : public utils::Castable<If, FlowNode> {
/// 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

8
src/tint/ir/load.cc
View File

@ -19,10 +19,10 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Load);
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;

11
src/tint/ir/load.h
View File

@ -35,13 +35,14 @@ class Load : public utils::Castable<Load, Instruction> {
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

12
src/tint/ir/load_test.cc
View File

@ -33,12 +33,12 @@ TEST_F(IR_InstructionTest, CreateLoad) {
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 @@ TEST_F(IR_InstructionTest, Load_Usage) {
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

32
src/tint/ir/loop.h
View File

@ -26,16 +26,32 @@ class Loop : public utils::Castable<Loop, FlowNode> {
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

5
src/tint/ir/root_terminator.h
View File

@ -25,7 +25,12 @@ class RootTerminator : public utils::Castable<RootTerminator, FlowNode> {
public:
/// Constructor
RootTerminator();
RootTerminator(const RootTerminator&) = delete;
RootTerminator(RootTerminator&&) = delete;
~RootTerminator() override;
RootTerminator& operator=(const RootTerminator&) = delete;
RootTerminator& operator=(RootTerminator&&) = delete;
};
} // namespace tint::ir

10
src/tint/ir/store.cc
View File

@ -19,11 +19,11 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Store);
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;

13
src/tint/ir/store.h
View File

@ -34,10 +34,15 @@ class Store : public utils::Castable<Store, Instruction> {
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

18
src/tint/ir/store_test.cc
View File

@ -33,10 +33,10 @@ TEST_F(IR_InstructionTest, CreateStore) {
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 @@ TEST_F(IR_InstructionTest, Store_Usage) {
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

2
src/tint/ir/switch.cc
View File

@ -18,7 +18,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Switch);
namespace tint::ir {
Switch::Switch(Value* cond) : Base(), condition(cond) {}
Switch::Switch(Value* cond) : Base(), condition_(cond) {}
Switch::~Switch() = default;

31
src/tint/ir/switch.h
View File

@ -41,21 +41,40 @@ class Switch : public utils::Castable<Switch, FlowNode> {
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

1
src/tint/ir/test_helper.h
View File

@ -33,7 +33,6 @@ template <typename BASE>
class TestHelperBase : public BASE, public ProgramBuilder {
public:
TestHelperBase() = default;
~TestHelperBase() override = default;
/// Build the module, cleaning up the program before returning.

73
src/tint/ir/to_program.cc
View File

@ -91,14 +91,14 @@ class State {
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 @@ class State {
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 @@ class State {
stmts.Push(s);
}
}
branch = &block->branch;
branch = &block->Branch();
return kContinue;
},
@ -145,8 +145,8 @@ class State {
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 @@ class State {
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 @@ class State {
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 @@ class State {
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 @@ class State {
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 @@ class State {
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 @@ class State {
}
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 @@ class State {
}
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 @@ class State {
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)); }

4
src/tint/ir/transform/add_empty_entry_point.cc
View File

@ -29,7 +29,7 @@ AddEmptyEntryPoint::~AddEmptyEntryPoint() = default;
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 @@ void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
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);
}

2
src/tint/ir/transform/add_empty_entry_point_test.cc
View File

@ -40,7 +40,7 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
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"(

4
src/tint/ir/unary.cc
View File

@ -19,8 +19,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Unary);
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);
}

15
src/tint/ir/unary.h
View File

@ -33,7 +33,7 @@ class Unary : public utils::Castable<Unary, Instruction> {
/// @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 @@ class Unary : public utils::Castable<Unary, Instruction> {
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

10
src/tint/ir/unary_test.cc
View File

@ -29,10 +29,10 @@ TEST_F(IR_InstructionTest, CreateComplement) {
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 @@ TEST_F(IR_InstructionTest, CreateNegation) {
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 @@ TEST_F(IR_InstructionTest, Unary_Usage) {
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);

2
src/tint/ir/user_call.cc
View File

@ -23,7 +23,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::UserCall);
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;

7
src/tint/ir/user_call.h
View File

@ -36,8 +36,11 @@ class UserCall : public utils::Castable<UserCall, Call> {
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

8
src/tint/ir/var.cc
View File

@ -19,8 +19,14 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Var);
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

14
src/tint/ir/var.h
View File

@ -36,13 +36,17 @@ class Var : public utils::Castable<Var, Instruction> {
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

21
src/tint/transform/manager_test.cc
View File

@ -52,7 +52,7 @@ class IR_AddFunction final : public ir::transform::Transform {
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::Module MakeIR() {
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 @@ TEST_F(TransformManagerTest, IR_MutateInPlace) {
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 @@ TEST_F(TransformManagerTest, IR_MixedTransforms_AST_Before_IR) {
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 @@ TEST_F(TransformManagerTest, IR_MixedTransforms_IR_Before_AST) {
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

54
src/tint/writer/spirv/ir/generator_impl_ir.cc
View File

@ -110,7 +110,7 @@ bool GeneratorImplIr::Generate() {
}
uint32_t GeneratorImplIr::Constant(const ir::Constant* constant) {
return Constant(constant->value);
return Constant(constant->Value());
}
uint32_t GeneratorImplIr::Constant(const constant::Value* constant) {
@ -214,15 +214,15 @@ void GeneratorImplIr::EmitFunction(const ir::Function* func) {
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.
@ -248,7 +248,7 @@ void GeneratorImplIr::EmitFunction(const ir::Function* func) {
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_);
@ -256,13 +256,13 @@ void GeneratorImplIr::EmitFunction(const ir::Function* func) {
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: {
@ -282,7 +282,7 @@ void GeneratorImplIr::EmitEntryPoint(const ir::Function* func, uint32_t id) {
}
// 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) {
@ -293,7 +293,7 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
}
// 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); },
@ -313,43 +313,43 @@ void GeneratorImplIr::EmitBlock(const ir::Block* block) {
// 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);
}
@ -357,7 +357,7 @@ void GeneratorImplIr::EmitIf(const ir::If* i) {
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) {
@ -376,7 +376,7 @@ uint32_t GeneratorImplIr::EmitBinary(const ir::Binary* binary) {
// 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;
@ -387,7 +387,7 @@ uint32_t GeneratorImplIr::EmitBinary(const ir::Binary* binary) {
}
default: {
TINT_ICE(Writer, diagnostics_)
<< "unimplemented binary instruction: " << static_cast<uint32_t>(binary->kind);
<< "unimplemented binary instruction: " << static_cast<uint32_t>(binary->Kind());
}
}
@ -400,12 +400,12 @@ uint32_t GeneratorImplIr::EmitBinary(const ir::Binary* binary) {
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)});
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) {
@ -417,8 +417,8 @@ uint32_t GeneratorImplIr::EmitVar(const ir::Var* var) {
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_)

69
src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc
View File

@ -21,10 +21,10 @@ namespace {
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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_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);
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 @@ OpFunctionEnd
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_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_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 @@ OpFunctionEnd
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"

14
src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc
View File

@ -67,7 +67,7 @@ TEST_F(SpvGeneratorImplTest, Constant_Vec4Bool) {
auto* f = b.Constant(false);
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
@ -82,7 +82,8 @@ TEST_F(SpvGeneratorImplTest, Constant_Vec2i) {
auto* i_42 = b.Constant(i32(42));
auto* i_n1 = b.Constant(i32(-1));
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);
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
@ -98,8 +99,8 @@ TEST_F(SpvGeneratorImplTest, Constant_Vec3u) {
auto* u_0 = b.Constant(u32(0));
auto* u_4b = b.Constant(u32(4000000000));
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);
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
@ -118,7 +119,7 @@ TEST_F(SpvGeneratorImplTest, Constant_Vec4f) {
auto* f_n1 = b.Constant(f32(-1));
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
@ -135,7 +136,8 @@ TEST_F(SpvGeneratorImplTest, Constant_Vec2h) {
auto* h_42 = b.Constant(f16(42));
auto* h_q = b.Constant(f16(0.25));
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);
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

16
src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc
View File

@ -19,7 +19,7 @@ namespace {
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 @@ OpFunctionEnd
// 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_DeduplicateType) {
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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);

44
src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc
View File

@ -23,11 +23,11 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) {
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 @@ TEST_F(SpvGeneratorImplTest, If_FalseEmpty) {
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 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty) {
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 @@ TEST_F(SpvGeneratorImplTest, If_BothBranchesReturn) {
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"

42
src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc
View File

@ -22,12 +22,11 @@ namespace {
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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 @@ OpFunctionEnd
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"
@ -140,14 +140,13 @@ OpFunctionEnd
TEST_F(SpvGeneratorImplTest, FunctionVar_Load) {
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* 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->start_target->instructions.Push(v);
func->start_target->instructions.Push(b.Load(v));
func->StartTarget()->SetInstructions(utils::Vector{v, b.Load(v)});
generator_.EmitFunction(func);
EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo"
@ -166,13 +165,12 @@ 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"