[spirv-writer] Fix emission of OpExecutionMode

All of the OpEntryPoint declarations must come before OpExecutionMode.
Currently if you have multiple fragment shaders we'll interleave the
OpEntryPoint and OpExeutionMode which will fail to validate.

Bug: tint:263
Change-Id: I7c925cf6b5345c03bfaf1aa15115caa1bdb9af4c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/29522
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: David Neto <dneto@google.com>

src/writer/spirv/builder.cc

@@ -311,6 +311,8 @@ uint32_t Builder::total_size() const {
 
   size += size_of(capabilities_);
   size += size_of(preamble_);
+  size += size_of(entry_points_);
+  size += size_of(execution_modes_);
   size += size_of(debug_);
   size += size_of(annotations_);
   size += size_of(types_);
@@ -328,6 +330,12 @@ void Builder::iterate(std::function<void(const Instruction&)> cb) const {
   for (const auto& inst : preamble_) {
     cb(inst);
   }
+  for (const auto& inst : entry_points_) {
+    cb(inst);
+  }
+  for (const auto& inst : execution_modes_) {
+    cb(inst);
+  }
   for (const auto& inst : debug_) {
     cb(inst);
   }
@@ -427,7 +435,7 @@ bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) {
 
     operands.push_back(Operand::Int(var_id));
   }
-  push_preamble(spv::Op::OpEntryPoint, operands);
+  push_entry_point(spv::Op::OpEntryPoint, operands);
 
   return true;
 }
