diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 1358b9c2ff..64fde4cea4 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -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", ] diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index a696399be3..8ec573a376 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -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() diff --git a/src/tint/ir/binary.cc b/src/tint/ir/binary.cc index af28306007..2b179acc36 100644 --- a/src/tint/ir/binary.cc +++ b/src/tint/ir/binary.cc @@ -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); diff --git a/src/tint/ir/binary.h b/src/tint/ir/binary.h index a6dc722262..7ea6c08de4 100644 --- a/src/tint/ir/binary.h +++ b/src/tint/ir/binary.h @@ -51,7 +51,7 @@ class Binary : public utils::Castable { /// @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& 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 { /// @returns the right-hand-side value for the instruction const Value* RHS() const { return rhs_; } - /// the kind of binary instruction - Kind kind = Kind::kAdd; - - /// the result type of the instruction - const type::Type* result_type = nullptr; - private: - Value* lhs_ = nullptr; - Value* rhs_ = nullptr; + enum Kind kind_; + const type::Type* result_type_; + Value* lhs_; + Value* rhs_; }; } // namespace tint::ir diff --git a/src/tint/ir/binary_test.cc b/src/tint/ir/binary_test.cc index 2968d283a3..dc359781c3 100644 --- a/src/tint/ir/binary_test.cc +++ b/src/tint/ir/binary_test.cc @@ -30,17 +30,16 @@ TEST_F(IR_InstructionTest, CreateAnd) { const auto* inst = b.And(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - 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()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -52,15 +51,15 @@ TEST_F(IR_InstructionTest, CreateOr) { const auto* inst = b.Or(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kOr); + EXPECT_EQ(inst->Kind(), Binary::Kind::kOr); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -72,15 +71,15 @@ TEST_F(IR_InstructionTest, CreateXor) { const auto* inst = b.Xor(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kXor); + EXPECT_EQ(inst->Kind(), Binary::Kind::kXor); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -92,15 +91,15 @@ TEST_F(IR_InstructionTest, CreateEqual) { const auto* inst = b.Equal(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kEqual); + EXPECT_EQ(inst->Kind(), Binary::Kind::kEqual); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -112,15 +111,15 @@ TEST_F(IR_InstructionTest, CreateNotEqual) { const auto* inst = b.NotEqual(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kNotEqual); + EXPECT_EQ(inst->Kind(), Binary::Kind::kNotEqual); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -132,15 +131,15 @@ TEST_F(IR_InstructionTest, CreateLessThan) { const auto* inst = b.LessThan(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kLessThan); + EXPECT_EQ(inst->Kind(), Binary::Kind::kLessThan); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -153,15 +152,15 @@ TEST_F(IR_InstructionTest, CreateGreaterThan) { b.GreaterThan(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThan); + EXPECT_EQ(inst->Kind(), Binary::Kind::kGreaterThan); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -174,15 +173,15 @@ TEST_F(IR_InstructionTest, CreateLessThanEqual) { b.LessThanEqual(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kLessThanEqual); + EXPECT_EQ(inst->Kind(), Binary::Kind::kLessThanEqual); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -195,15 +194,15 @@ TEST_F(IR_InstructionTest, CreateGreaterThanEqual) { b.GreaterThanEqual(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kGreaterThanEqual); + EXPECT_EQ(inst->Kind(), Binary::Kind::kGreaterThanEqual); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -214,15 +213,15 @@ TEST_F(IR_InstructionTest, CreateNot) { const auto* inst = b.Not(b.ir.types.Get(), b.Constant(true)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kEqual); + EXPECT_EQ(inst->Kind(), Binary::Kind::kEqual); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_TRUE(lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_FALSE(rhs->As>()->ValueAs()); } @@ -234,15 +233,15 @@ TEST_F(IR_InstructionTest, CreateShiftLeft) { const auto* inst = b.ShiftLeft(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kShiftLeft); + EXPECT_EQ(inst->Kind(), Binary::Kind::kShiftLeft); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -254,15 +253,15 @@ TEST_F(IR_InstructionTest, CreateShiftRight) { const auto* inst = b.ShiftRight(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kShiftRight); + EXPECT_EQ(inst->Kind(), Binary::Kind::kShiftRight); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -274,15 +273,15 @@ TEST_F(IR_InstructionTest, CreateAdd) { const auto* inst = b.Add(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kAdd); + EXPECT_EQ(inst->Kind(), Binary::Kind::kAdd); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -294,15 +293,15 @@ TEST_F(IR_InstructionTest, CreateSubtract) { const auto* inst = b.Subtract(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kSubtract); + EXPECT_EQ(inst->Kind(), Binary::Kind::kSubtract); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -314,15 +313,15 @@ TEST_F(IR_InstructionTest, CreateMultiply) { const auto* inst = b.Multiply(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kMultiply); + EXPECT_EQ(inst->Kind(), Binary::Kind::kMultiply); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -334,15 +333,15 @@ TEST_F(IR_InstructionTest, CreateDivide) { const auto* inst = b.Divide(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kDivide); + EXPECT_EQ(inst->Kind(), Binary::Kind::kDivide); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -354,15 +353,15 @@ TEST_F(IR_InstructionTest, CreateModulo) { const auto* inst = b.Modulo(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Binary::Kind::kModulo); + EXPECT_EQ(inst->Kind(), Binary::Kind::kModulo); ASSERT_TRUE(inst->LHS()->Is()); - auto lhs = inst->LHS()->As()->value; + auto lhs = inst->LHS()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); ASSERT_TRUE(inst->RHS()->Is()); - auto rhs = inst->RHS()->As()->value; + auto rhs = inst->RHS()->As()->Value(); ASSERT_TRUE(rhs->Is>()); EXPECT_EQ(2_i, rhs->As>()->ValueAs()); } @@ -372,7 +371,7 @@ TEST_F(IR_InstructionTest, Binary_Usage) { Builder b{mod}; const auto* inst = b.And(b.ir.types.Get(), 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(), val, val); - EXPECT_EQ(inst->kind, Binary::Kind::kAnd); + EXPECT_EQ(inst->Kind(), Binary::Kind::kAnd); ASSERT_EQ(inst->LHS(), inst->RHS()); ASSERT_NE(inst->LHS(), nullptr); diff --git a/src/tint/ir/bitcast_test.cc b/src/tint/ir/bitcast_test.cc index e78b734795..6eda56216c 100644 --- a/src/tint/ir/bitcast_test.cc +++ b/src/tint/ir/bitcast_test.cc @@ -32,9 +32,10 @@ TEST_F(IR_InstructionTest, Bitcast) { ASSERT_TRUE(inst->Is()); ASSERT_NE(inst->Type(), nullptr); - ASSERT_EQ(inst->args.Length(), 1u); - ASSERT_TRUE(inst->args[0]->Is()); - auto val = inst->args[0]->As()->value; + const auto args = inst->Args(); + ASSERT_EQ(args.Length(), 1u); + ASSERT_TRUE(args[0]->Is()); + auto val = args[0]->As()->Value(); ASSERT_TRUE(val->Is>()); EXPECT_EQ(4_i, val->As>()->ValueAs()); } @@ -44,10 +45,11 @@ TEST_F(IR_InstructionTest, Bitcast_Usage) { Builder b{mod}; const auto* inst = b.Bitcast(b.ir.types.Get(), b.Constant(4_i)); - ASSERT_EQ(inst->args.Length(), 1u); - ASSERT_NE(inst->args[0], nullptr); - ASSERT_EQ(inst->args[0]->Usage().Length(), 1u); - EXPECT_EQ(inst->args[0]->Usage()[0], inst); + const auto args = inst->Args(); + ASSERT_EQ(args.Length(), 1u); + ASSERT_NE(args[0], nullptr); + ASSERT_EQ(args[0]->Usage().Length(), 1u); + EXPECT_EQ(args[0]->Usage()[0], inst); } } // namespace diff --git a/src/tint/ir/block.cc b/src/tint/ir/block.cc index 2030c19c66..f5b5bfb077 100644 --- a/src/tint/ir/block.cc +++ b/src/tint/ir/block.cc @@ -22,4 +22,13 @@ Block::Block() : Base() {} Block::~Block() = default; +void Block::BranchTo(FlowNode* to, utils::VectorRef args) { + TINT_ASSERT(IR, to); + branch_.target = to; + branch_.args = args; + if (to) { + to->AddInboundBranch(this); + } +} + } // namespace tint::ir diff --git a/src/tint/ir/block.h b/src/tint/ir/block.h index 5af32fedc1..2b01fd00b1 100644 --- a/src/tint/ir/block.h +++ b/src/tint/ir/block.h @@ -15,6 +15,8 @@ #ifndef SRC_TINT_IR_BLOCK_H_ #define SRC_TINT_IR_BLOCK_H_ +#include + #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 { 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 args = {}); - /// The instructions in the block - utils::Vector 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 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 instructions) { + instructions_ = std::move(instructions); + } + + /// @returns the instructions in the block + utils::VectorRef Instructions() const { return instructions_; } + /// @returns the instructions in the block + utils::Vector& Instructions() { return instructions_; } + + /// Sets the params to the block + /// @param params the params for the block + void SetParams(utils::VectorRef params) { params_ = std::move(params); } + /// @returns the params to the block + utils::Vector& Params() { return params_; } + + /// @return the parameters passed into the block + utils::VectorRef Params() const { return params_; } + + private: + ir::Branch branch_ = {}; + utils::Vector instructions_; + utils::Vector params_; }; } // namespace tint::ir diff --git a/src/tint/ir/block_param.cc b/src/tint/ir/block_param.cc index f014d196ac..a3a0be5d74 100644 --- a/src/tint/ir/block_param.cc +++ b/src/tint/ir/block_param.cc @@ -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; diff --git a/src/tint/ir/block_param.h b/src/tint/ir/block_param.h index 8ba68a75d5..036ddbf22b 100644 --- a/src/tint/ir/block_param.h +++ b/src/tint/ir/block_param.h @@ -34,10 +34,11 @@ class BlockParam : public utils::Castable { BlockParam& operator=(BlockParam&& inst) = delete; /// @returns the type of the var - const type::Type* Type() const override { return type; } + const type::Type* Type() const override { return type_; } + private: /// the result type of the instruction - const type::Type* type = nullptr; + const type::Type* type_; }; } // namespace tint::ir diff --git a/src/tint/ir/builder.cc b/src/tint/ir/builder.cc index af3c93b053..b0b04d676d 100644 --- a/src/tint/ir/builder.cc +++ b/src/tint/ir/builder.cc @@ -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(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(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(); - 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(condition); - ir_switch->merge.target = CreateBlock(); + ir_switch->Merge().target = CreateBlock(); return ir_switch; } Block* Builder::CreateCase(Switch* s, utils::VectorRef 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* b = s->Cases().Back().Start().target->As(); // 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 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(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(kind, type, val); } diff --git a/src/tint/ir/builder.h b/src/tint/ir/builder.h index 6de608eda9..ebd7a87280 100644 --- a/src/tint/ir/builder.h +++ b/src/tint/ir/builder.h @@ -98,12 +98,6 @@ class Builder { /// @returns the start block for the case flow node Block* CreateCase(Switch* s, utils::VectorRef 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 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 diff --git a/src/tint/ir/call.cc b/src/tint/ir/call.cc index dd503eadaa..909079f6e7 100644 --- a/src/tint/ir/call.cc +++ b/src/tint/ir/call.cc @@ -21,8 +21,8 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::Call); namespace tint::ir { Call::Call(const type::Type* res_ty, utils::VectorRef arguments) - : result_type(res_ty), args(std::move(arguments)) { - for (auto* arg : args) { + : result_type_(res_ty), args_(std::move(arguments)) { + for (auto* arg : args_) { arg->AddUsage(this); } } diff --git a/src/tint/ir/call.h b/src/tint/ir/call.h index 5393810af5..f4e12f913b 100644 --- a/src/tint/ir/call.h +++ b/src/tint/ir/call.h @@ -31,13 +31,10 @@ class Call : public utils::Castable { 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 args; + /// @returns the call arguments + utils::VectorRef Args() const { return args_; } protected: /// Constructor @@ -46,6 +43,10 @@ class Call : public utils::Castable { /// @param result_type the result type /// @param args the constructor arguments Call(const type::Type* result_type, utils::VectorRef args); + + private: + const type::Type* result_type_; + utils::Vector args_; }; } // namespace tint::ir diff --git a/src/tint/ir/constant.cc b/src/tint/ir/constant.cc index 8b5260c667..a49d140241 100644 --- a/src/tint/ir/constant.cc +++ b/src/tint/ir/constant.cc @@ -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; diff --git a/src/tint/ir/constant.h b/src/tint/ir/constant.h index 68e0dc7b92..dcf3e5b2f8 100644 --- a/src/tint/ir/constant.h +++ b/src/tint/ir/constant.h @@ -26,13 +26,21 @@ class Constant : public utils::Castable { /// Constructor /// @param val the value stored in the constant explicit Constant(const constant::Value* val); + Constant(const Constant&) = delete; + Constant(Constant&&) = delete; ~Constant() override; - /// @returns the type of the constant - const type::Type* Type() const override { return value->Type(); } + Constant& operator=(const Constant&) = delete; + Constant& operator=(Constant&&) = delete; - /// The constants value - const constant::Value* const value; + /// @returns the constants value + const constant::Value* Value() const { return value_; } + + /// @returns the type of the constant + const type::Type* Type() const override { return value_->Type(); } + + private: + const constant::Value* const value_; }; } // namespace tint::ir diff --git a/src/tint/ir/constant_test.cc b/src/tint/ir/constant_test.cc index 7005751913..7c3024a322 100644 --- a/src/tint/ir/constant_test.cc +++ b/src/tint/ir/constant_test.cc @@ -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>()->ValueAs()); + EXPECT_EQ(1.2_f, c->Value()->As>()->ValueAs()); - EXPECT_TRUE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); + EXPECT_TRUE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); } 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>()->ValueAs()); + EXPECT_EQ(1.1_h, c->Value()->As>()->ValueAs()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_TRUE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_TRUE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); } 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>()->ValueAs()); + EXPECT_EQ(1_i, c->Value()->As>()->ValueAs()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_TRUE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_TRUE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); } 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>()->ValueAs()); + EXPECT_EQ(2_u, c->Value()->As>()->ValueAs()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_TRUE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_TRUE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); } 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>()->ValueAs()); + EXPECT_FALSE(c->Value()->As>()->ValueAs()); } { utils::StringStream str; auto c = b.Constant(true); - EXPECT_TRUE(c->value->As>()->ValueAs()); + EXPECT_TRUE(c->Value()->As>()->ValueAs()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_FALSE(c->value->Is>()); - EXPECT_TRUE(c->value->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_FALSE(c->Value()->Is>()); + EXPECT_TRUE(c->Value()->Is>()); } } diff --git a/src/tint/ir/debug.cc b/src/tint/ir/debug.cc index 655c45507e..85666ff628 100644 --- a/src/tint/ir/debug.cc +++ b/src/tint/ir/debug.cc @@ -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 << "}"; diff --git a/src/tint/ir/disassembler.cc b/src/tint/ir/disassembler.cc index ef5c20b79a..76af74091b 100644 --- a/src/tint/ir/disassembler.cc +++ b/src/tint/ir/disassembler.cc @@ -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()) { + if (b->Branch().target->Is()) { out_ << " -> %func_end"; suffix = "return"; - } else if (b->branch.target->Is()) { + } else if (b->Branch().target->Is()) { // 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()) { + if (!b->Branch().target->Is()) { 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; diff --git a/src/tint/ir/flow_node.h b/src/tint/ir/flow_node.h index 905f077605..289873b1c5 100644 --- a/src/tint/ir/flow_node.h +++ b/src/tint/ir/flow_node.h @@ -25,22 +25,30 @@ class FlowNode : public utils::Castable { 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 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 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 inbound_branches_; }; } // namespace tint::ir diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc index 11bd4028fb..595db584d7 100644 --- a/src/tint/ir/from_program.cc +++ b/src/tint/ir/from_program.cc @@ -98,17 +98,13 @@ namespace { using ResultType = utils::Result; -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()) { 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 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 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>()) { 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(); + current_flow_block_ = if_node->True().target->As(); 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(); + current_flow_block_ = if_node->False().target->As(); 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(); + if (IsConnected(if_node->Merge().target)) { + current_flow_block_ = if_node->Merge().target->As(); } } @@ -585,7 +577,7 @@ class Impl { { FlowStackScope scope(this, loop_node); - current_flow_block_ = loop_node->start.target->As(); + current_flow_block_ = loop_node->Start().target->As(); // 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(); + current_flow_block_ = loop_node->Continuing().target->As(); 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(); - 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(); + 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()); - builder_.Branch(loop_node->continuing.target->As(), loop_node->start.target, - utils::Empty); + TINT_ASSERT(IR, loop_node->Continuing().target->Is()); + loop_node->Continuing().target->As()->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(); + current_flow_block_ = loop_node->Start().target->As(); - // 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()); - builder_.Branch(if_node->true_.target->As(), if_node->merge.target, - utils::Empty); - - TINT_ASSERT(IR, if_node->false_.target->Is()); - builder_.Branch(if_node->false_.target->As(), loop_node->merge.target, - utils::Empty); + if_node->True().target->As()->BranchTo(if_node->Merge().target); + if_node->False().target->As()->BranchTo(loop_node->Merge().target); BranchTo(if_node); - current_flow_block_ = if_node->merge.target->As(); + current_flow_block_ = if_node->Merge().target->As(); 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(); + // 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(); } void EmitForLoop(const ast::ForLoopStatement* stmt) { auto* loop_node = builder_.CreateLoop(); - TINT_ASSERT(IR, loop_node->continuing.target->Is()); - builder_.Branch(loop_node->continuing.target->As(), loop_node->start.target, - utils::Empty); + loop_node->Continuing().target->As()->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(); + current_flow_block_ = loop_node->Start().target->As(); 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()); - builder_.Branch(if_node->true_.target->As(), if_node->merge.target, - utils::Empty); - - TINT_ASSERT(IR, if_node->false_.target->Is()); - builder_.Branch(if_node->false_.target->As(), loop_node->merge.target, - utils::Empty); + if_node->True().target->As()->BranchTo(if_node->Merge().target); + if_node->False().target->As()->BranchTo(loop_node->Merge().target); BranchTo(if_node); - current_flow_block_ = if_node->merge.target->As(); + current_flow_block_ = if_node->Merge().target->As(); } EmitBlock(stmt->body); - BranchToIfNeeded(loop_node->continuing.target); + BranchToIfNeeded(loop_node->Continuing().target); if (stmt->continuing) { - current_flow_block_ = loop_node->continuing.target->As(); + current_flow_block_ = loop_node->Continuing().target->As(); 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(); + // 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(); } 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(); + if (IsConnected(switch_node->Merge().target)) { + current_flow_block_ = switch_node->Merge().target->As(); } } @@ -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()) { - BranchTo(c->merge.target); + BranchTo(c->Merge().target); } else if (auto* s = current_control->As()) { - 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()) { - 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(); - current_flow_block_ = if_node->true_.target->As(); - BranchTo(loop->merge.target); + current_flow_block_ = if_node->True().target->As(); + BranchTo(loop->Merge().target); - current_flow_block_ = if_node->false_.target->As(); - BranchTo(if_node->merge.target); + current_flow_block_ = if_node->False().target->As(); + BranchTo(if_node->Merge().target); - current_flow_block_ = if_node->merge.target->As(); + current_flow_block_ = if_node->Merge().target->As(); // 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 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()) { 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()); - if_node->merge.target->As()->params.Push(result); + if_node->Merge().target->As()->SetParams(utils::Vector{result}); utils::Result 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(); - BranchTo(if_node->merge.target, std::move(alt_args)); + current_flow_block_ = if_node->False().target->As(); + BranchTo(if_node->Merge().target, std::move(alt_args)); - current_flow_block_ = if_node->true_.target->As(); + current_flow_block_ = if_node->True().target->As(); } 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(); - BranchTo(if_node->merge.target, std::move(alt_args)); + current_flow_block_ = if_node->True().target->As(); + BranchTo(if_node->Merge().target, std::move(alt_args)); - current_flow_block_ = if_node->false_.target->As(); + current_flow_block_ = if_node->False().target->As(); } rhs = EmitExpression(expr->rhs); @@ -1030,9 +1010,9 @@ class Impl { utils::Vector 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(); + current_flow_block_ = if_node->Merge().target->As(); 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; } diff --git a/src/tint/ir/from_program_literal_test.cc b/src/tint/ir/from_program_literal_test.cc index 39b6052207..f746bc00be 100644 --- a/src/tint/ir/from_program_literal_test.cc +++ b/src/tint/ir/from_program_literal_test.cc @@ -25,17 +25,19 @@ namespace tint::ir { namespace { -Value* GlobalVarInitializer(const Module& m) { - if (m.root_block->instructions.Length() == 0u) { +const Value* GlobalVarInitializer(const Module& m) { + const auto instr = m.root_block->Instructions(); + + if (instr.Length() == 0u) { ADD_FAILURE() << "m.root_block has no instruction"; return nullptr; } - auto* var = m.root_block->instructions[0]->As(); + auto* var = instr[0]->As(); 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(init)); - auto* val = init->As()->value; + auto* val = init->As()->Value(); EXPECT_TRUE(val->Is>()); EXPECT_TRUE(val->As>()->ValueAs()); } @@ -65,7 +67,7 @@ TEST_F(IR_BuilderImplTest, EmitLiteral_Bool_False) { auto* init = GlobalVarInitializer(m.Get()); ASSERT_TRUE(Is(init)); - auto* val = init->As()->value; + auto* val = init->As()->Value(); EXPECT_TRUE(val->Is>()); EXPECT_FALSE(val->As>()->ValueAs()); } @@ -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(); + auto instr = m.Get().root_block->Instructions(); + auto* var_a = instr[0]->As(); ASSERT_NE(var_a, nullptr); - auto* var_b = m.Get().root_block->instructions[1]->As(); + auto* var_b = instr[1]->As(); ASSERT_NE(var_b, nullptr); - auto* var_c = m.Get().root_block->instructions[2]->As(); + auto* var_c = instr[2]->As(); ASSERT_NE(var_c, nullptr); - auto* var_d = m.Get().root_block->instructions[3]->As(); + auto* var_d = instr[3]->As(); 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(init)); - auto* val = init->As()->value; + auto* val = init->As()->Value(); EXPECT_TRUE(val->Is>()); EXPECT_EQ(1.2_f, val->As>()->ValueAs()); } @@ -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(); + auto instr = m.Get().root_block->Instructions(); + auto* var_a = instr[0]->As(); ASSERT_NE(var_a, nullptr); - auto* var_b = m.Get().root_block->instructions[1]->As(); + auto* var_b = instr[1]->As(); ASSERT_NE(var_b, nullptr); - auto* var_c = m.Get().root_block->instructions[2]->As(); + auto* var_c = instr[2]->As(); 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(init)); - auto* val = init->As()->value; + auto* val = init->As()->Value(); EXPECT_TRUE(val->Is>()); EXPECT_EQ(1.2_h, val->As>()->ValueAs()); } @@ -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(); + auto instr = m.Get().root_block->Instructions(); + auto* var_a = instr[0]->As(); ASSERT_NE(var_a, nullptr); - auto* var_b = m.Get().root_block->instructions[1]->As(); + auto* var_b = instr[1]->As(); ASSERT_NE(var_b, nullptr); - auto* var_c = m.Get().root_block->instructions[2]->As(); + auto* var_c = instr[2]->As(); 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(init)); - auto* val = init->As()->value; + auto* val = init->As()->Value(); EXPECT_TRUE(val->Is>()); EXPECT_EQ(-2_i, val->As>()->ValueAs()); } @@ -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(); + auto instr = m.Get().root_block->Instructions(); + auto* var_a = instr[0]->As(); ASSERT_NE(var_a, nullptr); - auto* var_b = m.Get().root_block->instructions[1]->As(); + auto* var_b = instr[1]->As(); ASSERT_NE(var_b, nullptr); - auto* var_c = m.Get().root_block->instructions[2]->As(); + auto* var_c = instr[2]->As(); 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(init)); - auto* val = init->As()->value; + auto* val = init->As()->Value(); EXPECT_TRUE(val->Is>()); EXPECT_EQ(2_u, val->As>()->ValueAs()); } @@ -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(); + auto instr = m.Get().root_block->Instructions(); + auto* var_a = instr[0]->As(); ASSERT_NE(var_a, nullptr); - auto* var_b = m.Get().root_block->instructions[1]->As(); + auto* var_b = instr[1]->As(); ASSERT_NE(var_b, nullptr); - auto* var_c = m.Get().root_block->instructions[2]->As(); + auto* var_c = instr[2]->As(); 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 diff --git a/src/tint/ir/from_program_test.cc b/src/tint/ir/from_program_test.cc index 773fde6642..1049386c3e 100644 --- a/src/tint/ir/from_program_test.cc +++ b/src/tint/ir/from_program_test.cc @@ -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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(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(node)) { - return block->branch.target; + return block->Branch().target; } return nullptr; }; - auto* loop_flow_a = As(m->functions[0]->start_target->branch.target); + auto* loop_flow_a = As(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(block_exit(loop_flow_a->start.target)); + auto* loop_flow_b = As(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(block_exit(loop_flow_b->start.target)); + auto* if_flow_a = As(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(block_exit(if_flow_a->merge.target)); + auto* if_flow_b = As(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(block_exit(loop_flow_b->continuing.target)); + auto* loop_flow_c = As(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(block_exit(loop_flow_c->merge.target)); + auto* loop_flow_d = As(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(block_exit(loop_flow_d->continuing.target)); + auto* if_flow_c = As(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(block_exit(loop_flow_b->merge.target)); + auto* if_flow_d = As(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(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()->branch.target, nullptr); - ASSERT_TRUE(flow->start.target->As()->branch.target->Is()); - auto* if_flow = flow->start.target->As()->branch.target->As(); - 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()->Branch().target, nullptr); + ASSERT_TRUE(flow->Start().target->As()->Branch().target->Is()); + auto* if_flow = flow->Start().target->As()->Branch().target->As(); + 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(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()->branch.target, nullptr); - ASSERT_TRUE(flow->start.target->As()->branch.target->Is()); - auto* if_flow = flow->start.target->As()->branch.target->As(); - 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()->Branch().target, nullptr); + ASSERT_TRUE(flow->Start().target->As()->Branch().target->Is()); + auto* if_flow = flow->Start().target->As()->Branch().target->As(); + 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(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()->branch.target, nullptr); - ASSERT_TRUE(flow->start.target->As()->branch.target->Is()); - auto* if_flow = flow->start.target->As()->branch.target->As(); - 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()->Branch().target, nullptr); + ASSERT_TRUE(flow->Start().target->As()->Branch().target->Is()); + auto* if_flow = flow->Start().target->As()->Branch().target->As(); + 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(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(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>()); + 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>()); EXPECT_EQ(0_i, - flow->cases[0].selectors[0].val->value->As>()->ValueOf()); + cases[0].selectors[0].val->Value()->As>()->ValueOf()); - ASSERT_EQ(1u, flow->cases[1].selectors.Length()); - ASSERT_TRUE(flow->cases[1].selectors[0].val->value->Is>()); + ASSERT_EQ(1u, cases[1].selectors.Length()); + ASSERT_TRUE(cases[1].selectors[0].val->Value()->Is>()); EXPECT_EQ(1_i, - flow->cases[1].selectors[0].val->value->As>()->ValueOf()); + cases[1].selectors[0].val->Value()->As>()->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(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>()); + 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>()); EXPECT_EQ(0_i, - flow->cases[0].selectors[0].val->value->As>()->ValueOf()); + cases[0].selectors[0].val->Value()->As>()->ValueOf()); - ASSERT_TRUE(flow->cases[0].selectors[1].val->value->Is>()); + ASSERT_TRUE(cases[0].selectors[1].val->Value()->Is>()); EXPECT_EQ(1_i, - flow->cases[0].selectors[1].val->value->As>()->ValueOf()); + cases[0].selectors[1].val->Value()->As>()->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(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(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>()); + 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>()); EXPECT_EQ(0_i, - flow->cases[0].selectors[0].val->value->As>()->ValueOf()); + cases[0].selectors[0].val->Value()->As>()->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(m.Get()), nullptr); auto* flow = FindSingleFlowNode(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>()); + 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>()); EXPECT_EQ(0_i, - flow->cases[0].selectors[0].val->value->As>()->ValueOf()); + cases[0].selectors[0].val->Value()->As>()->ValueOf()); - ASSERT_EQ(1u, flow->cases[1].selectors.Length()); - EXPECT_TRUE(flow->cases[1].selectors[0].IsDefault()); + ASSERT_EQ(1u, cases[1].selectors.Length()); + EXPECT_TRUE(cases[1].selectors[0].IsDefault()); - EXPECT_EQ(1u, flow->inbound_branches.Length()); - EXPECT_EQ(1u, flow->cases[0].start.target->inbound_branches.Length()); - EXPECT_EQ(1u, flow->cases[1].start.target->inbound_branches.Length()); - EXPECT_EQ(0u, flow->merge.target->inbound_branches.Length()); - EXPECT_EQ(2u, func->end_target->inbound_branches.Length()); + EXPECT_EQ(1u, flow->InboundBranches().Length()); + EXPECT_EQ(1u, cases[0].Start().target->InboundBranches().Length()); + EXPECT_EQ(1u, cases[1].Start().target->InboundBranches().Length()); + EXPECT_EQ(0u, flow->Merge().target->InboundBranches().Length()); + EXPECT_EQ(2u, func->EndTarget()->InboundBranches().Length()); EXPECT_EQ(Disassemble(m.Get()), R"(%fn1 = func test_function():void [@compute @workgroup_size(1, 1, 1)] { diff --git a/src/tint/ir/function.cc b/src/tint/ir/function.cc index 2a5f8c3801..d6cd6fc17c 100644 --- a/src/tint/ir/function.cc +++ b/src/tint/ir/function.cc @@ -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> wg_size) - : Base(), name(n), pipeline_stage(stage), workgroup_size(wg_size), return_type(rt) {} + : Base(), name_(name), return_type_(rt), pipeline_stage_(stage), workgroup_size_(wg_size) {} Function::~Function() = default; diff --git a/src/tint/ir/function.h b/src/tint/ir/function.h index 1fc651c117..86465aaf6c 100644 --- a/src/tint/ir/function.h +++ b/src/tint/ir/function.h @@ -17,6 +17,7 @@ #include #include +#include #include "src/tint/ir/flow_node.h" #include "src/tint/ir/function_param.h" @@ -71,32 +72,81 @@ class Function : public utils::Castable { type::Type* rt, PipelineStage stage = PipelineStage::kUndefined, std::optional> 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> 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 return_attributes; - /// If the return attribute is `kLocation` this stores the location value. - std::optional return_location; + /// @returns the function pipeline stage + PipelineStage Stage() const { return pipeline_stage_; } - /// The parameters to the function - utils::Vector 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> 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 attrs) { + return_attributes_ = std::move(attrs); + } + /// @returns the return attributes + utils::VectorRef ReturnAttributes() const { return return_attributes_; } + + /// Sets the return location + /// @param loc the location to set + void SetReturnLocation(std::optional loc) { return_location_ = loc; } + /// @returns the return location + std::optional ReturnLocation() const { return return_location_; } + + /// Sets the function parameters + /// @param params the function paramters + void SetParams(utils::VectorRef params) { params_ = std::move(params); } + + /// @returns the function parameters + utils::VectorRef 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> workgroup_size_; + + utils::Vector return_attributes_; + std::optional return_location_; + + utils::Vector params_; + + Block* start_target_ = nullptr; + FunctionTerminator* end_target_ = nullptr; }; utils::StringStream& operator<<(utils::StringStream& out, Function::PipelineStage value); diff --git a/src/tint/ir/function_param.cc b/src/tint/ir/function_param.cc index 8a05d6a797..a9570faa8d 100644 --- a/src/tint/ir/function_param.cc +++ b/src/tint/ir/function_param.cc @@ -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; diff --git a/src/tint/ir/function_param.h b/src/tint/ir/function_param.h index 46a3d5ee66..1bbb812943 100644 --- a/src/tint/ir/function_param.h +++ b/src/tint/ir/function_param.h @@ -34,10 +34,11 @@ class FunctionParam : public utils::Castable { 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 diff --git a/src/tint/ir/if.cc b/src/tint/ir/if.cc index b59d87f56d..d235b5efd3 100644 --- a/src/tint/ir/if.cc +++ b/src/tint/ir/if.cc @@ -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; diff --git a/src/tint/ir/if.h b/src/tint/ir/if.h index aadc5c9912..d02c32a425 100644 --- a/src/tint/ir/if.h +++ b/src/tint/ir/if.h @@ -32,17 +32,36 @@ class If : public utils::Castable { /// Constructor /// @param cond the if condition explicit If(Value* cond); + If(const If&) = delete; + If(If&&) = delete; ~If() override; - /// The true branch block + If& operator=(const If&) = delete; + If& operator=(If&&) = delete; + + /// @returns the if condition + const Value* Condition() const { return condition_; } + + /// @returns the true branch block + const Branch& True() const { return true_; } + /// @returns the true branch block + Branch& True() { return true_; } + + /// @returns the false branch block + const Branch& False() const { return false_; } + /// @returns the false branch block + Branch& False() { return false_; } + + /// @returns the merge branch block + const Branch& Merge() const { return merge_; } + /// @returns the merge branch block + Branch& Merge() { return merge_; } + + private: Branch true_ = {}; - /// The false branch block Branch false_ = {}; - /// An block to converge the true/false branches. The block always exists, but there maybe no - /// branches into it. (e.g. if both branches `return`) - Branch merge = {}; - /// Value holding the condition result - const Value* condition = nullptr; + Branch merge_ = {}; + Value* condition_; }; } // namespace tint::ir diff --git a/src/tint/ir/load.cc b/src/tint/ir/load.cc index 1fe55c01bb..3b64cdd752 100644 --- a/src/tint/ir/load.cc +++ b/src/tint/ir/load.cc @@ -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; diff --git a/src/tint/ir/load.h b/src/tint/ir/load.h index b15ecedc2e..e1a365fcdf 100644 --- a/src/tint/ir/load.h +++ b/src/tint/ir/load.h @@ -35,13 +35,14 @@ class Load : public utils::Castable { Load& operator=(Load&& inst) = delete; /// @returns the type of the value - const type::Type* Type() const override { return result_type; } + const type::Type* Type() const override { return result_type_; } - /// the result type of the instruction - const type::Type* result_type = nullptr; + /// @returns the avlue being loaded from + Value* From() const { return from_; } - /// the value being loaded - Value* from = nullptr; + private: + const type::Type* result_type_; + Value* from_; }; } // namespace tint::ir diff --git a/src/tint/ir/load_test.cc b/src/tint/ir/load_test.cc index 2c6e5c1402..9881e98a2c 100644 --- a/src/tint/ir/load_test.cc +++ b/src/tint/ir/load_test.cc @@ -33,12 +33,12 @@ TEST_F(IR_InstructionTest, CreateLoad) { const auto* inst = b.Load(var); ASSERT_TRUE(inst->Is()); - ASSERT_EQ(inst->from, var); + ASSERT_EQ(inst->From(), var); EXPECT_EQ(inst->Type(), store_type); - ASSERT_TRUE(inst->from->Is()); - EXPECT_EQ(inst->from, var); + ASSERT_TRUE(inst->From()->Is()); + 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 diff --git a/src/tint/ir/loop.h b/src/tint/ir/loop.h index e0066f4465..590f49c567 100644 --- a/src/tint/ir/loop.h +++ b/src/tint/ir/loop.h @@ -26,16 +26,32 @@ class Loop : public utils::Castable { public: /// Constructor Loop(); + Loop(const Loop&) = delete; + Loop(Loop&&) = delete; ~Loop() override; - /// The start block is the first block in a loop. - Branch start = {}; - /// The continue target of the block. - Branch continuing = {}; - /// The loop merge target. If the `loop` does a `return` then this block may not actually - /// end up in the control flow. We need it if the loop does a `break` we know where to break - /// too. - Branch merge = {}; + Loop& operator=(const Loop&) = delete; + Loop& operator=(Loop&&) = delete; + + /// @returns the switch start branch + const Branch& Start() const { return start_; } + /// @returns the switch start branch + Branch& Start() { return start_; } + + /// @returns the switch continuing branch + const Branch& Continuing() const { return continuing_; } + /// @returns the switch continuing branch + Branch& Continuing() { return continuing_; } + + /// @returns the switch merge branch + const Branch& Merge() const { return merge_; } + /// @returns the switch merge branch + Branch& Merge() { return merge_; } + + private: + Branch start_ = {}; + Branch continuing_ = {}; + Branch merge_ = {}; }; } // namespace tint::ir diff --git a/src/tint/ir/root_terminator.h b/src/tint/ir/root_terminator.h index 361aa6d464..c6338ba4b4 100644 --- a/src/tint/ir/root_terminator.h +++ b/src/tint/ir/root_terminator.h @@ -25,7 +25,12 @@ class RootTerminator : public utils::Castable { public: /// Constructor RootTerminator(); + RootTerminator(const RootTerminator&) = delete; + RootTerminator(RootTerminator&&) = delete; ~RootTerminator() override; + + RootTerminator& operator=(const RootTerminator&) = delete; + RootTerminator& operator=(RootTerminator&&) = delete; }; } // namespace tint::ir diff --git a/src/tint/ir/store.cc b/src/tint/ir/store.cc index 87f3620da5..8b8de4681b 100644 --- a/src/tint/ir/store.cc +++ b/src/tint/ir/store.cc @@ -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; diff --git a/src/tint/ir/store.h b/src/tint/ir/store.h index 9095c41224..af5377df55 100644 --- a/src/tint/ir/store.h +++ b/src/tint/ir/store.h @@ -34,10 +34,15 @@ class Store : public utils::Castable { Store& operator=(const Store& inst) = delete; Store& operator=(Store&& inst) = delete; - /// the value being stored to - Value* to = nullptr; - /// the value being stored - Value* from = nullptr; + /// @returns the value being stored too + Value* To() const { return to_; } + + /// @returns the value being stored + Value* From() const { return from_; } + + private: + Value* to_; + Value* from_; }; } // namespace tint::ir diff --git a/src/tint/ir/store_test.cc b/src/tint/ir/store_test.cc index 902ca955d9..1906cc6ff1 100644 --- a/src/tint/ir/store_test.cc +++ b/src/tint/ir/store_test.cc @@ -33,10 +33,10 @@ TEST_F(IR_InstructionTest, CreateStore) { const auto* inst = b.Store(to, b.Constant(4_i)); ASSERT_TRUE(inst->Is()); - ASSERT_EQ(inst->to, to); + ASSERT_EQ(inst->To(), to); - ASSERT_TRUE(inst->from->Is()); - auto lhs = inst->from->As()->value; + ASSERT_TRUE(inst->From()->Is()); + auto lhs = inst->From()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); } @@ -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 diff --git a/src/tint/ir/switch.cc b/src/tint/ir/switch.cc index ad6a1453f7..3bccc83677 100644 --- a/src/tint/ir/switch.cc +++ b/src/tint/ir/switch.cc @@ -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; diff --git a/src/tint/ir/switch.h b/src/tint/ir/switch.h index 6be7b62efc..2e977eced7 100644 --- a/src/tint/ir/switch.h +++ b/src/tint/ir/switch.h @@ -41,21 +41,40 @@ class Switch : public utils::Castable { utils::Vector 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 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 Cases() const { return cases_; } + /// @returns the switch cases + utils::Vector& Cases() { return cases_; } + + /// @returns the condition + const Value* Condition() const { return condition_; } + + private: + Branch merge_ = {}; + utils::Vector cases_; + Value* condition_; }; } // namespace tint::ir diff --git a/src/tint/ir/test_helper.h b/src/tint/ir/test_helper.h index d9055cee86..b7279f3147 100644 --- a/src/tint/ir/test_helper.h +++ b/src/tint/ir/test_helper.h @@ -33,7 +33,6 @@ template class TestHelperBase : public BASE, public ProgramBuilder { public: TestHelperBase() = default; - ~TestHelperBase() override = default; /// Build the module, cleaning up the program before returning. diff --git a/src/tint/ir/to_program.cc b/src/tint/ir/to_program.cc index b1aa67fd62..9be189c216 100644 --- a/src/tint/ir/to_program.cc +++ b/src/tint/ir/to_program.cc @@ -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 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(NextNonEmptyNode(i->false_.target)); + if (auto* else_if = As(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()) { - 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()) { - for (auto* inst : block->instructions) { + for (const auto* inst : block->Instructions()) { // Load instructions will be inlined, so ignore them. if (!inst->Is()) { 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()); }, - [&](const type::U32*) { return b.Expr(c->value->ValueAs()); }, - [&](const type::F32*) { return b.Expr(c->value->ValueAs()); }, - [&](const type::F16*) { return b.Expr(c->value->ValueAs()); }, - [&](const type::Bool*) { return b.Expr(c->value->ValueAs()); }, + [&](const type::I32*) { return b.Expr(c->Value()->ValueAs()); }, + [&](const type::U32*) { return b.Expr(c->Value()->ValueAs()); }, + [&](const type::F32*) { return b.Expr(c->Value()->ValueAs()); }, + [&](const type::F16*) { return b.Expr(c->Value()->ValueAs()); }, + [&](const type::Bool*) { return b.Expr(c->Value()->ValueAs()); }, [&](Default) { UNHANDLED_CASE(c); return nullptr; }); } - const ast::Expression* LoadExpr(const ir::Load* l) { return Expr(l->from); } + const ast::Expression* LoadExpr(const ir::Load* l) { return Expr(l->From()); } const ast::Expression* VarExpr(const ir::Var* v) { return b.Expr(NameOf(v)); } diff --git a/src/tint/ir/transform/add_empty_entry_point.cc b/src/tint/ir/transform/add_empty_entry_point.cc index 809a6ad773..6788d7c375 100644 --- a/src/tint/ir/transform/add_empty_entry_point.cc +++ b/src/tint/ir/transform/add_empty_entry_point.cc @@ -29,7 +29,7 @@ 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(), Function::PipelineStage::kCompute, std::array{1u, 1u, 1u}); - builder.Branch(ep->start_target, ep->end_target); + ep->StartTarget()->BranchTo(ep->EndTarget()); ir->functions.Push(ep); } diff --git a/src/tint/ir/transform/add_empty_entry_point_test.cc b/src/tint/ir/transform/add_empty_entry_point_test.cc index baba8e0901..a363d83640 100644 --- a/src/tint/ir/transform/add_empty_entry_point_test.cc +++ b/src/tint/ir/transform/add_empty_entry_point_test.cc @@ -40,7 +40,7 @@ TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) { TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) { auto* ep = b.CreateFunction(mod.symbols.New("main"), mod.types.Get(), Function::PipelineStage::kFragment); - b.Branch(ep->start_target, ep->end_target); + ep->StartTarget()->BranchTo(ep->EndTarget()); mod.functions.Push(ep); auto* expect = R"( diff --git a/src/tint/ir/unary.cc b/src/tint/ir/unary.cc index db13d5d579..2a0ee944e8 100644 --- a/src/tint/ir/unary.cc +++ b/src/tint/ir/unary.cc @@ -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); } diff --git a/src/tint/ir/unary.h b/src/tint/ir/unary.h index 46edd450e9..98ad9b8b2c 100644 --- a/src/tint/ir/unary.h +++ b/src/tint/ir/unary.h @@ -33,7 +33,7 @@ class Unary : public utils::Castable { /// @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& operator=(Unary&& inst) = delete; /// @returns the type of the value - const type::Type* Type() const override { return result_type; } + const type::Type* Type() const override { return result_type_; } /// @returns the value for the instruction const Value* Val() const { return val_; } - /// the kind of unary instruction - Kind kind = Kind::kNegation; - - /// the result type of the instruction - const type::Type* result_type = nullptr; + /// @returns the kind of unary instruction + enum Kind Kind() const { return kind_; } private: - Value* val_ = nullptr; + enum Kind kind_; + const type::Type* result_type_; + Value* val_; }; } // namespace tint::ir diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc index 280c7de831..6cefed38a1 100644 --- a/src/tint/ir/unary_test.cc +++ b/src/tint/ir/unary_test.cc @@ -29,10 +29,10 @@ TEST_F(IR_InstructionTest, CreateComplement) { const auto* inst = b.Complement(b.ir.types.Get(), b.Constant(4_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Unary::Kind::kComplement); + EXPECT_EQ(inst->Kind(), Unary::Kind::kComplement); ASSERT_TRUE(inst->Val()->Is()); - auto lhs = inst->Val()->As()->value; + auto lhs = inst->Val()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); } @@ -43,10 +43,10 @@ TEST_F(IR_InstructionTest, CreateNegation) { const auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); ASSERT_TRUE(inst->Is()); - EXPECT_EQ(inst->kind, Unary::Kind::kNegation); + EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation); ASSERT_TRUE(inst->Val()->Is()); - auto lhs = inst->Val()->As()->value; + auto lhs = inst->Val()->As()->Value(); ASSERT_TRUE(lhs->Is>()); EXPECT_EQ(4_i, lhs->As>()->ValueAs()); } @@ -56,7 +56,7 @@ TEST_F(IR_InstructionTest, Unary_Usage) { Builder b{mod}; const auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); - EXPECT_EQ(inst->kind, Unary::Kind::kNegation); + EXPECT_EQ(inst->Kind(), Unary::Kind::kNegation); ASSERT_NE(inst->Val(), nullptr); ASSERT_EQ(inst->Val()->Usage().Length(), 1u); diff --git a/src/tint/ir/user_call.cc b/src/tint/ir/user_call.cc index f718284ee5..e44e28574b 100644 --- a/src/tint/ir/user_call.cc +++ b/src/tint/ir/user_call.cc @@ -23,7 +23,7 @@ TINT_INSTANTIATE_TYPEINFO(tint::ir::UserCall); namespace tint::ir { UserCall::UserCall(const type::Type* ty, Symbol n, utils::VectorRef arguments) - : Base(ty, std::move(arguments)), name(n) {} + : Base(ty, std::move(arguments)), name_(n) {} UserCall::~UserCall() = default; diff --git a/src/tint/ir/user_call.h b/src/tint/ir/user_call.h index 5ada8f8ddf..ba52e20d9b 100644 --- a/src/tint/ir/user_call.h +++ b/src/tint/ir/user_call.h @@ -36,8 +36,11 @@ class UserCall : public utils::Castable { UserCall& operator=(const UserCall& inst) = delete; UserCall& operator=(UserCall&& inst) = delete; - /// The function name - Symbol name; + /// @returns the called function name + Symbol Name() const { return name_; } + + private: + Symbol name_; }; } // namespace tint::ir diff --git a/src/tint/ir/var.cc b/src/tint/ir/var.cc index 9bf4b8da12..9b43329cd5 100644 --- a/src/tint/ir/var.cc +++ b/src/tint/ir/var.cc @@ -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 diff --git a/src/tint/ir/var.h b/src/tint/ir/var.h index c874a6235f..67ccf0d3d8 100644 --- a/src/tint/ir/var.h +++ b/src/tint/ir/var.h @@ -36,13 +36,17 @@ class Var : public utils::Castable { Var& operator=(Var&& inst) = delete; /// @returns the type of the var - const type::Type* Type() const override { return type; } + const type::Type* Type() const override { return type_; } - /// the result type of the instruction - const type::Type* type = nullptr; + /// Sets the var initializer + /// @param initializer the initializer + void SetInitializer(Value* initializer); + /// @returns the initializer + const Value* Initializer() const { return initializer_; } - /// The optional initializer - Value* initializer = nullptr; + private: + const type::Type* type_; + Value* initializer_ = nullptr; }; } // namespace tint::ir diff --git a/src/tint/transform/manager_test.cc b/src/tint/transform/manager_test.cc index 3b0ee25f19..a81f7bf07c 100644 --- a/src/tint/transform/manager_test.cc +++ b/src/tint/transform/manager_test.cc @@ -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()); - 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()); - 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(); 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 diff --git a/src/tint/writer/spirv/ir/generator_impl_ir.cc b/src/tint/writer/spirv/ir/generator_impl_ir.cc index 059e5273f3..bd9a735bc5 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir.cc @@ -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(); - auto* true_block = i->true_.target->As(); - auto* false_block = i->false_.target->As(); + auto* merge_block = i->Merge().target->As(); + auto* true_block = i->True().target->As(); + auto* false_block = i->False().target->As(); // 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(binary->kind); + << "unimplemented binary instruction: " << static_cast(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_) diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc index b893496c0e..e9231f32b4 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_binary_test.cc @@ -21,10 +21,10 @@ namespace { TEST_F(SpvGeneratorImplTest, Binary_Add_I32) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); - func->start_target->instructions.Push( - b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(2_i))); + func->StartTarget()->SetInstructions( + utils::Vector{b.Add(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); - func->start_target->instructions.Push( - b.Add(mod.types.Get(), b.Constant(1_u), b.Constant(2_u))); + func->StartTarget()->SetInstructions( + utils::Vector{b.Add(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); - func->start_target->instructions.Push( - b.Add(mod.types.Get(), b.Constant(1_f), b.Constant(2_f))); + func->StartTarget()->SetInstructions( + utils::Vector{b.Add(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); - func->start_target->instructions.Push( - b.Subtract(mod.types.Get(), b.Constant(1_i), b.Constant(2_i))); + func->StartTarget()->SetInstructions( + utils::Vector{b.Subtract(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); - func->start_target->instructions.Push( - b.Subtract(mod.types.Get(), b.Constant(1_u), b.Constant(2_u))); + func->StartTarget()->SetInstructions( + utils::Vector{b.Subtract(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); - func->start_target->instructions.Push( - b.Subtract(mod.types.Get(), b.Constant(1_f), b.Constant(2_f))); + func->StartTarget()->SetInstructions( + utils::Vector{b.Subtract(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* lhs = mod.constants_arena.Create( mod.types.Get(mod.types.Get(), 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( mod.types.Get(mod.types.Get(), 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(mod.types.Get(), 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(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* lhs = mod.constants_arena.Create( mod.types.Get(mod.types.Get(), 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( mod.types.Get(mod.types.Get(), 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(mod.types.Get(), 4u), b.Constant(lhs), - b.Constant(rhs))); + func->StartTarget()->SetInstructions( + utils::Vector{b.Subtract(mod.types.Get(mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* a = b.Subtract(mod.types.Get(), b.Constant(1_i), b.Constant(2_i)); - func->start_target->instructions.Push(a); - func->start_target->instructions.Push(b.Add(mod.types.Get(), a, a)); + func->StartTarget()->SetInstructions(utils::Vector{a, b.Add(mod.types.Get(), a, a)}); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc index c58859a863..6ab48aa2a2 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_constant_test.cc @@ -67,7 +67,7 @@ TEST_F(SpvGeneratorImplTest, Constant_Vec4Bool) { auto* f = b.Constant(false); auto* v = mod.constants_arena.Create( mod.types.Get(mod.types.Get(), 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( - mod.types.Get(i, 2u), utils::Vector{i_42->value, i_n1->value}, false, false); + mod.types.Get(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( - mod.types.Get(u, 3u), utils::Vector{u_42->value, u_0->value, u_4b->value}, - false, true); + mod.types.Get(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( mod.types.Get(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( - mod.types.Get(h, 2u), utils::Vector{h_42->value, h_q->value}, false, false); + mod.types.Get(h, 2u), utils::Vector{h_42->Value(), h_q->Value()}, false, + false); generator_.Constant(b.Constant(v)); EXPECT_EQ(DumpTypes(), R"(%3 = OpTypeFloat 16 %2 = OpTypeVector %3 2 diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc index 77b4a62a57..d2af246a79 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_function_test.cc @@ -19,7 +19,7 @@ namespace { TEST_F(SpvGeneratorImplTest, Function_Empty) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - 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()); - 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(), 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(), 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(), 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(), 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(), 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(), ir::Function::PipelineStage::kFragment); - b.Branch(f3->start_target, f3->end_target); + f3->StartTarget()->BranchTo(f3->EndTarget()); generator_.EmitFunction(f1); generator_.EmitFunction(f2); diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc index 7b41184ac3..fa8cc6bc11 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_if_test.cc @@ -23,11 +23,11 @@ TEST_F(SpvGeneratorImplTest, If_TrueEmpty_FalseEmpty) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); auto* i = b.CreateIf(b.Constant(true)); - b.Branch(i->true_.target->As(), i->merge.target); - b.Branch(i->false_.target->As(), i->merge.target); - b.Branch(i->merge.target->As(), func->end_target); + i->True().target->As()->BranchTo(i->Merge().target); + i->False().target->As()->BranchTo(i->Merge().target); + i->Merge().target->As()->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()); auto* i = b.CreateIf(b.Constant(true)); - b.Branch(i->false_.target->As(), i->merge.target); - b.Branch(i->merge.target->As(), func->end_target); + i->False().target->As()->BranchTo(i->Merge().target); + i->Merge().target->As()->BranchTo(func->EndTarget()); - auto* true_block = i->true_.target->As(); - true_block->instructions.Push( - b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(1_i))); - b.Branch(true_block, i->merge.target); + auto* true_block = i->True().target->As(); + true_block->SetInstructions( + utils::Vector{b.Add(mod.types.Get(), 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()); auto* i = b.CreateIf(b.Constant(true)); - b.Branch(i->true_.target->As(), i->merge.target); - b.Branch(i->merge.target->As(), func->end_target); + i->True().target->As()->BranchTo(i->Merge().target); + i->Merge().target->As()->BranchTo(func->EndTarget()); - auto* false_block = i->false_.target->As(); - false_block->instructions.Push( - b.Add(mod.types.Get(), b.Constant(1_i), b.Constant(1_i))); - b.Branch(false_block, i->merge.target); + auto* false_block = i->False().target->As(); + false_block->SetInstructions( + utils::Vector{b.Add(mod.types.Get(), 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()); auto* i = b.CreateIf(b.Constant(true)); - b.Branch(i->true_.target->As(), func->end_target); - b.Branch(i->false_.target->As(), func->end_target); - i->merge.target->As()->branch.target = nullptr; + i->True().target->As()->BranchTo(func->EndTarget()); + i->False().target->As()->BranchTo(func->EndTarget()); + i->Merge().target->As()->BranchTo(nullptr); - b.Branch(func->start_target, i); + func->StartTarget()->BranchTo(i); generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" diff --git a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc index d9bc60c67f..20e5167980 100644 --- a/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc +++ b/src/tint/writer/spirv/ir/generator_impl_ir_var_test.cc @@ -22,12 +22,11 @@ namespace { TEST_F(SpvGeneratorImplTest, FunctionVar_NoInit) { auto* func = b.CreateFunction(mod.symbols.Register("foo"), mod.types.Get()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), 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(), func->end_target); - b.Branch(i->merge.target->As(), func->end_target); + i->False().target->As()->BranchTo(func->EndTarget()); + i->Merge().target->As()->BranchTo(func->EndTarget()); - auto* true_block = i->true_.target->As(); - true_block->instructions.Push(v); - b.Branch(true_block, i->merge.target); + auto* true_block = i->True().target->As(); + 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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* store_ty = mod.types.Get(); auto* ty = mod.types.Get(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()); - b.Branch(func->start_target, func->end_target); + func->StartTarget()->BranchTo(func->EndTarget()); auto* ty = mod.types.Get( mod.types.Get(), 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"