[ir] Add AddEmptyEntryPoint transform

Use it in the SPIR-V writer.

Bug: tint:1718, tint:1906
Change-Id: If8f29300712c457a02ddc9eb2fd76e0b49ee8fea
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/132682
Reviewed-by: Dan Sinclair <dsinclair@chromium.org>
Commit-Queue: James Price <jrprice@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/tint/BUILD.gn

@@ -505,6 +505,20 @@ libtint_source_set("libtint_ast_transform_src") {
   ]
 }
 
+libtint_source_set("libtint_ir_transform_src") {
+  sources = [
+    "ir/transform/add_empty_entry_point.cc",
+    "ir/transform/add_empty_entry_point.h",
+  ]
+  deps = [
+    ":libtint_builtins_src",
+    ":libtint_ir_src",
+    ":libtint_symbols_src",
+    ":libtint_type_src",
+    ":libtint_utils_src",
+  ]
+}
+
 libtint_source_set("libtint_ast_hdrs") {
   sources = [
     "ast/accessor_expression.h",
@@ -1020,6 +1034,7 @@ libtint_source_set("libtint_spv_writer_src") {
     deps += [
       ":libtint_ir_builder_src",
       ":libtint_ir_src",
+      ":libtint_ir_transform_src",
     ]
   }
 }
@@ -1753,6 +1768,19 @@ if (tint_build_unittests) {
     ]
   }
 
+  tint_unittests_source_set("tint_unittests_ir_transform_src") {
+    sources = [
+      "ir/transform/add_empty_entry_point_test.cc",
+      "ir/transform/test_helper.h",
+    ]
+
+    deps = [
+      ":libtint_ir_src",
+      ":libtint_ir_transform_src",
+      ":libtint_transform_manager_src",
+    ]
+  }
+
   tint_unittests_source_set("tint_unittests_utils_src") {
     sources = [
       "debug_test.cc",
@@ -2296,7 +2324,10 @@ if (tint_build_unittests) {
     }
 
     if (tint_build_ir) {
-      deps += [ ":tint_unittests_ir_src" ]
+      deps += [
+        ":tint_unittests_ir_src",
+        ":tint_unittests_ir_transform_src",
+      ]
     }
 
     if (build_with_chromium) {

src/tint/CMakeLists.txt

@@ -768,6 +768,8 @@ if(${TINT_BUILD_IR})
     ir/value.h
     ir/var.cc
     ir/var.h
+    ir/transform/add_empty_entry_point.cc
+    ir/transform/add_empty_entry_point.h
     ir/transform/transform.cc
     ir/transform/transform.h
   )

src/tint/ir/transform/add_empty_entry_point.cc

@@ -0,0 +1,45 @@
+// Copyright 2023 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 "src/tint/ir/transform/add_empty_entry_point.h"
+
+#include <utility>
+
+#include "src/tint/ir/builder.h"
+#include "src/tint/ir/module.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::ir::transform::AddEmptyEntryPoint);
+
+namespace tint::ir::transform {
+
+AddEmptyEntryPoint::AddEmptyEntryPoint() = default;
+
+AddEmptyEntryPoint::~AddEmptyEntryPoint() = default;
+
+void AddEmptyEntryPoint::Run(ir::Module* ir, const DataMap&, DataMap&) const {
+    for (auto* func : ir->functions) {
+        if (func->pipeline_stage != Function::PipelineStage::kUndefined) {
+            return;
+        }
+    }
+
+    ir::Builder builder(*ir);
+    auto* ep =
+        builder.CreateFunction(ir->symbols.New("unused_entry_point"), ir->types.Get<type::Void>(),
+                               Function::PipelineStage::kCompute, std::array{1u, 1u, 1u});
+    builder.Branch(ep->start_target, ep->end_target);
+    ir->functions.Push(ep);
+}
+
+}  // namespace tint::ir::transform

src/tint/ir/transform/add_empty_entry_point.h

@@ -0,0 +1,36 @@
+// Copyright 2023 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.
+
+#ifndef SRC_TINT_IR_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_
+#define SRC_TINT_IR_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_
+
+#include "src/tint/ir/transform/transform.h"
+
+namespace tint::ir::transform {
+
+/// Add an empty entry point to the module, if no other entry points exist.
+class AddEmptyEntryPoint final : public utils::Castable<AddEmptyEntryPoint, Transform> {
+  public:
+    /// Constructor
+    AddEmptyEntryPoint();
+    /// Destructor
+    ~AddEmptyEntryPoint() override;
+
+    /// @copydoc Transform::Run
+    void Run(ir::Module* module, const DataMap& inputs, DataMap& outputs) const override;
+};
+
+}  // namespace tint::ir::transform
+
+#endif  // SRC_TINT_IR_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_

src/tint/ir/transform/add_empty_entry_point_test.cc

@@ -0,0 +1,60 @@
+// Copyright 2023 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 "src/tint/ir/transform/add_empty_entry_point.h"
+
+#include <utility>
+
+#include "src/tint/ir/transform/test_helper.h"
+
+namespace tint::ir::transform {
+namespace {
+
+using IR_AddEmptyEntryPointTest = TransformTest;
+
+TEST_F(IR_AddEmptyEntryPointTest, EmptyModule) {
+    auto* expect = R"(
+%fn1 = func unused_entry_point():void [@compute @workgroup_size(1, 1, 1)] {
+  %fn2 = block {
+  } -> %func_end # return
+} %func_end
+
+)";
+
+    Run<AddEmptyEntryPoint>();
+
+    EXPECT_EQ(expect, str());
+}
+
+TEST_F(IR_AddEmptyEntryPointTest, ExistingEntryPoint) {
+    auto* ep = b.CreateFunction(mod.symbols.New("main"), mod.types.Get<type::Void>(),
+                                Function::PipelineStage::kFragment);
+    b.Branch(ep->start_target, ep->end_target);
+    mod.functions.Push(ep);
+
+    auto* expect = R"(
+%fn1 = func main():void [@fragment] {
+  %fn2 = block {
+  } -> %func_end # return
+} %func_end
+
+)";
+
+    Run<AddEmptyEntryPoint>();
+
+    EXPECT_EQ(expect, str());
+}
+
+}  // namespace
+}  // namespace tint::ir::transform

