[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>
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 2b8ee2c..addfdd1 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -311,6 +311,8 @@
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 @@
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 @@
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::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 @@
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;
diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h
index b6e3f7c..f4b3388 100644
--- a/src/writer/spirv/builder.h
+++ b/src/writer/spirv/builder.h
@@ -103,6 +103,22 @@
}
/// @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 @@
uint32_t current_label_id_ = 0;
InstructionList capabilities_;
InstructionList preamble_;
+ InstructionList entry_points_;
+ InstructionList execution_modes_;
InstructionList debug_;
InstructionList types_;
InstructionList annotations_;
diff --git a/src/writer/spirv/builder_function_decoration_test.cc b/src/writer/spirv/builder_function_decoration_test.cc
index 63969cc..26cfcf2 100644
--- a/src/writer/spirv/builder_function_decoration_test.cc
+++ b/src/writer/spirv/builder_function_decoration_test.cc
@@ -48,7 +48,7 @@
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 @@
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 @@
%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 @@
%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 @@
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 @@
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 @@
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
