[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>

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",

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

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

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()]) {

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()) {

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;
 }
 

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

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<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

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

test/function.wgsl

@@ -21,4 +21,3 @@ fn main() -> f32 {
 fn ep() -> void {
   return;
 }
-entry_point compute = ep;