src/tint/ir/transform/test_helper.h

@@ -0,0 +1,72 @@
+// Copyright 2023 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.
+
+#ifndef SRC_TINT_IR_TRANSFORM_TEST_HELPER_H_
+#define SRC_TINT_IR_TRANSFORM_TEST_HELPER_H_
+
+#include <memory>
+#include <string>
+#include <utility>
+#include <vector>
+
+#include "gtest/gtest.h"
+#include "src/tint/ir/builder.h"
+#include "src/tint/ir/disassembler.h"
+#include "src/tint/ir/transform/transform.h"
+#include "src/tint/transform/manager.h"
+
+namespace tint::ir::transform {
+
+/// Helper class for testing IR transforms.
+template <typename BASE>
+class TransformTestBase : public BASE {
+  public:
+    /// Transforms the module, using transforms in `TRANSFORMS`.
+    /// @param data the optional Transform::DataMap to pass to Transform::Run()
+    /// @returns the transform outputs, if any
+    template <typename... TRANSFORMS>
+    Transform::DataMap Run(const Transform::DataMap& data = {}) {
+        tint::transform::Manager manager;
+        tint::transform::DataMap outputs;
+        for (auto* transform_ptr : std::initializer_list<Transform*>{new TRANSFORMS()...}) {
+            manager.append(std::unique_ptr<Transform>(transform_ptr));
+        }
+        manager.Run(&mod, data, outputs);
+        return outputs;
+    }
+
+    /// @returns the transformed module as a disassembled string
+    std::string str() {
+        ir::Disassembler dis(mod);
+        return "\n" + dis.Disassemble();
+    }
+
+  protected:
+    /// The test IR module.
+    ir::Module mod;
+    /// The test IR builder.
+    ir::Builder b{mod};
+
+  private:
+    std::vector<std::unique_ptr<Source::File>> files_;
+};
+
+using TransformTest = TransformTestBase<testing::Test>;
+
+template <typename T>
+using TransformTestWithParam = TransformTestBase<testing::TestWithParam<T>>;
+
+}  // namespace tint::ir::transform
+
+#endif  // SRC_TINT_IR_TRANSFORM_TEST_HELPER_H_

src/tint/writer/spirv/generator.cc

