[spirv-reader] Handle entry point declarations for GLCompute
Emit @compute and @workgroup_size attributes, and set the entry point
function name.
Bug: tint:1907
Change-Id: I800ac767b9b1798cf3cc85f67cc3d6bc07edb785
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/165560
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/tint/lang/spirv/reader/parser/function_test.cc b/src/tint/lang/spirv/reader/parser/function_test.cc
index be6c28c..66e7ae3 100644
--- a/src/tint/lang/spirv/reader/parser/function_test.cc
+++ b/src/tint/lang/spirv/reader/parser/function_test.cc
@@ -44,7 +44,7 @@
OpFunctionEnd
)",
R"(
-%1 = func():void -> %b1 {
+%main = @compute @workgroup_size(1, 1, 1) func():void -> %b1 {
%b1 = block {
ret
}
@@ -52,5 +52,62 @@
)");
}
+TEST_F(SpirvParserTest, LocalSize) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 3 4 5
+ %void = OpTypeVoid
+ %ep_type = OpTypeFunction %void
+ %main = OpFunction %void None %ep_type
+ %main_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%main = @compute @workgroup_size(3, 4, 5) func():void -> %b1 {
+ %b1 = block {
+ ret
+ }
+}
+)");
+}
+
+TEST_F(SpirvParserTest, MultipleEntryPoints) {
+ EXPECT_IR(R"(
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %foo "foo"
+ OpEntryPoint GLCompute %bar "bar"
+ OpExecutionMode %foo LocalSize 3 4 5
+ OpExecutionMode %bar LocalSize 6 7 8
+ %void = OpTypeVoid
+ %ep_type = OpTypeFunction %void
+
+ %foo = OpFunction %void None %ep_type
+ %foo_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+
+ %bar = OpFunction %void None %ep_type
+ %bar_start = OpLabel
+ OpReturn
+ OpFunctionEnd
+)",
+ R"(
+%foo = @compute @workgroup_size(3, 4, 5) func():void -> %b1 {
+ %b1 = block {
+ ret
+ }
+}
+%bar = @compute @workgroup_size(6, 7, 8) func():void -> %b2 {
+ %b2 = block {
+ ret
+ }
+}
+)");
+}
+
} // namespace
} // namespace tint::spirv::reader
diff --git a/src/tint/lang/spirv/reader/parser/parser.cc b/src/tint/lang/spirv/reader/parser/parser.cc
index aacba85..1559ab5 100644
--- a/src/tint/lang/spirv/reader/parser/parser.cc
+++ b/src/tint/lang/spirv/reader/parser/parser.cc
@@ -77,8 +77,8 @@
EmitFunctions();
- // TODO(crbug.com/tint/1907): Handle entry point declarations and execution modes.
- // TODO(crbug.com/tint/1907): Handle entry point declarations and execution modes.
+ EmitEntryPoints();
+
// TODO(crbug.com/tint/1907): Handle annotation instructions.
// TODO(crbug.com/tint/1907): Handle names.
@@ -109,10 +109,50 @@
// TODO(crbug.com/tint/1907): Emit function parameters as well.
current_function_ = b_.Function(
Type(func.type_id()), core::ir::Function::PipelineStage::kUndefined, std::nullopt);
+ functions_.Add(func.result_id(), current_function_);
EmitBlock(current_function_->Block(), *func.entry());
}
}
+ /// Emit entry point attributes.
+ void EmitEntryPoints() {
+ // Handle OpEntryPoint declarations.
+ for (auto& entry_point : spirv_context_->module()->entry_points()) {
+ auto model = entry_point.GetSingleWordInOperand(0);
+ auto* func = functions_.Get(entry_point.GetSingleWordInOperand(1)).value_or(nullptr);
+ TINT_ASSERT_OR_RETURN(func);
+
+ // Set the pipeline stage.
+ switch (spv::ExecutionModel(model)) {
+ case spv::ExecutionModel::GLCompute:
+ func->SetStage(core::ir::Function::PipelineStage::kCompute);
+ break;
+ default:
+ TINT_UNIMPLEMENTED() << "unhandled execution model: " << model;
+ }
+
+ // Set the entry point name.
+ ir_.SetName(func, entry_point.GetOperand(2).AsString());
+ }
+
+ // Handle OpExecutionMode declarations.
+ for (auto& execution_mode : spirv_context_->module()->execution_modes()) {
+ auto* func = functions_.Get(execution_mode.GetSingleWordInOperand(0)).value_or(nullptr);
+ auto mode = execution_mode.GetSingleWordInOperand(1);
+ TINT_ASSERT_OR_RETURN(func);
+
+ switch (spv::ExecutionMode(mode)) {
+ case spv::ExecutionMode::LocalSize:
+ func->SetWorkgroupSize(execution_mode.GetSingleWordInOperand(2),
+ execution_mode.GetSingleWordInOperand(3),
+ execution_mode.GetSingleWordInOperand(4));
+ break;
+ default:
+ TINT_UNIMPLEMENTED() << "unhandled execution mode: " << mode;
+ }
+ }
+ }
+
/// Emit the contents of SPIR-V block @p src into Tint IR block @p dst.
/// @param dst the Tint IR block to append to
/// @param src the SPIR-V block to emit
@@ -139,6 +179,8 @@
/// The Tint IR function that is currently being emitted.
core::ir::Function* current_function_ = nullptr;
+ /// A map from a SPIR-V function definition result ID to the corresponding Tint function object.
+ Hashmap<uint32_t, core::ir::Function*, 8> functions_;
/// The SPIR-V context containing the SPIR-V tools intermediate representation.
std::unique_ptr<spvtools::opt::IRContext> spirv_context_;
