[spirv-writer] Emit entrypoint from function decoration.
This CL updates the SPIRV-Writer to emit entry point information based
on the function stage as well as EntryPoint nodes.
Change-Id: I1fa937cbb2159b31516b0189216d679e03f0384d
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/28702
Reviewed-by: David Neto <dneto@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
diff --git a/BUILD.gn b/BUILD.gn
index fffbd2a..d1cec02 100644
--- a/BUILD.gn
+++ b/BUILD.gn
@@ -842,6 +842,7 @@
"src/writer/spirv/builder_discard_test.cc",
"src/writer/spirv/builder_entry_point_test.cc",
"src/writer/spirv/builder_format_conversion_test.cc",
+ "src/writer/spirv/builder_function_decoration_test.cc",
"src/writer/spirv/builder_function_test.cc",
"src/writer/spirv/builder_function_variable_test.cc",
"src/writer/spirv/builder_global_variable_test.cc",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index fdc2fdf..3aa2b7e 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -502,6 +502,7 @@
writer/spirv/builder_discard_test.cc
writer/spirv/builder_entry_point_test.cc
writer/spirv/builder_format_conversion_test.cc
+ writer/spirv/builder_function_decoration_test.cc
writer/spirv/builder_function_test.cc
writer/spirv/builder_function_variable_test.cc
writer/spirv/builder_global_variable_test.cc
diff --git a/src/ast/function.h b/src/ast/function.h
index 5cb75ba..0dd540b 100644
--- a/src/ast/function.h
+++ b/src/ast/function.h
@@ -104,6 +104,11 @@
/// @returns the functions pipeline stage or None if not set
ast::PipelineStage pipeline_stage() const;
+ /// @returns true if this function is an entry point
+ bool IsEntryPoint() const {
+ return pipeline_stage() != ast::PipelineStage::kNone;
+ }
+
/// Adds the given variable to the list of referenced module variables if it
/// is not already included.
/// @param var the module variable to add
diff --git a/src/type_determiner.cc b/src/type_determiner.cc
index a55ec6c..7e62fcf 100644
--- a/src/type_determiner.cc
+++ b/src/type_determiner.cc
@@ -220,7 +220,7 @@
// Walk over the caller to callee information and update functions with which
// entry points call those functions.
for (const auto& func : mod_->functions()) {
- if (func->pipeline_stage() == ast::PipelineStage::kNone) {
+ if (!func->IsEntryPoint()) {
continue;
}
for (const auto& callee : caller_to_callee_[func->name()]) {
diff --git a/src/validator_impl.cc b/src/validator_impl.cc
index 1b759e0..e64d9f1 100644
--- a/src/validator_impl.cc
+++ b/src/validator_impl.cc
@@ -97,7 +97,7 @@
return false;
}
- if (func->pipeline_stage() != ast::PipelineStage::kNone) {
+ if (func->IsEntryPoint()) {
pipeline_count++;
if (!func->return_type()->IsVoid()) {
diff --git a/src/writer/spirv/builder.cc b/src/writer/spirv/builder.cc
index 5bfdd9f..a952234 100644
--- a/src/writer/spirv/builder.cc
+++ b/src/writer/spirv/builder.cc
@@ -347,6 +347,37 @@
return true;
}
+bool Builder::GenerateEntryPoint(ast::Function* func, uint32_t id) {
+ auto stage = pipeline_stage_to_execution_model(func->pipeline_stage());
+ if (stage == SpvExecutionModelMax) {
+ error_ = "Unknown pipeline stage provided";
+ return false;
+ }
+
+ OperandList operands = {Operand::Int(stage), Operand::Int(id),
+ Operand::String(func->name())};
+
+ for (const auto* var : func->referenced_module_variables()) {
+ // For SPIR-V 1.3 we only output Input/output variables. If we update to
+ // SPIR-V 1.4 or later this should be all variables.
+ if (var->storage_class() != ast::StorageClass::kInput &&
+ var->storage_class() != ast::StorageClass::kOutput) {
+ continue;
+ }
+
+ uint32_t var_id;
+ if (!scope_stack_.get(var->name(), &var_id)) {
+ error_ = "unable to find ID for global variable: " + var->name();
+ return false;
+ }
+
+ operands.push_back(Operand::Int(var_id));
+ }
+ push_preamble(spv::Op::OpEntryPoint, operands);
+
+ return true;
+}
+
bool Builder::GenerateExecutionModes(ast::EntryPoint* ep) {
const auto id = id_for_entry_point(ep);
if (id == 0) {
@@ -373,6 +404,25 @@
return true;
}
+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(
+ spv::Op::OpExecutionMode,
+ {Operand::Int(id), Operand::Int(SpvExecutionModeOriginUpperLeft)});
+ } else if (func->pipeline_stage() == ast::PipelineStage::kCompute) {
+ uint32_t x = 0;
+ 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)});
+ }
+
+ return true;
+}
+
uint32_t Builder::GenerateExpression(ast::Expression* expr) {
if (expr->IsArrayAccessor()) {
return GenerateAccessorExpression(expr->AsArrayAccessor());
@@ -456,10 +506,20 @@
}
}
+ if (func->IsEntryPoint()) {
+ if (!GenerateEntryPoint(func, func_id)) {
+ return false;
+ }
+ if (!GenerateExecutionModes(func, func_id)) {
+ return false;
+ }
+ }
+
scope_stack_.pop_scope();
func_name_to_id_[func->name()] = func_id;
func_name_to_func_[func->name()] = func;
+
return true;
}
diff --git a/src/writer/spirv/builder.h b/src/writer/spirv/builder.h
index 3cc5848..42d762e 100644
--- a/src/writer/spirv/builder.h
+++ b/src/writer/spirv/builder.h
@@ -188,10 +188,20 @@
/// @param ep the entry point
/// @returns true if the instruction was generated, false otherwise
bool GenerateEntryPoint(ast::EntryPoint* ep);
+ /// Generates an entry point instruction
+ /// @param func the function
+ /// @param id the id of the function
+ /// @returns true if the instruction was generated, false otherwise
+ bool GenerateEntryPoint(ast::Function* func, uint32_t id);
/// Generates execution modes for an entry point
/// @param ep the entry point
/// @returns false on failure
bool GenerateExecutionModes(ast::EntryPoint* ep);
+ /// Generates execution modes for an entry point
+ /// @param func the function
+ /// @param id the id of the function
+ /// @returns false on failure
+ bool GenerateExecutionModes(ast::Function* func, uint32_t id);
/// Generates an expression
/// @param expr the expression to generate
/// @returns the resulting ID of the exp = {};ression or 0 on error
diff --git a/src/writer/spirv/builder_entry_point_test.cc b/src/writer/spirv/builder_entry_point_test.cc
index 85af7cf..c594edb 100644
--- a/src/writer/spirv/builder_entry_point_test.cc
+++ b/src/writer/spirv/builder_entry_point_test.cc
@@ -36,6 +36,9 @@
namespace spirv {
namespace {
+// TODO(dsinclair): These have all been ported to stage decorations and this
+// whole file can be deleted when we remove EntryPoint.
+
using BuilderTest = testing::Test;
TEST_F(BuilderTest, EntryPoint) {
@@ -233,55 +236,6 @@
)");
}
-TEST_F(BuilderTest, ExecutionModel_Fragment_OriginUpperLeft) {
- ast::type::VoidType void_type;
-
- ast::Function func("frag_main", {}, &void_type);
- ast::EntryPoint ep(ast::PipelineStage::kFragment, "main", "frag_main");
-
- ast::Module mod;
- Builder b(&mod);
- ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
- ASSERT_TRUE(b.GenerateExecutionModes(&ep));
-
- EXPECT_EQ(DumpInstructions(b.preamble()),
- R"(OpExecutionMode %3 OriginUpperLeft
-)");
-}
-
-TEST_F(BuilderTest, ExecutionModel_Compute_LocalSize) {
- ast::type::VoidType void_type;
-
- ast::Function func("main", {}, &void_type);
- ast::EntryPoint ep(ast::PipelineStage::kCompute, "main", "main");
-
- ast::Module mod;
- Builder b(&mod);
- ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
- ASSERT_TRUE(b.GenerateExecutionModes(&ep));
-
- EXPECT_EQ(DumpInstructions(b.preamble()),
- R"(OpExecutionMode %3 LocalSize 1 1 1
-)");
-}
-
-TEST_F(BuilderTest, ExecutionModel_Compute_LocalSize_WithWorkgroup) {
- ast::type::VoidType void_type;
-
- ast::Function func("main", {}, &void_type);
- func.add_decoration(std::make_unique<ast::WorkgroupDecoration>(2u, 4u, 6u));
- ast::EntryPoint ep(ast::PipelineStage::kCompute, "main", "main");
-
- ast::Module mod;
- Builder b(&mod);
- ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
- ASSERT_TRUE(b.GenerateExecutionModes(&ep));
-
- EXPECT_EQ(DumpInstructions(b.preamble()),
- R"(OpExecutionMode %3 LocalSize 2 4 6
-)");
-}
-
} // namespace
} // namespace spirv
} // namespace writer
diff --git a/src/writer/spirv/builder_function_decoration_test.cc b/src/writer/spirv/builder_function_decoration_test.cc
new file mode 100644
index 0000000..63969cc
--- /dev/null
+++ b/src/writer/spirv/builder_function_decoration_test.cc
@@ -0,0 +1,257 @@
+// Copyright 2020 The Tint Authors.
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+// http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+
+#include <string>
+
+#include "gtest/gtest.h"
+#include "spirv/unified1/spirv.h"
+#include "spirv/unified1/spirv.hpp11"
+#include "src/ast/assignment_statement.h"
+#include "src/ast/function.h"
+#include "src/ast/identifier_expression.h"
+#include "src/ast/pipeline_stage.h"
+#include "src/ast/stage_decoration.h"
+#include "src/ast/type/f32_type.h"
+#include "src/ast/type/void_type.h"
+#include "src/ast/variable.h"
+#include "src/ast/workgroup_decoration.h"
+#include "src/context.h"
+#include "src/type_determiner.h"
+#include "src/writer/spirv/builder.h"
+#include "src/writer/spirv/spv_dump.h"
+
+namespace tint {
+namespace writer {
+namespace spirv {
+namespace {
+
+using BuilderTest = testing::Test;
+
+TEST_F(BuilderTest, FunctionDecoration_Stage) {
+ ast::type::VoidType void_type;
+
+ ast::Function func("main", {}, &void_type);
+ func.add_decoration(
+ std::make_unique<ast::StageDecoration>(ast::PipelineStage::kVertex));
+
+ ast::Module mod;
+ Builder b(&mod);
+ ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.preamble()), R"(OpEntryPoint Vertex %3 "main"
+)");
+}
+
+struct FunctionStageData {
+ ast::PipelineStage stage;
+ SpvExecutionModel model;
+};
+inline std::ostream& operator<<(std::ostream& out, FunctionStageData data) {
+ out << data.stage;
+ return out;
+}
+using FunctionDecoration_StageTest = testing::TestWithParam<FunctionStageData>;
+TEST_P(FunctionDecoration_StageTest, Emit) {
+ auto params = GetParam();
+
+ ast::type::VoidType void_type;
+
+ ast::Function func("main", {}, &void_type);
+ func.add_decoration(std::make_unique<ast::StageDecoration>(params.stage));
+
+ ast::Module mod;
+ Builder b(&mod);
+ ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
+
+ auto preamble = b.preamble();
+ ASSERT_TRUE(preamble.size() >= 1u);
+ EXPECT_EQ(preamble[0].opcode(), spv::Op::OpEntryPoint);
+
+ ASSERT_GE(preamble[0].operands().size(), 3u);
+ EXPECT_EQ(preamble[0].operands()[0].to_i(), params.model);
+}
+INSTANTIATE_TEST_SUITE_P(
+ BuilderTest,
+ FunctionDecoration_StageTest,
+ testing::Values(FunctionStageData{ast::PipelineStage::kVertex,
+ SpvExecutionModelVertex},
+ FunctionStageData{ast::PipelineStage::kFragment,
+ SpvExecutionModelFragment},
+ FunctionStageData{ast::PipelineStage::kCompute,
+ SpvExecutionModelGLCompute}));
+
+TEST_F(BuilderTest, FunctionDecoration_Stage_WithUnusedInterfaceIds) {
+ ast::type::F32Type f32;
+ ast::type::VoidType void_type;
+
+ ast::Function func("main", {}, &void_type);
+ func.add_decoration(
+ std::make_unique<ast::StageDecoration>(ast::PipelineStage::kVertex));
+ auto v_in =
+ std::make_unique<ast::Variable>("my_in", ast::StorageClass::kInput, &f32);
+ auto v_out = std::make_unique<ast::Variable>(
+ "my_out", ast::StorageClass::kOutput, &f32);
+ auto v_wg = std::make_unique<ast::Variable>(
+ "my_wg", ast::StorageClass::kWorkgroup, &f32);
+
+ ast::Module mod;
+ Builder b(&mod);
+ EXPECT_TRUE(b.GenerateGlobalVariable(v_in.get())) << b.error();
+ EXPECT_TRUE(b.GenerateGlobalVariable(v_out.get())) << b.error();
+ EXPECT_TRUE(b.GenerateGlobalVariable(v_wg.get())) << b.error();
+
+ mod.AddGlobalVariable(std::move(v_in));
+ mod.AddGlobalVariable(std::move(v_out));
+ mod.AddGlobalVariable(std::move(v_wg));
+
+ ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %1 "my_in"
+OpName %4 "my_out"
+OpName %7 "my_wg"
+OpName %11 "main"
+)");
+ EXPECT_EQ(DumpInstructions(b.types()), R"(%3 = OpTypeFloat 32
+%2 = OpTypePointer Input %3
+%1 = OpVariable %2 Input
+%5 = OpTypePointer Output %3
+%6 = OpConstantNull %3
+%4 = OpVariable %5 Output %6
+%8 = OpTypePointer Workgroup %3
+%7 = OpVariable %8 Workgroup
+%10 = OpTypeVoid
+%9 = OpTypeFunction %10
+)");
+ EXPECT_EQ(DumpInstructions(b.preamble()),
+ R"(OpEntryPoint Vertex %11 "main"
+)");
+}
+
+TEST_F(BuilderTest, FunctionDecoration_Stage_WithUsedInterfaceIds) {
+ ast::type::F32Type f32;
+ ast::type::VoidType void_type;
+
+ ast::Function func("main", {}, &void_type);
+ func.add_decoration(
+ std::make_unique<ast::StageDecoration>(ast::PipelineStage::kVertex));
+
+ auto body = std::make_unique<ast::BlockStatement>();
+ body->append(std::make_unique<ast::AssignmentStatement>(
+ std::make_unique<ast::IdentifierExpression>("my_out"),
+ std::make_unique<ast::IdentifierExpression>("my_in")));
+ body->append(std::make_unique<ast::AssignmentStatement>(
+ std::make_unique<ast::IdentifierExpression>("my_wg"),
+ std::make_unique<ast::IdentifierExpression>("my_wg")));
+ // Add duplicate usages so we show they don't get output multiple times.
+ body->append(std::make_unique<ast::AssignmentStatement>(
+ std::make_unique<ast::IdentifierExpression>("my_out"),
+ std::make_unique<ast::IdentifierExpression>("my_in")));
+ func.set_body(std::move(body));
+
+ auto v_in =
+ std::make_unique<ast::Variable>("my_in", ast::StorageClass::kInput, &f32);
+ auto v_out = std::make_unique<ast::Variable>(
+ "my_out", ast::StorageClass::kOutput, &f32);
+ auto v_wg = std::make_unique<ast::Variable>(
+ "my_wg", ast::StorageClass::kWorkgroup, &f32);
+
+ Context ctx;
+ ast::Module mod;
+ TypeDeterminer td(&ctx, &mod);
+ td.RegisterVariableForTesting(v_in.get());
+ td.RegisterVariableForTesting(v_out.get());
+ td.RegisterVariableForTesting(v_wg.get());
+
+ ASSERT_TRUE(td.DetermineFunction(&func)) << td.error();
+
+ Builder b(&mod);
+
+ EXPECT_TRUE(b.GenerateGlobalVariable(v_in.get())) << b.error();
+ EXPECT_TRUE(b.GenerateGlobalVariable(v_out.get())) << b.error();
+ EXPECT_TRUE(b.GenerateGlobalVariable(v_wg.get())) << b.error();
+
+ mod.AddGlobalVariable(std::move(v_in));
+ mod.AddGlobalVariable(std::move(v_out));
+ mod.AddGlobalVariable(std::move(v_wg));
+
+ ASSERT_TRUE(b.GenerateFunction(&func)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.debug()), R"(OpName %1 "my_in"
+OpName %4 "my_out"
+OpName %7 "my_wg"
+OpName %11 "main"
+)");
+ EXPECT_EQ(DumpInstructions(b.types()), R"(%3 = OpTypeFloat 32
+%2 = OpTypePointer Input %3
+%1 = OpVariable %2 Input
+%5 = OpTypePointer Output %3
+%6 = OpConstantNull %3
+%4 = OpVariable %5 Output %6
+%8 = OpTypePointer Workgroup %3
+%7 = OpVariable %8 Workgroup
+%10 = OpTypeVoid
+%9 = OpTypeFunction %10
+)");
+ EXPECT_EQ(DumpInstructions(b.preamble()),
+ R"(OpEntryPoint Vertex %11 "main" %4 %1
+)");
+}
+
+TEST_F(BuilderTest, FunctionDecoration_ExecutionMode_Fragment_OriginUpperLeft) {
+ ast::type::VoidType void_type;
+
+ ast::Function func("main", {}, &void_type);
+ func.add_decoration(
+ std::make_unique<ast::StageDecoration>(ast::PipelineStage::kFragment));
+
+ ast::Module mod;
+ Builder b(&mod);
+ ASSERT_TRUE(b.GenerateExecutionModes(&func, 3)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.preamble()),
+ R"(OpExecutionMode %3 OriginUpperLeft
+)");
+}
+
+TEST_F(BuilderTest, FunctionDecoration_WorkgroupSize_Default) {
+ ast::type::VoidType void_type;
+
+ ast::Function func("main", {}, &void_type);
+ func.add_decoration(
+ std::make_unique<ast::StageDecoration>(ast::PipelineStage::kCompute));
+
+ ast::Module mod;
+ Builder b(&mod);
+ ASSERT_TRUE(b.GenerateExecutionModes(&func, 3)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.preamble()),
+ R"(OpExecutionMode %3 LocalSize 1 1 1
+)");
+}
+
+TEST_F(BuilderTest, FunctionDecoration_WorkgroupSize) {
+ ast::type::VoidType void_type;
+
+ ast::Function func("main", {}, &void_type);
+ func.add_decoration(std::make_unique<ast::WorkgroupDecoration>(2u, 4u, 6u));
+ func.add_decoration(
+ std::make_unique<ast::StageDecoration>(ast::PipelineStage::kCompute));
+
+ ast::Module mod;
+ Builder b(&mod);
+ ASSERT_TRUE(b.GenerateExecutionModes(&func, 3)) << b.error();
+ EXPECT_EQ(DumpInstructions(b.preamble()),
+ R"(OpExecutionMode %3 LocalSize 2 4 6
+)");
+}
+
+} // namespace
+} // namespace spirv
+} // namespace writer
+} // namespace tint
diff --git a/test/function.wgsl b/test/function.wgsl
index d1c92d0..f2b2e22 100644
--- a/test/function.wgsl
+++ b/test/function.wgsl
@@ -21,4 +21,3 @@
fn ep() -> void {
return;
}
-entry_point compute = ep;