@@ -41,14 +41,15 @@ Result Generate(const Program* program, const Options& options) {
 #if TINT_BUILD_IR
     if (options.use_tint_ir) {
         // Convert the AST program to an IR module.
-        auto ir = ir::FromProgram(program);
-        if (!ir) {
-            result.error = "IR converter: " + ir.Failure();
+        auto converted = ir::FromProgram(program);
+        if (!converted) {
+            result.error = "IR converter: " + converted.Failure();
             return result;
         }
 
         // Generate the SPIR-V code.
-        auto impl = std::make_unique<GeneratorImplIr>(&ir.Get(), zero_initialize_workgroup_memory);
+        auto ir = converted.Move();
+        auto impl = std::make_unique<GeneratorImplIr>(&ir, zero_initialize_workgroup_memory);
         result.success = impl->Generate();
         result.error = impl->Diagnostics().str();
         result.spirv = std::move(impl->Result());

src/tint/writer/spirv/generator_impl_ir.cc

@@ -19,7 +19,9 @@
 #include "src/tint/ir/block.h"
 #include "src/tint/ir/function_terminator.h"
 #include "src/tint/ir/module.h"
+#include "src/tint/ir/transform/add_empty_entry_point.h"
 #include "src/tint/switch.h"
+#include "src/tint/transform/manager.h"
 #include "src/tint/type/bool.h"
 #include "src/tint/type/f16.h"
 #include "src/tint/type/f32.h"
@@ -28,14 +30,32 @@
 #include "src/tint/type/u32.h"
 #include "src/tint/type/vector.h"
 #include "src/tint/type/void.h"
+#include "src/tint/writer/spirv/generator.h"
 #include "src/tint/writer/spirv/module.h"
 
 namespace tint::writer::spirv {
 
-GeneratorImplIr::GeneratorImplIr(const ir::Module* module, bool zero_init_workgroup_mem)
+namespace {
+
+void Sanitize(ir::Module* module) {
+    transform::Manager manager;
+    transform::DataMap data;
+
+    manager.Add<ir::transform::AddEmptyEntryPoint>();
+
+    transform::DataMap outputs;
+    manager.Run(module, data, outputs);
+}
+
+}  // namespace
+
+GeneratorImplIr::GeneratorImplIr(ir::Module* module, bool zero_init_workgroup_mem)
     : ir_(module), zero_init_workgroup_memory_(zero_init_workgroup_mem) {}
 
 bool GeneratorImplIr::Generate() {
+    // Run the IR transformations to prepare for SPIR-V emission.
+    Sanitize(ir_);
+
     // TODO(crbug.com/tint/1906): Check supported extensions.
 
     module_.PushCapability(SpvCapabilityShader);

src/tint/writer/spirv/generator_impl_ir.h

@@ -47,7 +47,7 @@ class GeneratorImplIr {
     /// @param module the Tint IR module to generate
     /// @param zero_init_workgroup_memory `true` to initialize all the variables in the Workgroup
     ///                                   storage class with OpConstantNull
-    GeneratorImplIr(const ir::Module* module, bool zero_init_workgroup_memory);
+    GeneratorImplIr(ir::Module* module, bool zero_init_workgroup_memory);
 
     /// @returns true on successful generation; false otherwise
     bool Generate();
@@ -100,7 +100,7 @@ class GeneratorImplIr {
     /// @returns the result ID of the constant
     uint32_t Constant(const constant::Value* constant);
 
-    const ir::Module* ir_;
+    ir::Module* ir_;
     spirv::Module module_;
     BinaryWriter writer_;
     diag::List diagnostics_;

src/tint/writer/spirv/generator_impl_ir_test.cc

@@ -14,15 +14,17 @@
 
 #include "src/tint/writer/spirv/test_helper_ir.h"
 
+#include "gmock/gmock.h"
+
 namespace tint::writer::spirv {
 namespace {
 
 TEST_F(SpvGeneratorImplTest, ModuleHeader) {
     ASSERT_TRUE(generator_.Generate()) << generator_.Diagnostics().str();
     auto got = Disassemble(generator_.Result());
-    EXPECT_EQ(got, R"(OpCapability Shader
+    EXPECT_THAT(got, testing::StartsWith(R"(OpCapability Shader
 OpMemoryModel Logical GLSL450
-)");
+)"));
 }
 
 }  // namespace