@@ -435,7 +443,7 @@ bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) {
 bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
   // WGSL fragment shader origin is upper left
   if (func->pipeline_stage() == ast::PipelineStage::kFragment) {
-    push_preamble(
+    push_execution_mode(
         spv::Op::OpExecutionMode,
         {Operand::Int(id), Operand::Int(SpvExecutionModeOriginUpperLeft)});
   } else if (func->pipeline_stage() == ast::PipelineStage::kCompute) {
@@ -443,9 +451,10 @@ bool Builder::GenerateExecutionModes(ast::Function* func, uint32_t id) {
     uint32_t y = 0;
     uint32_t z = 0;
     std::tie(x, y, z) = func->workgroup_size();
-    push_preamble(spv::Op::OpExecutionMode,
-                  {Operand::Int(id), Operand::Int(SpvExecutionModeLocalSize),
-                   Operand::Int(x), Operand::Int(y), Operand::Int(z)});
+    push_execution_mode(
+        spv::Op::OpExecutionMode,
+        {Operand::Int(id), Operand::Int(SpvExecutionModeLocalSize),
+         Operand::Int(x), Operand::Int(y), Operand::Int(z)});
   }
 
   return true;

src/writer/spirv/builder.h

@@ -103,6 +103,22 @@ class Builder {
   }
   /// @returns the preamble
   const InstructionList& preamble() const { return preamble_; }
+  /// Adds an instruction to the entry points
+  /// @param op the op to set
+  /// @param operands the operands for the instruction
+  void push_entry_point(spv::Op op, const OperandList& operands) {
+    entry_points_.push_back(Instruction{op, operands});
+  }
+  /// @returns the entry points
+  const InstructionList& entry_points() const { return entry_points_; }
+  /// Adds an instruction to the execution modes
+  /// @param op the op to set
+  /// @param operands the operands for the instruction
+  void push_execution_mode(spv::Op op, const OperandList& operands) {
+    execution_modes_.push_back(Instruction{op, operands});
+  }
+  /// @returns the execution modes
+  const InstructionList& execution_modes() const { return execution_modes_; }
   /// Adds an instruction to the debug
   /// @param op the op to set
   /// @param operands the operands for the instruction
@@ -426,6 +442,8 @@ class Builder {
   uint32_t current_label_id_ = 0;
   InstructionList capabilities_;
   InstructionList preamble_;
+  InstructionList entry_points_;
+  InstructionList execution_modes_;
   InstructionList debug_;
   InstructionList types_;
   InstructionList annotations_;

src/writer/spirv/builder_function_decoration_test.cc

@@ -48,7 +48,7 @@ TEST_F(BuilderTest, FunctionDecoration_Stage) {
   ast::Module mod;
   Builder b(&mod);
   ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
-  EXPECT_EQ(DumpInstructions(b.preamble()), R"(OpEntryPoint Vertex %3 "main"
+  EXPECT_EQ(DumpInstructions(b.entry_points()), R"(OpEntryPoint Vertex %3 "main"
 )");
 }
 
@@ -73,7 +73,7 @@ TEST_P(FunctionDecoration_StageTest, Emit) {
   Builder b(&mod);
   ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
 
-  auto preamble = b.preamble();
+  auto preamble = b.entry_points();
   ASSERT_TRUE(preamble.size() >= 1u);
   EXPECT_EQ(preamble[0].opcode(), spv::Op::OpEntryPoint);
 
@@ -131,7 +131,7 @@ OpName %11 "main"
 %10 = OpTypeVoid
 %9 = OpTypeFunction %10
 )");
-  EXPECT_EQ(DumpInstructions(b.preamble()),
+  EXPECT_EQ(DumpInstructions(b.entry_points()),
             R"(OpEntryPoint Vertex %11 "main"
 )");
 }
@@ -200,7 +200,7 @@ OpName %11 "main"
 %10 = OpTypeVoid
 %9 = OpTypeFunction %10
 )");
-  EXPECT_EQ(DumpInstructions(b.preamble()),
+  EXPECT_EQ(DumpInstructions(b.entry_points()),
             R"(OpEntryPoint Vertex %11 "main" %4 %1
 )");
 }
@@ -215,7 +215,7 @@ TEST_F(BuilderTest, FunctionDecoration_ExecutionMode_Fragment_OriginUpperLeft) {
   ast::Module mod;
   Builder b(&mod);
   ASSERT_TRUE(b.GenerateExecutionModes(&func, 3)) << b.error();
-  EXPECT_EQ(DumpInstructions(b.preamble()),
+  EXPECT_EQ(DumpInstructions(b.execution_modes()),
             R"(OpExecutionMode %3 OriginUpperLeft
 )");
 }
@@ -230,7 +230,7 @@ TEST_F(BuilderTest, FunctionDecoration_WorkgroupSize_Default) {
   ast::Module mod;
   Builder b(&mod);
   ASSERT_TRUE(b.GenerateExecutionModes(&func, 3)) << b.error();
-  EXPECT_EQ(DumpInstructions(b.preamble()),
+  EXPECT_EQ(DumpInstructions(b.execution_modes()),
             R"(OpExecutionMode %3 LocalSize 1 1 1
 )");
 }
@@ -246,11 +246,44 @@ TEST_F(BuilderTest, FunctionDecoration_WorkgroupSize) {
   ast::Module mod;
   Builder b(&mod);
   ASSERT_TRUE(b.GenerateExecutionModes(&func, 3)) << b.error();
-  EXPECT_EQ(DumpInstructions(b.preamble()),
+  EXPECT_EQ(DumpInstructions(b.execution_modes()),
             R"(OpExecutionMode %3 LocalSize 2 4 6
 )");
 }
 
+TEST_F(BuilderTest, FunctionDecoration_ExecutionMode_MultipleFragment) {
+  ast::type::VoidType void_type;
+
+  ast::Function func1("main1", {}, &void_type);
+  func1.add_decoration(
+      std::make_unique<ast::StageDecoration>(ast::PipelineStage::kFragment));
+
+  ast::Function func2("main2", {}, &void_type);
+  func2.add_decoration(
+      std::make_unique<ast::StageDecoration>(ast::PipelineStage::kFragment));
+
+  ast::Module mod;
+  Builder b(&mod);
+  ASSERT_TRUE(b.GenerateFunction(&func1)) << b.error();
+  ASSERT_TRUE(b.GenerateFunction(&func2)) << b.error();
+  EXPECT_EQ(DumpBuilder(b),
+            R"(OpEntryPoint Fragment %3 "main1"
+OpEntryPoint Fragment %5 "main2"
+OpExecutionMode %3 OriginUpperLeft
+OpExecutionMode %5 OriginUpperLeft
+OpName %3 "main1"
+OpName %5 "main2"
+%2 = OpTypeVoid
+%1 = OpTypeFunction %2
+%3 = OpFunction %2 None %1
+%4 = OpLabel
+OpFunctionEnd
+%5 = OpFunction %2 None %1
+%6 = OpLabel
+OpFunctionEnd
+)");
+}
+
 }  // namespace
 }  // namespace spirv
 }  // namespace writer