ir/spirv-writer: Emit entry point declarations

Emit the OpEntryPoint instruction with the pipeline stage. Interface
variables will be done later.

Emit OpExecutionMode instructions for the workgroup size and fragment
shader origin, depending on the pipeline stage.

Bug: tint:1906
Change-Id: Ieeeda5f17da48a8cf0d3344d3b254542c7198cb9
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/131381
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>

src/tint/writer/spirv/generator_impl_function_test.cc

@@ -46,5 +46,109 @@ TEST_F(SpvGeneratorImplTest, Function_DeduplicateType) {
 )");
 }
 
+TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Compute) {
+    auto* func = CreateFunction();
+    func->name = ir.symbols.Register("main");
+    func->return_type = ir.types.Get<type::Void>();
+    func->pipeline_stage = ir::Function::PipelineStage::kCompute;
+    func->workgroup_size = {32, 4, 1};
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main"
+OpExecutionMode %1 LocalSize 32 4 1
+OpName %1 "main"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Fragment) {
+    auto* func = CreateFunction();
+    func->name = ir.symbols.Register("main");
+    func->return_type = ir.types.Get<type::Void>();
+    func->pipeline_stage = ir::Function::PipelineStage::kFragment;
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Fragment %1 "main"
+OpExecutionMode %1 OriginUpperLeft
+OpName %1 "main"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Vertex) {
+    auto* func = CreateFunction();
+    func->name = ir.symbols.Register("main");
+    func->return_type = ir.types.Get<type::Void>();
+    func->pipeline_stage = ir::Function::PipelineStage::kVertex;
+
+    generator_.EmitFunction(func);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint Vertex %1 "main"
+OpName %1 "main"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
+TEST_F(SpvGeneratorImplTest, Function_EntryPoint_Multiple) {
+    auto* f1 = CreateFunction();
+    f1->name = ir.symbols.Register("main1");
+    f1->return_type = ir.types.Get<type::Void>();
+    f1->pipeline_stage = ir::Function::PipelineStage::kCompute;
+    f1->workgroup_size = {32, 4, 1};
+
+    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};
+
+    auto* f3 = CreateFunction();
+    f3->name = ir.symbols.Register("main3");
+    f3->return_type = ir.types.Get<type::Void>();
+    f3->pipeline_stage = ir::Function::PipelineStage::kFragment;
+
+    generator_.EmitFunction(f1);
+    generator_.EmitFunction(f2);
+    generator_.EmitFunction(f3);
+    EXPECT_EQ(DumpModule(generator_.Module()), R"(OpEntryPoint GLCompute %1 "main1"
+OpEntryPoint GLCompute %5 "main2"
+OpEntryPoint Fragment %7 "main3"
+OpExecutionMode %1 LocalSize 32 4 1
+OpExecutionMode %5 LocalSize 8 2 16
+OpExecutionMode %7 OriginUpperLeft
+OpName %1 "main1"
+OpName %5 "main2"
+OpName %7 "main3"
+%2 = OpTypeVoid
+%3 = OpTypeFunction %2
+%1 = OpFunction %2 None %3
+%4 = OpLabel
+OpReturn
+OpFunctionEnd
+%5 = OpFunction %2 None %3
+%6 = OpLabel
+OpReturn
+OpFunctionEnd
+%7 = OpFunction %2 None %3
+%8 = OpLabel
+OpReturn
+OpFunctionEnd
+)");
+}
+
 }  // namespace
 }  // namespace tint::writer::spirv

src/tint/writer/spirv/generator_impl_ir.cc

@@ -88,7 +88,10 @@ void GeneratorImplIr::EmitFunction(const ir::Function* func) {
     // Emit the function name.
     module_.PushDebug(spv::Op::OpName, {id, Operand(func->name.Name())});
 
-    // TODO(jrprice): Emit OpEntryPoint and OpExecutionMode declarations if needed.
+    // Emit OpEntryPoint and OpExecutionMode declarations if needed.
+    if (func->pipeline_stage != ir::Function::PipelineStage::kUndefined) {
+        EmitEntryPoint(func, id);
+    }
 
     // Get the ID for the return type.
     auto return_type_id = Type(func->return_type);
@@ -123,4 +126,34 @@ void GeneratorImplIr::EmitFunction(const ir::Function* func) {
     module_.PushFunction(current_function_);
 }
 
+void GeneratorImplIr::EmitEntryPoint(const ir::Function* func, uint32_t id) {
+    SpvExecutionModel stage;
+    switch (func->pipeline_stage) {
+        case ir::Function::PipelineStage::kCompute: {
+            stage = SpvExecutionModelGLCompute;
+            module_.PushExecutionMode(spv::Op::OpExecutionMode,
+                                      {id, SpvExecutionModeLocalSize, func->workgroup_size->at(0),
+                                       func->workgroup_size->at(1), func->workgroup_size->at(2)});
+            break;
+        }
+        case ir::Function::PipelineStage::kFragment: {
+            stage = SpvExecutionModelFragment;
+            module_.PushExecutionMode(spv::Op::OpExecutionMode,
+                                      {id, SpvExecutionModeOriginUpperLeft});
+            // TODO(jrprice): Add DepthReplacing execution mode if FragDepth is used.
+            break;
+        }
+        case ir::Function::PipelineStage::kVertex: {
+            stage = SpvExecutionModelVertex;
+            break;
+        }
+        case ir::Function::PipelineStage::kUndefined:
+            TINT_ICE(Writer, diagnostics_) << "undefined pipeline stage for entry point";
+            return;
+    }
+
+    // TODO(jrprice): Add the interface list of all referenced global variables.
+    module_.PushEntryPoint(spv::Op::OpEntryPoint, {stage, id, func->name.Name()});
+}
+
 }  // namespace tint::writer::spirv

src/tint/writer/spirv/generator_impl_ir.h

@@ -64,6 +64,11 @@ class GeneratorImplIr {
     /// @param func the function to emit
     void EmitFunction(const ir::Function* func);
 
+    /// Emit entry point declarations for a function.
+    /// @param func the function to emit entry point declarations for
+    /// @param id the result ID of the function declaration
+    void EmitEntryPoint(const ir::Function* func, uint32_t id);
+
   private:
     const ir::Module* ir_;
     spirv::Module module_;