[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 <jrprice@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

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<type::Void>();
+    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<type::Void>();
+    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<type::Void>();
     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<type::Void>();
     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<type::Void>();
     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<type::Void>();
     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<type::Void>();
     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<type::Void>();
     f3->pipeline_stage = ir::Function::PipelineStage::kFragment;
+    f3->start_target->branch.target = f3->end_target;
 
     generator_.EmitFunction(f1);
     generator_.EmitFunction(f2);

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

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<const ir::Constant*, uint32_t, 16, ConstantHasher, ConstantEquals> constants_;
 
+    /// The map of instructions to their result IDs.
+    utils::Hashmap<const ir::Instruction*, uint32_t, 8> instructions_;
+
+    /// The current function that is being emitted.
+    Function current_function_;
+
     bool zero_init_workgroup_memory_ = false;
 };