From 5e7807f02bf79fa68910a051a868531e89cc84dc Mon Sep 17 00:00:00 2001 From: James Price Date: Tue, 9 May 2023 22:20:16 +0000 Subject: [PATCH] [ir][spirv-writer] Emit blocks All instructions are unimplemented, and the only supported branch target is the function terminator. No new tests, as this only handles returning from an empty function. Bug: tint:1906 Change-Id: I05cf5d252d96d37757c3ac1f84b79d0586c36b97 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/132283 Commit-Queue: James Price Reviewed-by: Ben Clayton Kokoro: Kokoro --- .../spirv/generator_impl_function_test.cc | 8 +++++ src/tint/writer/spirv/generator_impl_ir.cc | 33 ++++++++++++++++--- src/tint/writer/spirv/generator_impl_ir.h | 12 +++++++ 3 files changed, 48 insertions(+), 5 deletions(-) diff --git a/src/tint/writer/spirv/generator_impl_function_test.cc b/src/tint/writer/spirv/generator_impl_function_test.cc index bb7008ce3b..a1ce3cc86a 100644 --- a/src/tint/writer/spirv/generator_impl_function_test.cc +++ b/src/tint/writer/spirv/generator_impl_function_test.cc @@ -21,6 +21,7 @@ TEST_F(SpvGeneratorImplTest, Function_Empty) { auto* func = CreateFunction(); func->name = ir.symbols.Register("foo"); func->return_type = ir.types.Get(); + func->start_target->branch.target = func->end_target; generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpName %1 "foo" @@ -37,6 +38,7 @@ OpFunctionEnd TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) { auto* func = CreateFunction(); func->return_type = ir.types.Get(); + func->start_target->branch.target = func->end_target; generator_.EmitFunction(func); generator_.EmitFunction(func); @@ -52,6 +54,7 @@ TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Compute) { func->return_type = ir.types.Get(); func->pipeline_stage = ir::Function::PipelineStage::kCompute; func->workgroup_size = {32, 4, 1}; + func->start_target->branch.target = func->end_target; generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main" @@ -71,6 +74,7 @@ TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Fragment) { func->name = ir.symbols.Register("main"); func->return_type = ir.types.Get(); func->pipeline_stage = ir::Function::PipelineStage::kFragment; + func->start_target->branch.target = func->end_target; generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Fragment %1 "main" @@ -90,6 +94,7 @@ TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Vertex) { func->name = ir.symbols.Register("main"); func->return_type = ir.types.Get(); func->pipeline_stage = ir::Function::PipelineStage::kVertex; + func->start_target->branch.target = func->end_target; generator_.EmitFunction(func); EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Vertex %1 "main" @@ -109,17 +114,20 @@ TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Multiple) { f1->return_type = ir.types.Get(); f1->pipeline_stage = ir::Function::PipelineStage::kCompute; f1->workgroup_size = {32, 4, 1}; + f1->start_target->branch.target = f1->end_target; auto* f2 = CreateFunction(); f2->name = ir.symbols.Register("main2"); f2->return_type = ir.types.Get(); f2->pipeline_stage = ir::Function::PipelineStage::kCompute; f2->workgroup_size = {8, 2, 16}; + f2->start_target->branch.target = f2->end_target; auto* f3 = CreateFunction(); f3->name = ir.symbols.Register("main3"); f3->return_type = ir.types.Get(); f3->pipeline_stage = ir::Function::PipelineStage::kFragment; + f3->start_target->branch.target = f3->end_target; generator_.EmitFunction(f1); generator_.EmitFunction(f2); diff --git a/src/tint/writer/spirv/generator_impl_ir.cc b/src/tint/writer/spirv/generator_impl_ir.cc index 76de13a805..b38ecba1b6 100644 --- a/src/tint/writer/spirv/generator_impl_ir.cc +++ b/src/tint/writer/spirv/generator_impl_ir.cc @@ -15,6 +15,8 @@ #include "src/tint/writer/spirv/generator_impl_ir.h" #include "spirv/unified1/spirv.h" +#include "src/tint/ir/block.h" +#include "src/tint/ir/function_terminator.h" #include "src/tint/ir/module.h" #include "src/tint/switch.h" #include "src/tint/type/bool.h" @@ -150,12 +152,11 @@ void GeneratorImplIr::EmitFunction(const ir::Function* func) { // Create a function that we will add instructions to. // TODO(jrprice): Add the parameter declarations when they are supported in the IR. auto entry_block = module_.NextId(); - Function current_function_(decl, entry_block, {}); + current_function_ = Function(decl, entry_block, {}); + TINT_DEFER(current_function_ = Function()); - // TODO(jrprice): Emit the body of the function. - - // TODO(jrprice): Remove this when we start emitting OpReturn for branches to the terminator. - current_function_.push_inst(spv::Op::OpReturn, {}); + // Emit the body of the function. + EmitBlock(func->start_target); // Add the function to the module. module_.PushFunction(current_function_); @@ -192,4 +193,26 @@ void GeneratorImplIr::EmitEntryPoint(const ir::Function* func, uint32_t id) { module_.PushEntryPoint(spv::Op::OpEntryPoint, {U32Operand(stage), id, func->name.Name()}); } +void GeneratorImplIr::EmitBlock(const ir::Block* block) { + // Emit the instructions. + for (auto* inst : block->instructions) { + auto result = Switch(inst, // + [&](Default) { + TINT_ICE(Writer, diagnostics_) + << "unimplemented instruction: " << inst->TypeInfo().name; + return 0u; + }); + instructions_.Add(inst, result); + } + + // Handle the branch at the end of the block. + Switch( + block->branch.target, + [&](const ir::FunctionTerminator*) { + // TODO(jrprice): Handle the return value, which will be a branch argument. + current_function_.push_inst(spv::Op::OpReturn, {}); + }, + [&](Default) { TINT_ICE(Writer, diagnostics_) << "unimplemented branch target"; }); +} + } // namespace tint::writer::spirv diff --git a/src/tint/writer/spirv/generator_impl_ir.h b/src/tint/writer/spirv/generator_impl_ir.h index f7a6898f82..c8965dcd93 100644 --- a/src/tint/writer/spirv/generator_impl_ir.h +++ b/src/tint/writer/spirv/generator_impl_ir.h @@ -23,10 +23,12 @@ #include "src/tint/utils/hashmap.h" #include "src/tint/utils/vector.h" #include "src/tint/writer/spirv/binary_writer.h" +#include "src/tint/writer/spirv/function.h" #include "src/tint/writer/spirv/module.h" // Forward declarations namespace tint::ir { +class Block; class Function; class Module; } // namespace tint::ir @@ -76,6 +78,10 @@ class GeneratorImplIr { /// @param id the result ID of the function declaration void EmitEntryPoint(const ir::Function* func, uint32_t id); + /// Emit a block. + /// @param block the block to emit + void EmitBlock(const ir::Block* block); + private: const ir::Module* ir_; spirv::Module module_; @@ -135,6 +141,12 @@ class GeneratorImplIr { /// The map of constants to their result IDs. utils::Hashmap constants_; + /// The map of instructions to their result IDs. + utils::Hashmap instructions_; + + /// The current function that is being emitted. + Function current_function_; + bool zero_init_workgroup_memory_ = false; };