From a8274b2fef8038cd8580e312aaeb407504a8942e Mon Sep 17 00:00:00 2001 From: dan sinclair Date: Mon, 21 Sep 2020 18:49:01 +0000 Subject: [PATCH] [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 Commit-Queue: dan sinclair --- BUILD.gn | 1 + src/CMakeLists.txt | 1 + src/ast/function.h | 5 + src/type_determiner.cc | 2 +- src/validator_impl.cc | 2 +- src/writer/spirv/builder.cc | 60 ++++ src/writer/spirv/builder.h | 10 + src/writer/spirv/builder_entry_point_test.cc | 52 +--- .../spirv/builder_function_decoration_test.cc | 257 ++++++++++++++++++ test/function.wgsl | 1 - 10 files changed, 339 insertions(+), 52 deletions(-) create mode 100644 src/writer/spirv/builder_function_decoration_test.cc diff --git a/BUILD.gn b/BUILD.gn index fffbd2a661..d1cec02c21 100644 --- a/BUILD.gn +++ b/BUILD.gn @@ -842,6 +842,7 @@ source_set("tint_unittests_spv_writer_src") { "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 fdc2fdfd3e..3aa2b7ee99 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -502,6 +502,7 @@ if(${TINT_BUILD_SPV_WRITER}) 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 5cb75bafe7..0dd540b020 100644 --- a/src/ast/function.h +++ b/src/ast/function.h @@ -104,6 +104,11 @@ class Function : public Node { /// @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 a55ec6c3ec..7e62fcf308 100644 --- a/src/type_determiner.cc +++ b/src/type_determiner.cc @@ -220,7 +220,7 @@ bool TypeDeterminer::Determine() { // 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 1b759e0fc4..e64d9f1c41 100644 --- a/src/validator_impl.cc +++ b/src/validator_impl.cc @@ -97,7 +97,7 @@ bool ValidatorImpl::ValidateFunctions(const ast::Module* mod, 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 5bfdd9f38d..a952234956 100644 --- a/src/writer/spirv/builder.cc +++ b/src/writer/spirv/builder.cc @@ -347,6 +347,37 @@ bool Builder::GenerateEntryPoint(ast::EntryPoint* ep) { 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 @@ bool Builder::GenerateExecutionModes(ast::EntryPoint* ep) { 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 @@ bool Builder::GenerateFunction(ast::Function* func) { } } + 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 3cc5848d5c..42d762eb5a 100644 --- a/src/writer/spirv/builder.h +++ b/src/writer/spirv/builder.h @@ -188,10 +188,20 @@ class Builder { /// @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 85af7cfbb2..c594edbb06 100644 --- a/src/writer/spirv/builder_entry_point_test.cc +++ b/src/writer/spirv/builder_entry_point_test.cc @@ -36,6 +36,9 @@ namespace writer { 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 @@ OpName %11 "main" )"); } -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(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 0000000000..63969cc471 --- /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 + +#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::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; +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(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::PipelineStage::kVertex)); + auto v_in = + std::make_unique("my_in", ast::StorageClass::kInput, &f32); + auto v_out = std::make_unique( + "my_out", ast::StorageClass::kOutput, &f32); + auto v_wg = std::make_unique( + "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::PipelineStage::kVertex)); + + auto body = std::make_unique(); + body->append(std::make_unique( + std::make_unique("my_out"), + std::make_unique("my_in"))); + body->append(std::make_unique( + std::make_unique("my_wg"), + std::make_unique("my_wg"))); + // Add duplicate usages so we show they don't get output multiple times. + body->append(std::make_unique( + std::make_unique("my_out"), + std::make_unique("my_in"))); + func.set_body(std::move(body)); + + auto v_in = + std::make_unique("my_in", ast::StorageClass::kInput, &f32); + auto v_out = std::make_unique( + "my_out", ast::StorageClass::kOutput, &f32); + auto v_wg = std::make_unique( + "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::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::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(2u, 4u, 6u)); + func.add_decoration( + std::make_unique(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 d1c92d0fc4..f2b2e2281d 100644 --- a/test/function.wgsl +++ b/test/function.wgsl @@ -21,4 +21,3 @@ fn main() -> f32 { fn ep() -> void { return; } -entry_point compute = ep;