diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 5d04077773..3073d81a9d 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -1121,8 +1121,6 @@ libtint_source_set("libtint_syntax_tree_writer_src") { libtint_source_set("libtint_ir_builder_src") { sources = [ - "ir/builder_impl.cc", - "ir/builder_impl.h", "ir/from_program.cc", "ir/from_program.h", ] @@ -2151,16 +2149,16 @@ if (tint_build_unittests) { sources = [ "ir/binary_test.cc", "ir/bitcast_test.cc", - "ir/builder_impl_binary_test.cc", - "ir/builder_impl_call_test.cc", - "ir/builder_impl_literal_test.cc", - "ir/builder_impl_materialize_test.cc", - "ir/builder_impl_store_test.cc", - "ir/builder_impl_test.cc", - "ir/builder_impl_unary_test.cc", - "ir/builder_impl_var_test.cc", "ir/constant_test.cc", "ir/discard_test.cc", + "ir/from_program_binary_test.cc", + "ir/from_program_call_test.cc", + "ir/from_program_literal_test.cc", + "ir/from_program_materialize_test.cc", + "ir/from_program_store_test.cc", + "ir/from_program_test.cc", + "ir/from_program_unary_test.cc", + "ir/from_program_var_test.cc", "ir/module_test.cc", "ir/store_test.cc", "ir/test_helper.h", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index ea69d503b0..7feda1dba1 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -716,8 +716,6 @@ if(${TINT_BUILD_IR}) ir/block.h ir/builder.cc ir/builder.h - ir/builder_impl.cc - ir/builder_impl.h ir/builtin.cc ir/builtin.h ir/call.cc @@ -1447,16 +1445,16 @@ if(TINT_BUILD_TESTS) list(APPEND TINT_TEST_SRCS ir/binary_test.cc ir/bitcast_test.cc - ir/builder_impl_binary_test.cc - ir/builder_impl_call_test.cc - ir/builder_impl_literal_test.cc - ir/builder_impl_materialize_test.cc - ir/builder_impl_store_test.cc - ir/builder_impl_test.cc - ir/builder_impl_unary_test.cc - ir/builder_impl_var_test.cc ir/constant_test.cc ir/discard_test.cc + ir/from_program_binary_test.cc + ir/from_program_call_test.cc + ir/from_program_literal_test.cc + ir/from_program_materialize_test.cc + ir/from_program_store_test.cc + ir/from_program_test.cc + ir/from_program_unary_test.cc + ir/from_program_var_test.cc ir/module_test.cc ir/store_test.cc ir/test_helper.h diff --git a/src/tint/ir/binary_test.cc b/src/tint/ir/binary_test.cc index 6705f63e1e..2968d283a3 100644 --- a/src/tint/ir/binary_test.cc +++ b/src/tint/ir/binary_test.cc @@ -24,7 +24,8 @@ using namespace tint::number_suffixes; // NOLINT using IR_InstructionTest = TestHelper; TEST_F(IR_InstructionTest, CreateAnd) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.And(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -45,7 +46,8 @@ TEST_F(IR_InstructionTest, CreateAnd) { } TEST_F(IR_InstructionTest, CreateOr) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Or(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -64,7 +66,8 @@ TEST_F(IR_InstructionTest, CreateOr) { } TEST_F(IR_InstructionTest, CreateXor) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Xor(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -83,7 +86,8 @@ TEST_F(IR_InstructionTest, CreateXor) { } TEST_F(IR_InstructionTest, CreateEqual) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Equal(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -102,7 +106,8 @@ TEST_F(IR_InstructionTest, CreateEqual) { } TEST_F(IR_InstructionTest, CreateNotEqual) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.NotEqual(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -121,7 +126,8 @@ TEST_F(IR_InstructionTest, CreateNotEqual) { } TEST_F(IR_InstructionTest, CreateLessThan) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.LessThan(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -140,7 +146,8 @@ TEST_F(IR_InstructionTest, CreateLessThan) { } TEST_F(IR_InstructionTest, CreateGreaterThan) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.GreaterThan(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -160,7 +167,8 @@ TEST_F(IR_InstructionTest, CreateGreaterThan) { } TEST_F(IR_InstructionTest, CreateLessThanEqual) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.LessThanEqual(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -180,7 +188,8 @@ TEST_F(IR_InstructionTest, CreateLessThanEqual) { } TEST_F(IR_InstructionTest, CreateGreaterThanEqual) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.GreaterThanEqual(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -200,7 +209,8 @@ TEST_F(IR_InstructionTest, CreateGreaterThanEqual) { } TEST_F(IR_InstructionTest, CreateNot) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Not(b.ir.types.Get(), b.Constant(true)); ASSERT_TRUE(inst->Is()); @@ -218,7 +228,8 @@ TEST_F(IR_InstructionTest, CreateNot) { } TEST_F(IR_InstructionTest, CreateShiftLeft) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.ShiftLeft(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -237,7 +248,8 @@ TEST_F(IR_InstructionTest, CreateShiftLeft) { } TEST_F(IR_InstructionTest, CreateShiftRight) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.ShiftRight(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -256,7 +268,8 @@ TEST_F(IR_InstructionTest, CreateShiftRight) { } TEST_F(IR_InstructionTest, CreateAdd) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Add(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -275,7 +288,8 @@ TEST_F(IR_InstructionTest, CreateAdd) { } TEST_F(IR_InstructionTest, CreateSubtract) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Subtract(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -294,7 +308,8 @@ TEST_F(IR_InstructionTest, CreateSubtract) { } TEST_F(IR_InstructionTest, CreateMultiply) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Multiply(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -313,7 +328,8 @@ TEST_F(IR_InstructionTest, CreateMultiply) { } TEST_F(IR_InstructionTest, CreateDivide) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Divide(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -332,7 +348,8 @@ TEST_F(IR_InstructionTest, CreateDivide) { } TEST_F(IR_InstructionTest, CreateModulo) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Modulo(b.ir.types.Get(), b.Constant(4_i), b.Constant(2_i)); @@ -351,7 +368,8 @@ TEST_F(IR_InstructionTest, CreateModulo) { } TEST_F(IR_InstructionTest, Binary_Usage) { - Builder b; + Module mod; + 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); @@ -366,7 +384,8 @@ TEST_F(IR_InstructionTest, Binary_Usage) { } TEST_F(IR_InstructionTest, Binary_Usage_DuplicateValue) { - Builder b; + Module mod; + Builder b{mod}; auto val = b.Constant(4_i); const auto* inst = b.And(b.ir.types.Get(), val, val); diff --git a/src/tint/ir/bitcast_test.cc b/src/tint/ir/bitcast_test.cc index 0062deb60c..e78b734795 100644 --- a/src/tint/ir/bitcast_test.cc +++ b/src/tint/ir/bitcast_test.cc @@ -25,7 +25,8 @@ using namespace tint::number_suffixes; // NOLINT using IR_InstructionTest = TestHelper; TEST_F(IR_InstructionTest, Bitcast) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Bitcast(b.ir.types.Get(), b.Constant(4_i)); ASSERT_TRUE(inst->Is()); @@ -39,7 +40,8 @@ TEST_F(IR_InstructionTest, Bitcast) { } TEST_F(IR_InstructionTest, Bitcast_Usage) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Bitcast(b.ir.types.Get(), b.Constant(4_i)); ASSERT_EQ(inst->args.Length(), 1u); diff --git a/src/tint/ir/builder_impl.cc b/src/tint/ir/builder_impl.cc deleted file mode 100644 index 57cd02c465..0000000000 --- a/src/tint/ir/builder_impl.cc +++ /dev/null @@ -1,1139 +0,0 @@ -// Copyright 2022 The Tint Authors. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#include "src/tint/ir/builder_impl.h" - -#include - -#include "src/tint/ast/alias.h" -#include "src/tint/ast/assignment_statement.h" -#include "src/tint/ast/binary_expression.h" -#include "src/tint/ast/bitcast_expression.h" -#include "src/tint/ast/block_statement.h" -#include "src/tint/ast/bool_literal_expression.h" -#include "src/tint/ast/break_if_statement.h" -#include "src/tint/ast/break_statement.h" -#include "src/tint/ast/call_expression.h" -#include "src/tint/ast/call_statement.h" -#include "src/tint/ast/compound_assignment_statement.h" -#include "src/tint/ast/const.h" -#include "src/tint/ast/const_assert.h" -#include "src/tint/ast/continue_statement.h" -#include "src/tint/ast/discard_statement.h" -#include "src/tint/ast/enable.h" -#include "src/tint/ast/float_literal_expression.h" -#include "src/tint/ast/for_loop_statement.h" -#include "src/tint/ast/function.h" -#include "src/tint/ast/id_attribute.h" -#include "src/tint/ast/identifier.h" -#include "src/tint/ast/identifier_expression.h" -#include "src/tint/ast/if_statement.h" -#include "src/tint/ast/int_literal_expression.h" -#include "src/tint/ast/invariant_attribute.h" -#include "src/tint/ast/let.h" -#include "src/tint/ast/literal_expression.h" -#include "src/tint/ast/loop_statement.h" -#include "src/tint/ast/override.h" -#include "src/tint/ast/return_statement.h" -#include "src/tint/ast/statement.h" -#include "src/tint/ast/struct.h" -#include "src/tint/ast/struct_member_align_attribute.h" -#include "src/tint/ast/struct_member_size_attribute.h" -#include "src/tint/ast/switch_statement.h" -#include "src/tint/ast/templated_identifier.h" -#include "src/tint/ast/unary_op_expression.h" -#include "src/tint/ast/var.h" -#include "src/tint/ast/variable_decl_statement.h" -#include "src/tint/ast/while_statement.h" -#include "src/tint/ir/function.h" -#include "src/tint/ir/if.h" -#include "src/tint/ir/loop.h" -#include "src/tint/ir/module.h" -#include "src/tint/ir/store.h" -#include "src/tint/ir/switch.h" -#include "src/tint/ir/value.h" -#include "src/tint/program.h" -#include "src/tint/sem/builtin.h" -#include "src/tint/sem/call.h" -#include "src/tint/sem/function.h" -#include "src/tint/sem/materialize.h" -#include "src/tint/sem/module.h" -#include "src/tint/sem/switch_statement.h" -#include "src/tint/sem/value_constructor.h" -#include "src/tint/sem/value_conversion.h" -#include "src/tint/sem/value_expression.h" -#include "src/tint/sem/variable.h" -#include "src/tint/switch.h" -#include "src/tint/type/void.h" -#include "src/tint/utils/defer.h" -#include "src/tint/utils/scoped_assignment.h" - -namespace tint::ir { -namespace { - -using ResultType = utils::Result; - -class FlowStackScope { - public: - FlowStackScope(BuilderImpl* impl, FlowNode* node) : impl_(impl) { - impl_->flow_stack.Push(node); - } - - ~FlowStackScope() { impl_->flow_stack.Pop(); } - - private: - BuilderImpl* impl_; -}; - -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) { - if (IsConnected(parent)) { - return true; - } - } - // Getting here means all the incoming branches are disconnected. - return false; -} - -} // namespace - -BuilderImpl::BuilderImpl(const Program* program) - : program_(program), - clone_ctx_{ - type::CloneContext{{&program->Symbols()}, {&builder.ir.symbols, &builder.ir.types}}, - {&builder.ir.constants}} {} - -BuilderImpl::~BuilderImpl() = default; - -void BuilderImpl::add_error(const Source& s, const std::string& err) { - diagnostics_.add_error(tint::diag::System::IR, err, s); -} - -void BuilderImpl::BranchTo(FlowNode* node, utils::VectorRef args) { - TINT_ASSERT(IR, current_flow_block); - TINT_ASSERT(IR, !IsBranched(current_flow_block)); - - builder.Branch(current_flow_block, node, args); - current_flow_block = nullptr; -} - -void BuilderImpl::BranchToIfNeeded(FlowNode* node) { - if (!current_flow_block || IsBranched(current_flow_block)) { - return; - } - BranchTo(node); -} - -FlowNode* BuilderImpl::FindEnclosingControl(ControlFlags flags) { - for (auto it = flow_stack.rbegin(); it != flow_stack.rend(); ++it) { - if ((*it)->Is()) { - return *it; - } - if (flags == ControlFlags::kExcludeSwitch) { - continue; - } - if ((*it)->Is()) { - return *it; - } - } - return nullptr; -} - -Symbol BuilderImpl::CloneSymbol(Symbol sym) const { - return clone_ctx_.type_ctx.dst.st->Register(sym.Name()); -} - -ResultType BuilderImpl::Build() { - auto* sem = program_->Sem().Module(); - - for (auto* decl : sem->DependencyOrderedDeclarations()) { - tint::Switch( - decl, // - [&](const ast::Struct*) { - // Will be encoded into the `type::Struct` when used. We will then hoist all - // used structs up to module scope when converting IR. - }, - [&](const ast::Alias*) { - // Folded away and doesn't appear in the IR. - }, - [&](const ast::Variable* var) { - // Setup the current flow node to be the root block for the module. The builder will - // handle creating it if it doesn't exist already. - TINT_SCOPED_ASSIGNMENT(current_flow_block, builder.CreateRootBlockIfNeeded()); - EmitVariable(var); - }, - [&](const ast::Function* func) { EmitFunction(func); }, - [&](const ast::Enable*) { - // TODO(dsinclair): Implement? I think these need to be passed along so further - // stages know what is enabled. - }, - [&](const ast::ConstAssert*) { - // Evaluated by the resolver, drop from the IR. - }, - [&](Default) { - add_error(decl->source, "unknown type: " + std::string(decl->TypeInfo().name)); - }); - } - if (!diagnostics_.empty()) { - return utils::Failure; - } - - return ResultType{std::move(mod)}; -} - -void BuilderImpl::EmitFunction(const ast::Function* ast_func) { - // The flow stack should have been emptied when the previous function finished building. - TINT_ASSERT(IR, flow_stack.IsEmpty()); - - const auto* sem = program_->Sem().Get(ast_func); - - auto* ir_func = builder.CreateFunction(CloneSymbol(ast_func->name->symbol), - sem->ReturnType()->Clone(clone_ctx_.type_ctx)); - - current_function_ = ir_func; - builder.ir.functions.Push(ir_func); - - ast_to_flow_[ast_func] = ir_func; - - if (ast_func->IsEntryPoint()) { - builder.ir.entry_points.Push(ir_func); - - switch (ast_func->PipelineStage()) { - case ast::PipelineStage::kVertex: - ir_func->pipeline_stage = Function::PipelineStage::kVertex; - break; - case ast::PipelineStage::kFragment: - ir_func->pipeline_stage = Function::PipelineStage::kFragment; - break; - case ast::PipelineStage::kCompute: { - ir_func->pipeline_stage = 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), - }; - break; - } - default: { - TINT_ICE(IR, diagnostics_) << "Invalid pipeline stage"; - return; - } - } - - for (auto* attr : ast_func->return_type_attributes) { - tint::Switch( - attr, // - [&](const ast::LocationAttribute*) { - ir_func->return_attributes.Push(Function::ReturnAttribute::kLocation); - }, - [&](const ast::InvariantAttribute*) { - ir_func->return_attributes.Push(Function::ReturnAttribute::kInvariant); - }, - [&](const ast::BuiltinAttribute* b) { - if (auto* ident_sem = - program_->Sem() - .Get(b) - ->As>()) { - switch (ident_sem->Value()) { - case builtin::BuiltinValue::kPosition: - ir_func->return_attributes.Push( - Function::ReturnAttribute::kPosition); - break; - case builtin::BuiltinValue::kFragDepth: - ir_func->return_attributes.Push( - Function::ReturnAttribute::kFragDepth); - break; - case builtin::BuiltinValue::kSampleMask: - ir_func->return_attributes.Push( - Function::ReturnAttribute::kSampleMask); - break; - default: - TINT_ICE(IR, diagnostics_) - << "Unknown builtin value in return attributes " - << ident_sem->Value(); - return; - } - } else { - TINT_ICE(IR, diagnostics_) << "Builtin attribute sem invalid"; - return; - } - }); - } - } - ir_func->return_location = sem->ReturnLocation(); - - { - FlowStackScope scope(this, ir_func); - - current_flow_block = ir_func->start_target; - EmitBlock(ast_func->body); - - // TODO(dsinclair): Store return type and attributes - // TODO(dsinclair): Store parameters - - // If the branch target has already been set then a `return` was called. Only set in the - // case where `return` wasn't called. - BranchToIfNeeded(current_function_->end_target); - } - - TINT_ASSERT(IR, flow_stack.IsEmpty()); - current_flow_block = nullptr; - current_function_ = nullptr; -} - -void BuilderImpl::EmitStatements(utils::VectorRef stmts) { - for (auto* s : stmts) { - EmitStatement(s); - - // 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)) { - break; - } - } -} - -void BuilderImpl::EmitStatement(const ast::Statement* stmt) { - tint::Switch( - stmt, // - [&](const ast::AssignmentStatement* a) { EmitAssignment(a); }, - [&](const ast::BlockStatement* b) { EmitBlock(b); }, - [&](const ast::BreakStatement* b) { EmitBreak(b); }, - [&](const ast::BreakIfStatement* b) { EmitBreakIf(b); }, - [&](const ast::CallStatement* c) { EmitCall(c); }, - [&](const ast::CompoundAssignmentStatement* c) { EmitCompoundAssignment(c); }, - [&](const ast::ContinueStatement* c) { EmitContinue(c); }, - [&](const ast::DiscardStatement* d) { EmitDiscard(d); }, - [&](const ast::IfStatement* i) { EmitIf(i); }, - [&](const ast::LoopStatement* l) { EmitLoop(l); }, - [&](const ast::ForLoopStatement* l) { EmitForLoop(l); }, - [&](const ast::WhileStatement* l) { EmitWhile(l); }, - [&](const ast::ReturnStatement* r) { EmitReturn(r); }, - [&](const ast::SwitchStatement* s) { EmitSwitch(s); }, - [&](const ast::VariableDeclStatement* v) { EmitVariable(v->variable); }, - [&](const ast::ConstAssert*) { - // Not emitted - }, - [&](Default) { - add_error(stmt->source, - "unknown statement type: " + std::string(stmt->TypeInfo().name)); - }); -} - -void BuilderImpl::EmitAssignment(const ast::AssignmentStatement* stmt) { - auto lhs = EmitExpression(stmt->lhs); - if (!lhs) { - return; - } - - auto rhs = EmitExpression(stmt->rhs); - if (!rhs) { - return; - } - auto store = builder.Store(lhs.Get(), rhs.Get()); - current_flow_block->instructions.Push(store); -} - -void BuilderImpl::EmitCompoundAssignment(const ast::CompoundAssignmentStatement* stmt) { - auto lhs = EmitExpression(stmt->lhs); - if (!lhs) { - return; - } - - auto rhs = EmitExpression(stmt->rhs); - if (!rhs) { - return; - } - - auto* ty = lhs.Get()->Type(); - Binary* inst = nullptr; - switch (stmt->op) { - case ast::BinaryOp::kAnd: - inst = builder.And(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kOr: - inst = builder.Or(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kXor: - inst = builder.Xor(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kShiftLeft: - inst = builder.ShiftLeft(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kShiftRight: - inst = builder.ShiftRight(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kAdd: - inst = builder.Add(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kSubtract: - inst = builder.Subtract(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kMultiply: - inst = builder.Multiply(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kDivide: - inst = builder.Divide(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kModulo: - inst = builder.Modulo(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kLessThanEqual: - case ast::BinaryOp::kGreaterThanEqual: - case ast::BinaryOp::kGreaterThan: - case ast::BinaryOp::kLessThan: - case ast::BinaryOp::kNotEqual: - case ast::BinaryOp::kEqual: - case ast::BinaryOp::kLogicalAnd: - case ast::BinaryOp::kLogicalOr: - TINT_ICE(IR, diagnostics_) << "invalid compound assignment"; - return; - case ast::BinaryOp::kNone: - TINT_ICE(IR, diagnostics_) << "missing binary operand type"; - return; - } - current_flow_block->instructions.Push(inst); - - auto store = builder.Store(lhs.Get(), inst); - current_flow_block->instructions.Push(store); -} - -void BuilderImpl::EmitBlock(const ast::BlockStatement* block) { - scopes_.Push(); - TINT_DEFER(scopes_.Pop()); - - // Note, this doesn't need to emit a Block as the current block flow node should be sufficient - // as the blocks all get flattened. Each flow control node will inject the basic blocks it - // requires. - EmitStatements(block->statements); -} - -void BuilderImpl::EmitIf(const ast::IfStatement* stmt) { - // Emit the if condition into the end of the preceding block - auto reg = EmitExpression(stmt->condition); - if (!reg) { - return; - } - auto* if_node = builder.CreateIf(reg.Get()); - - BranchTo(if_node); - - ast_to_flow_[stmt] = if_node; - - { - FlowStackScope scope(this, if_node); - - 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); - - 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); - } - 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(); - } -} - -void BuilderImpl::EmitLoop(const ast::LoopStatement* stmt) { - auto* loop_node = builder.CreateLoop(); - - BranchTo(loop_node); - - ast_to_flow_[stmt] = loop_node; - - { - FlowStackScope scope(this, loop_node); - - current_flow_block = loop_node->start.target->As(); - EmitBlock(stmt->body); - - // The current block didn't `break`, `return` or `continue`, go to the continuing block. - BranchToIfNeeded(loop_node->continuing.target); - - 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); - } - - // 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)) { - current_flow_block = nullptr; - } -} - -void BuilderImpl::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); - - BranchTo(loop_node); - - ast_to_flow_[stmt] = loop_node; - - { - FlowStackScope scope(this, loop_node); - - current_flow_block = loop_node->start.target->As(); - - // Emit the while condition into the start target of the loop - auto reg = EmitExpression(stmt->condition); - if (!reg) { - return; - } - - // 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); - - BranchTo(if_node); - - current_flow_block = if_node->merge.target->As(); - EmitBlock(stmt->body); - - 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(); -} - -void BuilderImpl::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); - - // Make sure the initializer ends up in a contained scope - scopes_.Push(); - TINT_DEFER(scopes_.Pop()); - - if (stmt->initializer) { - // Emit the for initializer before branching to the loop - EmitStatement(stmt->initializer); - } - - BranchTo(loop_node); - - ast_to_flow_[stmt] = loop_node; - - { - FlowStackScope scope(this, loop_node); - - current_flow_block = loop_node->start.target->As(); - - if (stmt->condition) { - // Emit the condition into the target target of the loop - auto reg = EmitExpression(stmt->condition); - if (!reg) { - return; - } - - // 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); - - BranchTo(if_node); - current_flow_block = if_node->merge.target->As(); - } - - EmitBlock(stmt->body); - BranchToIfNeeded(loop_node->continuing.target); - - if (stmt->continuing) { - 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(); -} - -void BuilderImpl::EmitSwitch(const ast::SwitchStatement* stmt) { - // Emit the condition into the preceding block - auto reg = EmitExpression(stmt->condition); - if (!reg) { - return; - } - auto* switch_node = builder.CreateSwitch(reg.Get()); - - BranchTo(switch_node); - - ast_to_flow_[stmt] = switch_node; - - { - FlowStackScope scope(this, switch_node); - - const auto* sem = program_->Sem().Get(stmt); - for (const auto* c : sem->Cases()) { - utils::Vector selectors; - for (const auto* selector : c->Selectors()) { - if (selector->IsDefault()) { - selectors.Push({nullptr}); - } else { - selectors.Push({builder.Constant(selector->Value()->Clone(clone_ctx_))}); - } - } - - current_flow_block = builder.CreateCase(switch_node, selectors); - EmitBlock(c->Body()->Declaration()); - - BranchToIfNeeded(switch_node->merge.target); - } - } - current_flow_block = nullptr; - - if (IsConnected(switch_node->merge.target)) { - current_flow_block = switch_node->merge.target->As(); - } -} - -void BuilderImpl::EmitReturn(const ast::ReturnStatement* stmt) { - utils::Vector ret_value; - if (stmt->value) { - auto ret = EmitExpression(stmt->value); - if (!ret) { - return; - } - ret_value.Push(ret.Get()); - } - - BranchTo(current_function_->end_target, std::move(ret_value)); -} - -void BuilderImpl::EmitBreak(const ast::BreakStatement*) { - auto* current_control = FindEnclosingControl(ControlFlags::kNone); - TINT_ASSERT(IR, current_control); - - if (auto* c = current_control->As()) { - BranchTo(c->merge.target); - } else if (auto* s = current_control->As()) { - BranchTo(s->merge.target); - } else { - TINT_UNREACHABLE(IR, diagnostics_); - } -} - -void BuilderImpl::EmitContinue(const ast::ContinueStatement*) { - auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch); - TINT_ASSERT(IR, current_control); - - if (auto* c = current_control->As()) { - BranchTo(c->continuing.target); - } else { - TINT_UNREACHABLE(IR, diagnostics_); - } -} - -// Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so the -// code has to continue as before it just predicates writes. If WGSL grows some kind of terminating -// discard that would probably make sense as a FlowNode but would then require figuring out the -// multi-level exit that is triggered. -void BuilderImpl::EmitDiscard(const ast::DiscardStatement*) { - auto* inst = builder.Discard(); - current_flow_block->instructions.Push(inst); -} - -void BuilderImpl::EmitBreakIf(const ast::BreakIfStatement* stmt) { - // Emit the break-if condition into the end of the preceding block - auto reg = EmitExpression(stmt->condition); - if (!reg) { - return; - } - auto* if_node = builder.CreateIf(reg.Get()); - - BranchTo(if_node); - - ast_to_flow_[stmt] = if_node; - - auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch); - TINT_ASSERT(IR, current_control); - TINT_ASSERT(IR, current_control->Is()); - - auto* loop = current_control->As(); - - 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->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); -} - -utils::Result BuilderImpl::EmitExpression(const ast::Expression* expr) { - // If this is a value that has been const-eval'd return the result. - if (auto* sem = program_->Sem().Get(expr)->As()) { - if (auto* v = sem->ConstantValue()) { - if (auto* cv = v->Clone(clone_ctx_)) { - return builder.Constant(cv); - } - } - } - - return tint::Switch( - expr, - // [&](const ast::IndexAccessorExpression* a) { - // TODO(dsinclair): Implement - // }, - [&](const ast::BinaryExpression* b) { return EmitBinary(b); }, - [&](const ast::BitcastExpression* b) { return EmitBitcast(b); }, - [&](const ast::CallExpression* c) { return EmitCall(c); }, - [&](const ast::IdentifierExpression* i) { - auto* v = scopes_.Get(i->identifier->symbol); - return utils::Result{v}; - }, - [&](const ast::LiteralExpression* l) { return EmitLiteral(l); }, - // [&](const ast::MemberAccessorExpression* m) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::PhonyExpression*) { - // TODO(dsinclair): Implement. The call may have side effects so has to be made. - // }, - [&](const ast::UnaryOpExpression* u) { return EmitUnary(u); }, - [&](Default) { - add_error(expr->source, - "unknown expression type: " + std::string(expr->TypeInfo().name)); - return utils::Failure; - }); -} - -void BuilderImpl::EmitVariable(const ast::Variable* var) { - auto* sem = program_->Sem().Get(var); - - return tint::Switch( // - var, - [&](const ast::Var* v) { - auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); - auto* val = builder.Declare(ty, sem->AddressSpace(), sem->Access()); - current_flow_block->instructions.Push(val); - - if (v->initializer) { - auto init = EmitExpression(v->initializer); - if (!init) { - return; - } - val->initializer = init.Get(); - } - // Store the declaration so we can get the instruction to store too - scopes_.Set(v->name->symbol, val); - - // Record the original name of the var - builder.ir.SetName(val, v->name->symbol.Name()); - }, - [&](const ast::Let* l) { - // A `let` doesn't exist as a standalone item in the IR, it's just the result of the - // initializer. - auto init = EmitExpression(l->initializer); - if (!init) { - return; - } - - // Store the results of the initialization - scopes_.Set(l->name->symbol, init.Get()); - - // Record the original name of the let - builder.ir.SetName(init.Get(), l->name->symbol.Name()); - }, - [&](const ast::Override*) { - add_error(var->source, - "found an `Override` variable. The SubstituteOverrides " - "transform must be run before converting to IR"); - }, - [&](const ast::Const*) { - // Skip. This should be handled by const-eval already, so the const will be a - // `constant::` value at the usage sites. Can just ignore the `const` variable as it - // should never be used. - // - // TODO(dsinclair): Probably want to store the const variable somewhere and then in - // identifier expression log an error if we ever see a const identifier. Add this when - // identifiers and variables are supported. - }, - [&](Default) { - add_error(var->source, "unknown variable: " + std::string(var->TypeInfo().name)); - }); -} - -utils::Result BuilderImpl::EmitUnary(const ast::UnaryOpExpression* expr) { - auto val = EmitExpression(expr->expr); - if (!val) { - return utils::Failure; - } - - auto* sem = program_->Sem().Get(expr); - auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); - - Instruction* inst = nullptr; - switch (expr->op) { - case ast::UnaryOp::kAddressOf: - inst = builder.AddressOf(ty, val.Get()); - break; - case ast::UnaryOp::kComplement: - inst = builder.Complement(ty, val.Get()); - break; - case ast::UnaryOp::kIndirection: - inst = builder.Indirection(ty, val.Get()); - break; - case ast::UnaryOp::kNegation: - inst = builder.Negation(ty, val.Get()); - break; - case ast::UnaryOp::kNot: - inst = builder.Not(ty, val.Get()); - break; - } - - current_flow_block->instructions.Push(inst); - return inst; -} - -// A short-circut needs special treatment. The short-circuit is decomposed into the relevant if -// statements and declarations. -utils::Result BuilderImpl::EmitShortCircuit(const ast::BinaryExpression* expr) { - switch (expr->op) { - case ast::BinaryOp::kLogicalAnd: - case ast::BinaryOp::kLogicalOr: - break; - default: - TINT_ICE(IR, diagnostics_) << "invalid operation type for short-circut decomposition"; - return utils::Failure; - } - - // Evaluate the LHS of the short-circuit - auto lhs = EmitExpression(expr->lhs); - if (!lhs) { - return utils::Failure; - } - - // Generate a variable to store the short-circut into - auto* ty = builder.ir.types.Get(); - auto* result_var = - builder.Declare(ty, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); - current_flow_block->instructions.Push(result_var); - - auto* lhs_store = builder.Store(result_var, lhs.Get()); - current_flow_block->instructions.Push(lhs_store); - - auto* if_node = builder.CreateIf(lhs.Get()); - BranchTo(if_node); - - utils::Result rhs; - { - FlowStackScope scope(this, if_node); - - // If this is an `&&` then we only evaluate the RHS expression in the true block. - // If this is an `||` then we only evaluate the RHS expression in the false block. - if (expr->op == ast::BinaryOp::kLogicalAnd) { - current_flow_block = if_node->true_.target->As(); - } else { - current_flow_block = if_node->false_.target->As(); - } - - rhs = EmitExpression(expr->rhs); - if (!rhs) { - return utils::Failure; - } - auto* rhs_store = builder.Store(result_var, rhs.Get()); - current_flow_block->instructions.Push(rhs_store); - - BranchTo(if_node->merge.target); - } - current_flow_block = if_node->merge.target->As(); - - return result_var; -} - -utils::Result BuilderImpl::EmitBinary(const ast::BinaryExpression* expr) { - if (expr->op == ast::BinaryOp::kLogicalAnd || expr->op == ast::BinaryOp::kLogicalOr) { - return EmitShortCircuit(expr); - } - - auto lhs = EmitExpression(expr->lhs); - if (!lhs) { - return utils::Failure; - } - - auto rhs = EmitExpression(expr->rhs); - if (!rhs) { - return utils::Failure; - } - - auto* sem = program_->Sem().Get(expr); - auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); - - Binary* inst = nullptr; - switch (expr->op) { - case ast::BinaryOp::kAnd: - inst = builder.And(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kOr: - inst = builder.Or(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kXor: - inst = builder.Xor(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kEqual: - inst = builder.Equal(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kNotEqual: - inst = builder.NotEqual(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kLessThan: - inst = builder.LessThan(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kGreaterThan: - inst = builder.GreaterThan(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kLessThanEqual: - inst = builder.LessThanEqual(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kGreaterThanEqual: - inst = builder.GreaterThanEqual(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kShiftLeft: - inst = builder.ShiftLeft(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kShiftRight: - inst = builder.ShiftRight(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kAdd: - inst = builder.Add(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kSubtract: - inst = builder.Subtract(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kMultiply: - inst = builder.Multiply(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kDivide: - inst = builder.Divide(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kModulo: - inst = builder.Modulo(ty, lhs.Get(), rhs.Get()); - break; - case ast::BinaryOp::kLogicalAnd: - case ast::BinaryOp::kLogicalOr: - TINT_ICE(IR, diagnostics_) << "short circuit op should have already been handled"; - return utils::Failure; - case ast::BinaryOp::kNone: - TINT_ICE(IR, diagnostics_) << "missing binary operand type"; - return utils::Failure; - } - - current_flow_block->instructions.Push(inst); - return inst; -} - -utils::Result BuilderImpl::EmitBitcast(const ast::BitcastExpression* expr) { - auto val = EmitExpression(expr->expr); - if (!val) { - return utils::Failure; - } - - auto* sem = program_->Sem().Get(expr); - auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); - auto* inst = builder.Bitcast(ty, val.Get()); - - current_flow_block->instructions.Push(inst); - return inst; -} - -void BuilderImpl::EmitCall(const ast::CallStatement* stmt) { - (void)EmitCall(stmt->expr); -} - -utils::Result BuilderImpl::EmitCall(const ast::CallExpression* expr) { - // If this is a materialized semantic node, just use the constant value. - if (auto* mat = program_->Sem().Get(expr)) { - if (mat->ConstantValue()) { - auto* cv = mat->ConstantValue()->Clone(clone_ctx_); - if (!cv) { - add_error(expr->source, "failed to get constant value for call " + - std::string(expr->TypeInfo().name)); - return utils::Failure; - } - return builder.Constant(cv); - } - } - - utils::Vector args; - args.Reserve(expr->args.Length()); - - // Emit the arguments - for (const auto* arg : expr->args) { - auto value = EmitExpression(arg); - if (!value) { - add_error(arg->source, "failed to convert arguments"); - return utils::Failure; - } - args.Push(value.Get()); - } - - auto* sem = program_->Sem().Get(expr); - if (!sem) { - add_error(expr->source, "failed to get semantic information for call " + - std::string(expr->TypeInfo().name)); - return utils::Failure; - } - - auto* ty = sem->Target()->ReturnType()->Clone(clone_ctx_.type_ctx); - - Instruction* inst = nullptr; - - // If this is a builtin function, emit the specific builtin value - if (auto* b = sem->Target()->As()) { - inst = builder.Builtin(ty, b->Type(), args); - } else if (sem->Target()->As()) { - inst = builder.Construct(ty, std::move(args)); - } else if (auto* conv = sem->Target()->As()) { - auto* from = conv->Source()->Clone(clone_ctx_.type_ctx); - inst = builder.Convert(ty, from, std::move(args)); - } else if (expr->target->identifier->Is()) { - TINT_UNIMPLEMENTED(IR, diagnostics_) << "missing templated ident support"; - return utils::Failure; - } else { - // Not a builtin and not a templated call, so this is a user function. - auto name = CloneSymbol(expr->target->identifier->symbol); - inst = builder.UserCall(ty, name, std::move(args)); - } - if (inst == nullptr) { - return utils::Failure; - } - current_flow_block->instructions.Push(inst); - return inst; -} - -utils::Result BuilderImpl::EmitLiteral(const ast::LiteralExpression* lit) { - auto* sem = program_->Sem().Get(lit); - if (!sem) { - add_error(lit->source, "failed to get semantic information for node " + - std::string(lit->TypeInfo().name)); - return utils::Failure; - } - - auto* cv = sem->ConstantValue()->Clone(clone_ctx_); - if (!cv) { - add_error(lit->source, - "failed to get constant value for node " + std::string(lit->TypeInfo().name)); - return utils::Failure; - } - return builder.Constant(cv); -} - -void BuilderImpl::EmitAttributes(utils::VectorRef attrs) { - for (auto* attr : attrs) { - EmitAttribute(attr); - } -} - -void BuilderImpl::EmitAttribute(const ast::Attribute* attr) { - tint::Switch( // - attr, - // [&](const ast::WorkgroupAttribute* wg) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::StageAttribute* s) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::BindingAttribute* b) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::GroupAttribute* g) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::LocationAttribute* l) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::BuiltinAttribute* b) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::InterpolateAttribute* i) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::InvariantAttribute* i) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::MustUseAttribute* i) { - // TODO(dsinclair): Implement - // }, - [&](const ast::IdAttribute*) { - add_error(attr->source, - "found an `Id` attribute. The SubstituteOverrides transform " - "must be run before converting to IR"); - }, - [&](const ast::StructMemberSizeAttribute*) { - TINT_ICE(IR, diagnostics_) - << "StructMemberSizeAttribute encountered during IR conversion"; - }, - [&](const ast::StructMemberAlignAttribute*) { - TINT_ICE(IR, diagnostics_) - << "StructMemberAlignAttribute encountered during IR conversion"; - }, - // [&](const ast::StrideAttribute* s) { - // TODO(dsinclair): Implement - // }, - // [&](const ast::InternalAttribute *i) { - // TODO(dsinclair): Implement - // }, - [&](Default) { - add_error(attr->source, "unknown attribute: " + std::string(attr->TypeInfo().name)); - }); -} - -} // namespace tint::ir diff --git a/src/tint/ir/builder_impl.h b/src/tint/ir/builder_impl.h deleted file mode 100644 index a83b76b711..0000000000 --- a/src/tint/ir/builder_impl.h +++ /dev/null @@ -1,245 +0,0 @@ -// Copyright 2022 The Tint Authors. -// -// Licensed under the Apache License, Version 2.0 (the "License"); -// you may not use this file except in compliance with the License. -// You may obtain a copy of the License at -// -// http://www.apache.org/licenses/LICENSE-2.0 -// -// Unless required by applicable law or agreed to in writing, software -// distributed under the License is distributed on an "AS IS" BASIS, -// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. -// See the License for the specific language governing permissions and -// limitations under the License. - -#ifndef SRC_TINT_IR_BUILDER_IMPL_H_ -#define SRC_TINT_IR_BUILDER_IMPL_H_ - -#include -#include -#include - -#include "src/tint/ast/type.h" -#include "src/tint/constant/clone_context.h" -#include "src/tint/diagnostic/diagnostic.h" -#include "src/tint/ir/builder.h" -#include "src/tint/ir/flow_node.h" -#include "src/tint/ir/module.h" -#include "src/tint/ir/value.h" -#include "src/tint/scope_stack.h" -#include "src/tint/utils/result.h" - -// Forward Declarations -namespace tint { -class Program; -} // namespace tint -namespace tint::ast { -class Attribute; -class AssignmentStatement; -class BinaryExpression; -class BitcastExpression; -class BlockStatement; -class BreakIfStatement; -class BreakStatement; -class CallExpression; -class CallStatement; -class CompoundAssignmentStatement; -class ContinueStatement; -class DiscardStatement; -class Expression; -class ForLoopStatement; -class Function; -class IfStatement; -class LoopStatement; -class LiteralExpression; -class Node; -class ReturnStatement; -class Statement; -class SwitchStatement; -class UnaryOpExpression; -class WhileStatement; -class Variable; -} // namespace tint::ast -namespace tint::sem { -class Builtin; -} // namespace tint::sem - -namespace tint::ir { - -/// Builds an ir::Module from a given ast::Program -class BuilderImpl { - public: - /// Constructor - /// @param program the program to create from - explicit BuilderImpl(const Program* program); - /// Destructor - ~BuilderImpl(); - - /// Builds an ir::Module from the given Program - /// @returns true on success, false otherwise - utils::Result Build(); - - /// @returns the diagnostics - diag::List Diagnostics() const { return diagnostics_; } - - /// Emits a function to the IR. - /// @param func the function to emit - void EmitFunction(const ast::Function* func); - - /// Emits a set of statements to the IR. - /// @param stmts the statements to emit - void EmitStatements(utils::VectorRef stmts); - - /// Emits a statement to the IR - /// @param stmt the statment to emit - void EmitStatement(const ast::Statement* stmt); - - /// Emits a block statement to the IR. - /// @param block the block to emit - void EmitBlock(const ast::BlockStatement* block); - - /// Emits an if control node to the IR. - /// @param stmt the if statement - void EmitIf(const ast::IfStatement* stmt); - - /// Emits a return node to the IR. - /// @param stmt the return AST statement - void EmitReturn(const ast::ReturnStatement* stmt); - - /// Emits a loop control node to the IR. - /// @param stmt the loop statement - void EmitLoop(const ast::LoopStatement* stmt); - - /// Emits a loop control node to the IR. - /// @param stmt the while statement - void EmitWhile(const ast::WhileStatement* stmt); - - /// Emits a loop control node to the IR. - /// @param stmt the for loop statement - void EmitForLoop(const ast::ForLoopStatement* stmt); - - /// Emits a switch statement - /// @param stmt the switch statement - void EmitSwitch(const ast::SwitchStatement* stmt); - - /// Emits a break statement - /// @param stmt the break statement - void EmitBreak(const ast::BreakStatement* stmt); - - /// Emits a continue statement - /// @param stmt the continue statement - void EmitContinue(const ast::ContinueStatement* stmt); - - /// Emits a discard statement - void EmitDiscard(const ast::DiscardStatement*); - - /// Emits a break-if statement - /// @param stmt the break-if statement - void EmitBreakIf(const ast::BreakIfStatement* stmt); - - /// Emits an assignment statement - /// @param stmt the statement - void EmitAssignment(const ast::AssignmentStatement* stmt); - - /// Emits a compound assignment statement - /// @param stmt the statement - void EmitCompoundAssignment(const ast::CompoundAssignmentStatement* stmt); - - /// Emits an expression - /// @param expr the expression to emit - /// @returns true if successful, false otherwise - utils::Result EmitExpression(const ast::Expression* expr); - - /// Emits a variable - /// @param var the variable to emit - void EmitVariable(const ast::Variable* var); - - /// Emits a Unary expression - /// @param expr the unary expression - /// @returns the value storing the result if successful, utils::Failure otherwise - utils::Result EmitUnary(const ast::UnaryOpExpression* expr); - - /// Emits a short-circult binary expression - /// @param expr the binary expression - /// @returns the value storing the result if successful, utils::Failure otherwise - utils::Result EmitShortCircuit(const ast::BinaryExpression* expr); - - /// Emits a binary expression - /// @param expr the binary expression - /// @returns the value storing the result if successful, utils::Failure otherwise - utils::Result EmitBinary(const ast::BinaryExpression* expr); - - /// Emits a bitcast expression - /// @param expr the bitcast expression - /// @returns the value storing the result if successful, utils::Failure otherwise - utils::Result EmitBitcast(const ast::BitcastExpression* expr); - - /// Emits a call expression - /// @param stmt the call statement - void EmitCall(const ast::CallStatement* stmt); - - /// Emits a call expression - /// @param expr the call expression - /// @returns the value storing the result if successful, utils::Failure otherwise - utils::Result EmitCall(const ast::CallExpression* expr); - - /// Emits a literal expression - /// @param lit the literal to emit - /// @returns true if successful, false otherwise - utils::Result EmitLiteral(const ast::LiteralExpression* lit); - - /// Emits a set of attributes - /// @param attrs the attributes to emit - void EmitAttributes(utils::VectorRef attrs); - - /// Emits an attribute - /// @param attr the attribute to emit - void EmitAttribute(const ast::Attribute* attr); - - /// Retrieve the IR Flow node for a given AST node. - /// @param n the node to lookup - /// @returns the FlowNode for the given ast::Node or nullptr if it doesn't exist. - const ir::FlowNode* FlowNodeForAstNode(const ast::Node* n) const { - if (ast_to_flow_.count(n) == 0) { - return nullptr; - } - return ast_to_flow_.at(n); - } - - /// The stack of flow control blocks. - utils::Vector flow_stack; - - /// The IR module being built - Module mod; - /// The IR builder being used by the impl. - Builder builder{mod}; - - /// The current flow block for expressions - Block* current_flow_block = nullptr; - - private: - enum class ControlFlags { kNone, kExcludeSwitch }; - - void BranchTo(ir::FlowNode* node, utils::VectorRef args = {}); - void BranchToIfNeeded(ir::FlowNode* node); - - FlowNode* FindEnclosingControl(ControlFlags flags); - - void add_error(const Source& s, const std::string& err); - - Symbol CloneSymbol(Symbol sym) const; - - const Program* program_ = nullptr; - Function* current_function_ = nullptr; - ScopeStack scopes_; - constant::CloneContext clone_ctx_; - diag::List diagnostics_; - - /// Map from ast nodes to flow nodes, used to retrieve the flow node for a given AST node. - /// Used for testing purposes. - std::unordered_map ast_to_flow_; -}; - -} // namespace tint::ir - -#endif // SRC_TINT_IR_BUILDER_IMPL_H_ diff --git a/src/tint/ir/constant_test.cc b/src/tint/ir/constant_test.cc index f5641efc7a..7005751913 100644 --- a/src/tint/ir/constant_test.cc +++ b/src/tint/ir/constant_test.cc @@ -24,7 +24,8 @@ using namespace tint::number_suffixes; // NOLINT using IR_ConstantTest = TestHelper; TEST_F(IR_ConstantTest, f32) { - Builder b; + Module mod; + Builder b{mod}; utils::StringStream str; @@ -39,7 +40,8 @@ TEST_F(IR_ConstantTest, f32) { } TEST_F(IR_ConstantTest, f16) { - Builder b; + Module mod; + Builder b{mod}; utils::StringStream str; @@ -54,7 +56,8 @@ TEST_F(IR_ConstantTest, f16) { } TEST_F(IR_ConstantTest, i32) { - Builder b; + Module mod; + Builder b{mod}; utils::StringStream str; @@ -69,7 +72,8 @@ TEST_F(IR_ConstantTest, i32) { } TEST_F(IR_ConstantTest, u32) { - Builder b; + Module mod; + Builder b{mod}; utils::StringStream str; @@ -84,7 +88,8 @@ TEST_F(IR_ConstantTest, u32) { } TEST_F(IR_ConstantTest, bool) { - Builder b; + Module mod; + Builder b{mod}; { utils::StringStream str; diff --git a/src/tint/ir/discard_test.cc b/src/tint/ir/discard_test.cc index 942f02c78a..054727f2e2 100644 --- a/src/tint/ir/discard_test.cc +++ b/src/tint/ir/discard_test.cc @@ -22,7 +22,8 @@ namespace { using IR_InstructionTest = TestHelper; TEST_F(IR_InstructionTest, Discard) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Discard(); ASSERT_TRUE(inst->Is()); diff --git a/src/tint/ir/from_program.cc b/src/tint/ir/from_program.cc index d093103acd..d5c6b416ec 100644 --- a/src/tint/ir/from_program.cc +++ b/src/tint/ir/from_program.cc @@ -14,17 +14,1168 @@ #include "src/tint/ir/from_program.h" -#include "src/tint/ir/builder_impl.h" +#include +#include +#include + +#include "src/tint/ast/alias.h" +#include "src/tint/ast/assignment_statement.h" +#include "src/tint/ast/binary_expression.h" +#include "src/tint/ast/bitcast_expression.h" +#include "src/tint/ast/block_statement.h" +#include "src/tint/ast/bool_literal_expression.h" +#include "src/tint/ast/break_if_statement.h" +#include "src/tint/ast/break_statement.h" +#include "src/tint/ast/call_expression.h" +#include "src/tint/ast/call_statement.h" +#include "src/tint/ast/compound_assignment_statement.h" +#include "src/tint/ast/const.h" +#include "src/tint/ast/const_assert.h" +#include "src/tint/ast/continue_statement.h" +#include "src/tint/ast/discard_statement.h" +#include "src/tint/ast/enable.h" +#include "src/tint/ast/float_literal_expression.h" +#include "src/tint/ast/for_loop_statement.h" +#include "src/tint/ast/function.h" +#include "src/tint/ast/id_attribute.h" +#include "src/tint/ast/identifier.h" +#include "src/tint/ast/identifier_expression.h" +#include "src/tint/ast/if_statement.h" +#include "src/tint/ast/int_literal_expression.h" +#include "src/tint/ast/invariant_attribute.h" +#include "src/tint/ast/let.h" +#include "src/tint/ast/literal_expression.h" +#include "src/tint/ast/loop_statement.h" +#include "src/tint/ast/override.h" +#include "src/tint/ast/return_statement.h" +#include "src/tint/ast/statement.h" +#include "src/tint/ast/struct.h" +#include "src/tint/ast/struct_member_align_attribute.h" +#include "src/tint/ast/struct_member_size_attribute.h" +#include "src/tint/ast/switch_statement.h" +#include "src/tint/ast/templated_identifier.h" +#include "src/tint/ast/unary_op_expression.h" +#include "src/tint/ast/var.h" +#include "src/tint/ast/variable_decl_statement.h" +#include "src/tint/ast/while_statement.h" +#include "src/tint/ir/builder.h" +#include "src/tint/ir/function.h" +#include "src/tint/ir/if.h" +#include "src/tint/ir/loop.h" +#include "src/tint/ir/module.h" +#include "src/tint/ir/store.h" +#include "src/tint/ir/switch.h" +#include "src/tint/ir/value.h" #include "src/tint/program.h" +#include "src/tint/scope_stack.h" +#include "src/tint/sem/builtin.h" +#include "src/tint/sem/call.h" +#include "src/tint/sem/function.h" +#include "src/tint/sem/materialize.h" +#include "src/tint/sem/module.h" +#include "src/tint/sem/switch_statement.h" +#include "src/tint/sem/value_constructor.h" +#include "src/tint/sem/value_conversion.h" +#include "src/tint/sem/value_expression.h" +#include "src/tint/sem/variable.h" +#include "src/tint/switch.h" +#include "src/tint/type/void.h" +#include "src/tint/utils/defer.h" +#include "src/tint/utils/result.h" +#include "src/tint/utils/scoped_assignment.h" namespace tint::ir { +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) { + if (IsConnected(parent)) { + return true; + } + } + // Getting here means all the incoming branches are disconnected. + return false; +} + +class Impl { + public: + explicit Impl(const Program* program) + : program_(program), + clone_ctx_{ + type::CloneContext{{&program->Symbols()}, {&builder.ir.symbols, &builder.ir.types}}, + {&builder.ir.constants}} {} + + ResultType Build() { + auto* sem = program_->Sem().Module(); + + for (auto* decl : sem->DependencyOrderedDeclarations()) { + tint::Switch( + decl, // + [&](const ast::Struct*) { + // Will be encoded into the `type::Struct` when used. We will then hoist all + // used structs up to module scope when converting IR. + }, + [&](const ast::Alias*) { + // Folded away and doesn't appear in the IR. + }, + [&](const ast::Variable* var) { + // Setup the current flow node to be the root block for the module. The builder + // will handle creating it if it doesn't exist already. + TINT_SCOPED_ASSIGNMENT(current_flow_block, builder.CreateRootBlockIfNeeded()); + EmitVariable(var); + }, + [&](const ast::Function* func) { EmitFunction(func); }, + [&](const ast::Enable*) { + // TODO(dsinclair): Implement? I think these need to be passed along so further + // stages know what is enabled. + }, + [&](const ast::ConstAssert*) { + // Evaluated by the resolver, drop from the IR. + }, + [&](Default) { + add_error(decl->source, "unknown type: " + std::string(decl->TypeInfo().name)); + }); + } + if (!diagnostics_.empty()) { + return utils::Failure; + } + + return ResultType{std::move(mod)}; + } + + /// @returns the diagnostics + diag::List Diagnostics() const { return diagnostics_; } + + private: + enum class ControlFlags { kNone, kExcludeSwitch }; + + /// The stack of flow control blocks. + utils::Vector flow_stack; + + /// The IR module being built + Module mod; + + /// The IR builder being used by the impl. + Builder builder{mod}; + + /// The current flow block for expressions + Block* current_flow_block = nullptr; + + const Program* program_ = nullptr; + Function* current_function_ = nullptr; + ScopeStack scopes_; + constant::CloneContext clone_ctx_; + diag::List diagnostics_; + + /// Map from ast nodes to flow nodes, used to retrieve the flow node for a given AST node. + /// Used for testing purposes. + std::unordered_map ast_to_flow_; + + class FlowStackScope { + public: + FlowStackScope(Impl* impl, FlowNode* node) : impl_(impl) { impl_->flow_stack.Push(node); } + + ~FlowStackScope() { impl_->flow_stack.Pop(); } + + private: + Impl* impl_; + }; + + void add_error(const Source& s, const std::string& err) { + diagnostics_.add_error(tint::diag::System::IR, err, s); + } + + void BranchTo(FlowNode* node, utils::VectorRef args = {}) { + TINT_ASSERT(IR, current_flow_block); + TINT_ASSERT(IR, !IsBranched(current_flow_block)); + + builder.Branch(current_flow_block, node, args); + current_flow_block = nullptr; + } + + void BranchToIfNeeded(FlowNode* node) { + if (!current_flow_block || IsBranched(current_flow_block)) { + return; + } + BranchTo(node); + } + + FlowNode* FindEnclosingControl(ControlFlags flags) { + for (auto it = flow_stack.rbegin(); it != flow_stack.rend(); ++it) { + if ((*it)->Is()) { + return *it; + } + if (flags == ControlFlags::kExcludeSwitch) { + continue; + } + if ((*it)->Is()) { + return *it; + } + } + return nullptr; + } + + Symbol CloneSymbol(Symbol sym) const { + return clone_ctx_.type_ctx.dst.st->Register(sym.Name()); + } + + void EmitFunction(const ast::Function* ast_func) { + // The flow stack should have been emptied when the previous function finished building. + TINT_ASSERT(IR, flow_stack.IsEmpty()); + + const auto* sem = program_->Sem().Get(ast_func); + + auto* ir_func = builder.CreateFunction(CloneSymbol(ast_func->name->symbol), + sem->ReturnType()->Clone(clone_ctx_.type_ctx)); + current_function_ = ir_func; + builder.ir.functions.Push(ir_func); + + ast_to_flow_[ast_func] = ir_func; + + if (ast_func->IsEntryPoint()) { + builder.ir.entry_points.Push(ir_func); + + switch (ast_func->PipelineStage()) { + case ast::PipelineStage::kVertex: + ir_func->pipeline_stage = Function::PipelineStage::kVertex; + break; + case ast::PipelineStage::kFragment: + ir_func->pipeline_stage = Function::PipelineStage::kFragment; + break; + case ast::PipelineStage::kCompute: { + ir_func->pipeline_stage = 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), + }; + break; + } + default: { + TINT_ICE(IR, diagnostics_) << "Invalid pipeline stage"; + return; + } + } + + for (auto* attr : ast_func->return_type_attributes) { + tint::Switch( + attr, // + [&](const ast::LocationAttribute*) { + ir_func->return_attributes.Push(Function::ReturnAttribute::kLocation); + }, + [&](const ast::InvariantAttribute*) { + ir_func->return_attributes.Push(Function::ReturnAttribute::kInvariant); + }, + [&](const ast::BuiltinAttribute* b) { + if (auto* ident_sem = + program_->Sem() + .Get(b) + ->As>()) { + switch (ident_sem->Value()) { + case builtin::BuiltinValue::kPosition: + ir_func->return_attributes.Push( + Function::ReturnAttribute::kPosition); + break; + case builtin::BuiltinValue::kFragDepth: + ir_func->return_attributes.Push( + Function::ReturnAttribute::kFragDepth); + break; + case builtin::BuiltinValue::kSampleMask: + ir_func->return_attributes.Push( + Function::ReturnAttribute::kSampleMask); + break; + default: + TINT_ICE(IR, diagnostics_) + << "Unknown builtin value in return attributes " + << ident_sem->Value(); + return; + } + } else { + TINT_ICE(IR, diagnostics_) << "Builtin attribute sem invalid"; + return; + } + }); + } + } + ir_func->return_location = sem->ReturnLocation(); + + { + FlowStackScope scope(this, ir_func); + + current_flow_block = ir_func->start_target; + EmitBlock(ast_func->body); + + // TODO(dsinclair): Store return type and attributes + // TODO(dsinclair): Store parameters + + // If the branch target has already been set then a `return` was called. Only set in the + // case where `return` wasn't called. + BranchToIfNeeded(current_function_->end_target); + } + + TINT_ASSERT(IR, flow_stack.IsEmpty()); + current_flow_block = nullptr; + current_function_ = nullptr; + } + + void EmitStatements(utils::VectorRef stmts) { + for (auto* s : stmts) { + EmitStatement(s); + + // 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)) { + break; + } + } + } + + void EmitStatement(const ast::Statement* stmt) { + tint::Switch( + stmt, // + [&](const ast::AssignmentStatement* a) { EmitAssignment(a); }, + [&](const ast::BlockStatement* b) { EmitBlock(b); }, + [&](const ast::BreakStatement* b) { EmitBreak(b); }, + [&](const ast::BreakIfStatement* b) { EmitBreakIf(b); }, + [&](const ast::CallStatement* c) { EmitCall(c); }, + [&](const ast::CompoundAssignmentStatement* c) { EmitCompoundAssignment(c); }, + [&](const ast::ContinueStatement* c) { EmitContinue(c); }, + [&](const ast::DiscardStatement* d) { EmitDiscard(d); }, + [&](const ast::IfStatement* i) { EmitIf(i); }, + [&](const ast::LoopStatement* l) { EmitLoop(l); }, + [&](const ast::ForLoopStatement* l) { EmitForLoop(l); }, + [&](const ast::WhileStatement* l) { EmitWhile(l); }, + [&](const ast::ReturnStatement* r) { EmitReturn(r); }, + [&](const ast::SwitchStatement* s) { EmitSwitch(s); }, + [&](const ast::VariableDeclStatement* v) { EmitVariable(v->variable); }, + [&](const ast::ConstAssert*) { + // Not emitted + }, + [&](Default) { + add_error(stmt->source, + "unknown statement type: " + std::string(stmt->TypeInfo().name)); + }); + } + + void EmitAssignment(const ast::AssignmentStatement* stmt) { + auto lhs = EmitExpression(stmt->lhs); + if (!lhs) { + return; + } + + auto rhs = EmitExpression(stmt->rhs); + if (!rhs) { + return; + } + auto store = builder.Store(lhs.Get(), rhs.Get()); + current_flow_block->instructions.Push(store); + } + + void EmitCompoundAssignment(const ast::CompoundAssignmentStatement* stmt) { + auto lhs = EmitExpression(stmt->lhs); + if (!lhs) { + return; + } + + auto rhs = EmitExpression(stmt->rhs); + if (!rhs) { + return; + } + + auto* ty = lhs.Get()->Type(); + Binary* inst = nullptr; + switch (stmt->op) { + case ast::BinaryOp::kAnd: + inst = builder.And(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kOr: + inst = builder.Or(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kXor: + inst = builder.Xor(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kShiftLeft: + inst = builder.ShiftLeft(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kShiftRight: + inst = builder.ShiftRight(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kAdd: + inst = builder.Add(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kSubtract: + inst = builder.Subtract(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kMultiply: + inst = builder.Multiply(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kDivide: + inst = builder.Divide(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kModulo: + inst = builder.Modulo(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kLessThanEqual: + case ast::BinaryOp::kGreaterThanEqual: + case ast::BinaryOp::kGreaterThan: + case ast::BinaryOp::kLessThan: + case ast::BinaryOp::kNotEqual: + case ast::BinaryOp::kEqual: + case ast::BinaryOp::kLogicalAnd: + case ast::BinaryOp::kLogicalOr: + TINT_ICE(IR, diagnostics_) << "invalid compound assignment"; + return; + case ast::BinaryOp::kNone: + TINT_ICE(IR, diagnostics_) << "missing binary operand type"; + return; + } + current_flow_block->instructions.Push(inst); + + auto store = builder.Store(lhs.Get(), inst); + current_flow_block->instructions.Push(store); + } + + void EmitBlock(const ast::BlockStatement* block) { + scopes_.Push(); + TINT_DEFER(scopes_.Pop()); + + // Note, this doesn't need to emit a Block as the current block flow node should be + // sufficient as the blocks all get flattened. Each flow control node will inject the basic + // blocks it requires. + EmitStatements(block->statements); + } + + void EmitIf(const ast::IfStatement* stmt) { + // Emit the if condition into the end of the preceding block + auto reg = EmitExpression(stmt->condition); + if (!reg) { + return; + } + auto* if_node = builder.CreateIf(reg.Get()); + + BranchTo(if_node); + + ast_to_flow_[stmt] = if_node; + + { + FlowStackScope scope(this, if_node); + + 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); + + 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); + } + 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(); + } + } + + void EmitLoop(const ast::LoopStatement* stmt) { + auto* loop_node = builder.CreateLoop(); + + BranchTo(loop_node); + + ast_to_flow_[stmt] = loop_node; + + { + FlowStackScope scope(this, loop_node); + + current_flow_block = loop_node->start.target->As(); + EmitBlock(stmt->body); + + // The current block didn't `break`, `return` or `continue`, go to the continuing block. + BranchToIfNeeded(loop_node->continuing.target); + + 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); + } + + // 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)) { + current_flow_block = nullptr; + } + } + + 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); + + BranchTo(loop_node); + + ast_to_flow_[stmt] = loop_node; + + { + FlowStackScope scope(this, loop_node); + + current_flow_block = loop_node->start.target->As(); + + // Emit the while condition into the start target of the loop + auto reg = EmitExpression(stmt->condition); + if (!reg) { + return; + } + + // 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); + + BranchTo(if_node); + + current_flow_block = if_node->merge.target->As(); + EmitBlock(stmt->body); + + 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(); + } + + 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); + + // Make sure the initializer ends up in a contained scope + scopes_.Push(); + TINT_DEFER(scopes_.Pop()); + + if (stmt->initializer) { + // Emit the for initializer before branching to the loop + EmitStatement(stmt->initializer); + } + + BranchTo(loop_node); + + ast_to_flow_[stmt] = loop_node; + + { + FlowStackScope scope(this, loop_node); + + current_flow_block = loop_node->start.target->As(); + + if (stmt->condition) { + // Emit the condition into the target target of the loop + auto reg = EmitExpression(stmt->condition); + if (!reg) { + return; + } + + // 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); + + BranchTo(if_node); + current_flow_block = if_node->merge.target->As(); + } + + EmitBlock(stmt->body); + BranchToIfNeeded(loop_node->continuing.target); + + if (stmt->continuing) { + 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(); + } + + void EmitSwitch(const ast::SwitchStatement* stmt) { + // Emit the condition into the preceding block + auto reg = EmitExpression(stmt->condition); + if (!reg) { + return; + } + auto* switch_node = builder.CreateSwitch(reg.Get()); + + BranchTo(switch_node); + + ast_to_flow_[stmt] = switch_node; + + { + FlowStackScope scope(this, switch_node); + + const auto* sem = program_->Sem().Get(stmt); + for (const auto* c : sem->Cases()) { + utils::Vector selectors; + for (const auto* selector : c->Selectors()) { + if (selector->IsDefault()) { + selectors.Push({nullptr}); + } else { + selectors.Push({builder.Constant(selector->Value()->Clone(clone_ctx_))}); + } + } + + current_flow_block = builder.CreateCase(switch_node, selectors); + EmitBlock(c->Body()->Declaration()); + + BranchToIfNeeded(switch_node->merge.target); + } + } + current_flow_block = nullptr; + + if (IsConnected(switch_node->merge.target)) { + current_flow_block = switch_node->merge.target->As(); + } + } + + void EmitReturn(const ast::ReturnStatement* stmt) { + utils::Vector ret_value; + if (stmt->value) { + auto ret = EmitExpression(stmt->value); + if (!ret) { + return; + } + ret_value.Push(ret.Get()); + } + + BranchTo(current_function_->end_target, std::move(ret_value)); + } + + void EmitBreak(const ast::BreakStatement*) { + auto* current_control = FindEnclosingControl(ControlFlags::kNone); + TINT_ASSERT(IR, current_control); + + if (auto* c = current_control->As()) { + BranchTo(c->merge.target); + } else if (auto* s = current_control->As()) { + BranchTo(s->merge.target); + } else { + TINT_UNREACHABLE(IR, diagnostics_); + } + } + + void EmitContinue(const ast::ContinueStatement*) { + auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch); + TINT_ASSERT(IR, current_control); + + if (auto* c = current_control->As()) { + BranchTo(c->continuing.target); + } else { + TINT_UNREACHABLE(IR, diagnostics_); + } + } + + // Discard is being treated as an instruction. The semantics in WGSL is demote_to_helper, so the + // code has to continue as before it just predicates writes. If WGSL grows some kind of + // terminating discard that would probably make sense as a FlowNode but would then require + // figuring out the multi-level exit that is triggered. + void EmitDiscard(const ast::DiscardStatement*) { + auto* inst = builder.Discard(); + current_flow_block->instructions.Push(inst); + } + + void EmitBreakIf(const ast::BreakIfStatement* stmt) { + // Emit the break-if condition into the end of the preceding block + auto reg = EmitExpression(stmt->condition); + if (!reg) { + return; + } + auto* if_node = builder.CreateIf(reg.Get()); + + BranchTo(if_node); + + ast_to_flow_[stmt] = if_node; + + auto* current_control = FindEnclosingControl(ControlFlags::kExcludeSwitch); + TINT_ASSERT(IR, current_control); + TINT_ASSERT(IR, current_control->Is()); + + auto* loop = current_control->As(); + + 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->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); + } + + utils::Result EmitExpression(const ast::Expression* expr) { + // If this is a value that has been const-eval'd return the result. + if (auto* sem = program_->Sem().Get(expr)->As()) { + if (auto* v = sem->ConstantValue()) { + if (auto* cv = v->Clone(clone_ctx_)) { + return builder.Constant(cv); + } + } + } + + return tint::Switch( + expr, + // [&](const ast::IndexAccessorExpression* a) { + // TODO(dsinclair): Implement + // }, + [&](const ast::BinaryExpression* b) { return EmitBinary(b); }, + [&](const ast::BitcastExpression* b) { return EmitBitcast(b); }, + [&](const ast::CallExpression* c) { return EmitCall(c); }, + [&](const ast::IdentifierExpression* i) { + auto* v = scopes_.Get(i->identifier->symbol); + return utils::Result{v}; + }, + [&](const ast::LiteralExpression* l) { return EmitLiteral(l); }, + // [&](const ast::MemberAccessorExpression* m) { + // TODO(dsinclair): Implement + // }, + // [&](const ast::PhonyExpression*) { + // TODO(dsinclair): Implement. The call may have side effects so has to be made. + // }, + [&](const ast::UnaryOpExpression* u) { return EmitUnary(u); }, + [&](Default) { + add_error(expr->source, + "unknown expression type: " + std::string(expr->TypeInfo().name)); + return utils::Failure; + }); + } + + void EmitVariable(const ast::Variable* var) { + auto* sem = program_->Sem().Get(var); + + return tint::Switch( // + var, + [&](const ast::Var* v) { + auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); + auto* val = builder.Declare(ty, sem->AddressSpace(), sem->Access()); + current_flow_block->instructions.Push(val); + + if (v->initializer) { + auto init = EmitExpression(v->initializer); + if (!init) { + return; + } + val->initializer = init.Get(); + } + // Store the declaration so we can get the instruction to store too + scopes_.Set(v->name->symbol, val); + + // Record the original name of the var + builder.ir.SetName(val, v->name->symbol.Name()); + }, + [&](const ast::Let* l) { + // A `let` doesn't exist as a standalone item in the IR, it's just the result of the + // initializer. + auto init = EmitExpression(l->initializer); + if (!init) { + return; + } + + // Store the results of the initialization + scopes_.Set(l->name->symbol, init.Get()); + + // Record the original name of the let + builder.ir.SetName(init.Get(), l->name->symbol.Name()); + }, + [&](const ast::Override*) { + add_error(var->source, + "found an `Override` variable. The SubstituteOverrides " + "transform must be run before converting to IR"); + }, + [&](const ast::Const*) { + // Skip. This should be handled by const-eval already, so the const will be a + // `constant::` value at the usage sites. Can just ignore the `const` variable as it + // should never be used. + // + // TODO(dsinclair): Probably want to store the const variable somewhere and then in + // identifier expression log an error if we ever see a const identifier. Add this + // when identifiers and variables are supported. + }, + [&](Default) { + add_error(var->source, "unknown variable: " + std::string(var->TypeInfo().name)); + }); + } + + utils::Result EmitUnary(const ast::UnaryOpExpression* expr) { + auto val = EmitExpression(expr->expr); + if (!val) { + return utils::Failure; + } + + auto* sem = program_->Sem().Get(expr); + auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); + + Instruction* inst = nullptr; + switch (expr->op) { + case ast::UnaryOp::kAddressOf: + inst = builder.AddressOf(ty, val.Get()); + break; + case ast::UnaryOp::kComplement: + inst = builder.Complement(ty, val.Get()); + break; + case ast::UnaryOp::kIndirection: + inst = builder.Indirection(ty, val.Get()); + break; + case ast::UnaryOp::kNegation: + inst = builder.Negation(ty, val.Get()); + break; + case ast::UnaryOp::kNot: + inst = builder.Not(ty, val.Get()); + break; + } + + current_flow_block->instructions.Push(inst); + return inst; + } + + // A short-circut needs special treatment. The short-circuit is decomposed into the relevant if + // statements and declarations. + utils::Result EmitShortCircuit(const ast::BinaryExpression* expr) { + switch (expr->op) { + case ast::BinaryOp::kLogicalAnd: + case ast::BinaryOp::kLogicalOr: + break; + default: + TINT_ICE(IR, diagnostics_) + << "invalid operation type for short-circut decomposition"; + return utils::Failure; + } + + // Evaluate the LHS of the short-circuit + auto lhs = EmitExpression(expr->lhs); + if (!lhs) { + return utils::Failure; + } + + // Generate a variable to store the short-circut into + auto* ty = builder.ir.types.Get(); + auto* result_var = + builder.Declare(ty, builtin::AddressSpace::kFunction, builtin::Access::kReadWrite); + current_flow_block->instructions.Push(result_var); + + auto* lhs_store = builder.Store(result_var, lhs.Get()); + current_flow_block->instructions.Push(lhs_store); + + auto* if_node = builder.CreateIf(lhs.Get()); + BranchTo(if_node); + + utils::Result rhs; + { + FlowStackScope scope(this, if_node); + + // If this is an `&&` then we only evaluate the RHS expression in the true block. + // If this is an `||` then we only evaluate the RHS expression in the false block. + if (expr->op == ast::BinaryOp::kLogicalAnd) { + current_flow_block = if_node->true_.target->As(); + } else { + current_flow_block = if_node->false_.target->As(); + } + + rhs = EmitExpression(expr->rhs); + if (!rhs) { + return utils::Failure; + } + auto* rhs_store = builder.Store(result_var, rhs.Get()); + current_flow_block->instructions.Push(rhs_store); + + BranchTo(if_node->merge.target); + } + current_flow_block = if_node->merge.target->As(); + + return result_var; + } + + utils::Result EmitBinary(const ast::BinaryExpression* expr) { + if (expr->op == ast::BinaryOp::kLogicalAnd || expr->op == ast::BinaryOp::kLogicalOr) { + return EmitShortCircuit(expr); + } + + auto lhs = EmitExpression(expr->lhs); + if (!lhs) { + return utils::Failure; + } + + auto rhs = EmitExpression(expr->rhs); + if (!rhs) { + return utils::Failure; + } + + auto* sem = program_->Sem().Get(expr); + auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); + + Binary* inst = nullptr; + switch (expr->op) { + case ast::BinaryOp::kAnd: + inst = builder.And(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kOr: + inst = builder.Or(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kXor: + inst = builder.Xor(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kEqual: + inst = builder.Equal(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kNotEqual: + inst = builder.NotEqual(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kLessThan: + inst = builder.LessThan(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kGreaterThan: + inst = builder.GreaterThan(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kLessThanEqual: + inst = builder.LessThanEqual(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kGreaterThanEqual: + inst = builder.GreaterThanEqual(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kShiftLeft: + inst = builder.ShiftLeft(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kShiftRight: + inst = builder.ShiftRight(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kAdd: + inst = builder.Add(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kSubtract: + inst = builder.Subtract(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kMultiply: + inst = builder.Multiply(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kDivide: + inst = builder.Divide(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kModulo: + inst = builder.Modulo(ty, lhs.Get(), rhs.Get()); + break; + case ast::BinaryOp::kLogicalAnd: + case ast::BinaryOp::kLogicalOr: + TINT_ICE(IR, diagnostics_) << "short circuit op should have already been handled"; + return utils::Failure; + case ast::BinaryOp::kNone: + TINT_ICE(IR, diagnostics_) << "missing binary operand type"; + return utils::Failure; + } + + current_flow_block->instructions.Push(inst); + return inst; + } + + utils::Result EmitBitcast(const ast::BitcastExpression* expr) { + auto val = EmitExpression(expr->expr); + if (!val) { + return utils::Failure; + } + + auto* sem = program_->Sem().Get(expr); + auto* ty = sem->Type()->Clone(clone_ctx_.type_ctx); + auto* inst = builder.Bitcast(ty, val.Get()); + + current_flow_block->instructions.Push(inst); + return inst; + } + + void EmitCall(const ast::CallStatement* stmt) { (void)EmitCall(stmt->expr); } + + utils::Result EmitCall(const ast::CallExpression* expr) { + // If this is a materialized semantic node, just use the constant value. + if (auto* mat = program_->Sem().Get(expr)) { + if (mat->ConstantValue()) { + auto* cv = mat->ConstantValue()->Clone(clone_ctx_); + if (!cv) { + add_error(expr->source, "failed to get constant value for call " + + std::string(expr->TypeInfo().name)); + return utils::Failure; + } + return builder.Constant(cv); + } + } + + utils::Vector args; + args.Reserve(expr->args.Length()); + + // Emit the arguments + for (const auto* arg : expr->args) { + auto value = EmitExpression(arg); + if (!value) { + add_error(arg->source, "failed to convert arguments"); + return utils::Failure; + } + args.Push(value.Get()); + } + + auto* sem = program_->Sem().Get(expr); + if (!sem) { + add_error(expr->source, "failed to get semantic information for call " + + std::string(expr->TypeInfo().name)); + return utils::Failure; + } + + auto* ty = sem->Target()->ReturnType()->Clone(clone_ctx_.type_ctx); + + Instruction* inst = nullptr; + + // If this is a builtin function, emit the specific builtin value + if (auto* b = sem->Target()->As()) { + inst = builder.Builtin(ty, b->Type(), args); + } else if (sem->Target()->As()) { + inst = builder.Construct(ty, std::move(args)); + } else if (auto* conv = sem->Target()->As()) { + auto* from = conv->Source()->Clone(clone_ctx_.type_ctx); + inst = builder.Convert(ty, from, std::move(args)); + } else if (expr->target->identifier->Is()) { + TINT_UNIMPLEMENTED(IR, diagnostics_) << "missing templated ident support"; + return utils::Failure; + } else { + // Not a builtin and not a templated call, so this is a user function. + auto name = CloneSymbol(expr->target->identifier->symbol); + inst = builder.UserCall(ty, name, std::move(args)); + } + if (inst == nullptr) { + return utils::Failure; + } + current_flow_block->instructions.Push(inst); + return inst; + } + + utils::Result EmitLiteral(const ast::LiteralExpression* lit) { + auto* sem = program_->Sem().Get(lit); + if (!sem) { + add_error(lit->source, "failed to get semantic information for node " + + std::string(lit->TypeInfo().name)); + return utils::Failure; + } + + auto* cv = sem->ConstantValue()->Clone(clone_ctx_); + if (!cv) { + add_error(lit->source, + "failed to get constant value for node " + std::string(lit->TypeInfo().name)); + return utils::Failure; + } + return builder.Constant(cv); + } + + // void EmitAttributes(utils::VectorRef attrs) { + // for (auto* attr : attrs) { + // EmitAttribute(attr); + // } + // } + // + // void EmitAttribute(const ast::Attribute* attr) { + // tint::Switch( // + // attr, + // [&](const ast::WorkgroupAttribute* wg) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::StageAttribute* s) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::BindingAttribute* b) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::GroupAttribute* g) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::LocationAttribute* l) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::BuiltinAttribute* b) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::InterpolateAttribute* i) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::InvariantAttribute* i) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::MustUseAttribute* i) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::IdAttribute*) { + // add_error(attr->source, + // "found an `Id` attribute. The SubstituteOverrides transform " + // "must be run before converting to IR"); + // }, + // [&](const ast::StructMemberSizeAttribute*) { + // TINT_ICE(IR, diagnostics_) + // << "StructMemberSizeAttribute encountered during IR conversion"; + // }, + // [&](const ast::StructMemberAlignAttribute*) { + // TINT_ICE(IR, diagnostics_) + // << "StructMemberAlignAttribute encountered during IR conversion"; + // }, + // [&](const ast::StrideAttribute* s) { + // // TODO(dsinclair): Implement + // }, + // [&](const ast::InternalAttribute *i) { + // // TODO(dsinclair): Implement + // }, + // [&](Default) { + // add_error(attr->source, "unknown attribute: " + + // std::string(attr->TypeInfo().name)); + // }); + // } +}; + +} // namespace + utils::Result FromProgram(const Program* program) { if (!program->IsValid()) { return std::string("input program is not valid"); } - BuilderImpl b(program); + Impl b(program); auto r = b.Build(); if (!r) { return b.Diagnostics().str(); diff --git a/src/tint/ir/builder_impl_binary_test.cc b/src/tint/ir/from_program_binary_test.cc similarity index 100% rename from src/tint/ir/builder_impl_binary_test.cc rename to src/tint/ir/from_program_binary_test.cc diff --git a/src/tint/ir/builder_impl_call_test.cc b/src/tint/ir/from_program_call_test.cc similarity index 100% rename from src/tint/ir/builder_impl_call_test.cc rename to src/tint/ir/from_program_call_test.cc diff --git a/src/tint/ir/builder_impl_literal_test.cc b/src/tint/ir/from_program_literal_test.cc similarity index 100% rename from src/tint/ir/builder_impl_literal_test.cc rename to src/tint/ir/from_program_literal_test.cc diff --git a/src/tint/ir/builder_impl_materialize_test.cc b/src/tint/ir/from_program_materialize_test.cc similarity index 100% rename from src/tint/ir/builder_impl_materialize_test.cc rename to src/tint/ir/from_program_materialize_test.cc diff --git a/src/tint/ir/builder_impl_store_test.cc b/src/tint/ir/from_program_store_test.cc similarity index 100% rename from src/tint/ir/builder_impl_store_test.cc rename to src/tint/ir/from_program_store_test.cc diff --git a/src/tint/ir/builder_impl_test.cc b/src/tint/ir/from_program_test.cc similarity index 100% rename from src/tint/ir/builder_impl_test.cc rename to src/tint/ir/from_program_test.cc diff --git a/src/tint/ir/builder_impl_unary_test.cc b/src/tint/ir/from_program_unary_test.cc similarity index 100% rename from src/tint/ir/builder_impl_unary_test.cc rename to src/tint/ir/from_program_unary_test.cc diff --git a/src/tint/ir/builder_impl_var_test.cc b/src/tint/ir/from_program_var_test.cc similarity index 100% rename from src/tint/ir/builder_impl_var_test.cc rename to src/tint/ir/from_program_var_test.cc diff --git a/src/tint/ir/store_test.cc b/src/tint/ir/store_test.cc index 8288b9a3cc..902ca955d9 100644 --- a/src/tint/ir/store_test.cc +++ b/src/tint/ir/store_test.cc @@ -24,7 +24,8 @@ using namespace tint::number_suffixes; // NOLINT using IR_InstructionTest = TestHelper; TEST_F(IR_InstructionTest, CreateStore) { - Builder b; + Module mod; + Builder b{mod}; // TODO(dsinclair): This is wrong, but we don't have anything correct to store too at the // moment. @@ -41,7 +42,8 @@ TEST_F(IR_InstructionTest, CreateStore) { } TEST_F(IR_InstructionTest, Store_Usage) { - Builder b; + Module mod; + Builder b{mod}; auto* to = b.Discard(); const auto* inst = b.Store(to, b.Constant(4_i)); diff --git a/src/tint/ir/unary_test.cc b/src/tint/ir/unary_test.cc index 198dc9ed8c..392a75e340 100644 --- a/src/tint/ir/unary_test.cc +++ b/src/tint/ir/unary_test.cc @@ -24,7 +24,8 @@ using namespace tint::number_suffixes; // NOLINT using IR_InstructionTest = TestHelper; TEST_F(IR_InstructionTest, CreateAddressOf) { - Builder b; + Module mod; + Builder b{mod}; // TODO(dsinclair): This would be better as an identifier, but works for now. const auto* inst = b.AddressOf( @@ -44,7 +45,8 @@ TEST_F(IR_InstructionTest, CreateAddressOf) { } TEST_F(IR_InstructionTest, CreateComplement) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Complement(b.ir.types.Get(), b.Constant(4_i)); ASSERT_TRUE(inst->Is()); @@ -57,7 +59,8 @@ TEST_F(IR_InstructionTest, CreateComplement) { } TEST_F(IR_InstructionTest, CreateIndirection) { - Builder b; + Module mod; + Builder b{mod}; // TODO(dsinclair): This would be better as an identifier, but works for now. const auto* inst = b.Indirection(b.ir.types.Get(), b.Constant(4_i)); @@ -72,7 +75,8 @@ TEST_F(IR_InstructionTest, CreateIndirection) { } TEST_F(IR_InstructionTest, CreateNegation) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); ASSERT_TRUE(inst->Is()); @@ -85,7 +89,8 @@ TEST_F(IR_InstructionTest, CreateNegation) { } TEST_F(IR_InstructionTest, Unary_Usage) { - Builder b; + Module mod; + Builder b{mod}; const auto* inst = b.Negation(b.ir.types.Get(), b.Constant(4_i)); EXPECT_EQ(inst->kind, Unary::Kind::kNegation); diff --git a/src/tint/writer/spirv/generator_impl_type_test.cc b/src/tint/writer/spirv/generator_impl_type_test.cc index 46c54e82f3..c86d363b10 100644 --- a/src/tint/writer/spirv/generator_impl_type_test.cc +++ b/src/tint/writer/spirv/generator_impl_type_test.cc @@ -61,7 +61,7 @@ TEST_F(SpvGeneratorImplTest, Type_F16) { } TEST_F(SpvGeneratorImplTest, Type_Vec2i) { - auto* vec = ir.types.Get(ir.types.Get(), 2u); + auto* vec = b.ir.types.Get(b.ir.types.Get(), 2u); auto id = generator_.Type(vec); EXPECT_EQ(id, 1u); EXPECT_EQ(DumpTypes(), @@ -70,7 +70,7 @@ TEST_F(SpvGeneratorImplTest, Type_Vec2i) { } TEST_F(SpvGeneratorImplTest, Type_Vec3u) { - auto* vec = ir.types.Get(ir.types.Get(), 3u); + auto* vec = b.ir.types.Get(b.ir.types.Get(), 3u); auto id = generator_.Type(vec); EXPECT_EQ(id, 1u); EXPECT_EQ(DumpTypes(), @@ -79,7 +79,7 @@ TEST_F(SpvGeneratorImplTest, Type_Vec3u) { } TEST_F(SpvGeneratorImplTest, Type_Vec4f) { - auto* vec = ir.types.Get(ir.types.Get(), 4u); + auto* vec = b.ir.types.Get(b.ir.types.Get(), 4u); auto id = generator_.Type(vec); EXPECT_EQ(id, 1u); EXPECT_EQ(DumpTypes(), @@ -88,7 +88,7 @@ TEST_F(SpvGeneratorImplTest, Type_Vec4f) { } TEST_F(SpvGeneratorImplTest, Type_Vec4h) { - auto* vec = ir.types.Get(ir.types.Get(), 2u); + auto* vec = b.ir.types.Get(b.ir.types.Get(), 2u); auto id = generator_.Type(vec); EXPECT_EQ(id, 1u); EXPECT_EQ(DumpTypes(), @@ -97,7 +97,7 @@ TEST_F(SpvGeneratorImplTest, Type_Vec4h) { } TEST_F(SpvGeneratorImplTest, Type_Vec4Bool) { - auto* vec = ir.types.Get(ir.types.Get(), 4u); + auto* vec = b.ir.types.Get(b.ir.types.Get(), 4u); auto id = generator_.Type(vec); EXPECT_EQ(id, 1u); EXPECT_EQ(DumpTypes(),