transform: Add transform to add empty entry point

Use this from the HLSL and SPIR-V sanitizers, instead of duplicating
this logic for them.

This is step towards removing the sanitizers completely.

Change-Id: Ifa9f23d84fd3505d30a928c260181a699c5f1783
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/63582
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>

src/BUILD.gn

@@ -416,6 +416,8 @@ libtint_source_set("libtint_core_all_src") {
     "symbol_table.cc",
     "symbol_table.h",
     "traits.h",
+    "transform/add_empty_entry_point.cc",
+    "transform/add_empty_entry_point.h",
     "transform/array_length_from_uniform.cc",
     "transform/array_length_from_uniform.h",
     "transform/binding_remapper.cc",

src/CMakeLists.txt

@@ -286,6 +286,8 @@ set(TINT_LIB_SRCS
   symbol.cc
   symbol.h
   traits.h
+  transform/add_empty_entry_point.cc
+  transform/add_empty_entry_point.h
   transform/array_length_from_uniform.cc
   transform/array_length_from_uniform.h
   transform/binding_remapper.cc
@@ -921,6 +923,7 @@ if(${TINT_BUILD_TESTS})
 
   if(${TINT_BUILD_WGSL_READER} AND ${TINT_BUILD_WGSL_WRITER})
     list(APPEND TINT_TEST_SRCS
+      transform/add_empty_entry_point_test.cc
       transform/array_length_from_uniform_test.cc
       transform/binding_remapper_test.cc
       transform/calculate_array_length_test.cc

src/transform/add_empty_entry_point.cc

@@ -0,0 +1,45 @@
+// Copyright 2021 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/transform/add_empty_entry_point.h"
+
+#include <utility>
+
+#include "src/program_builder.h"
+
+TINT_INSTANTIATE_TYPEINFO(tint::transform::AddEmptyEntryPoint);
+
+namespace tint {
+namespace transform {
+
+AddEmptyEntryPoint::AddEmptyEntryPoint() = default;
+
+AddEmptyEntryPoint::~AddEmptyEntryPoint() = default;
+
+void AddEmptyEntryPoint::Run(CloneContext& ctx, const DataMap&, DataMap&) {
+  for (auto* func : ctx.src->AST().Functions()) {
+    if (func->IsEntryPoint()) {
+      ctx.Clone();
+      return;
+    }
+  }
+  ctx.dst->Func(ctx.dst->Symbols().New("unused_entry_point"), {},
+                ctx.dst->ty.void_(), {},
+                {ctx.dst->Stage(ast::PipelineStage::kCompute),
+                 ctx.dst->WorkgroupSize(1)});
+  ctx.Clone();
+}
+
+}  // namespace transform
+}  // namespace tint

src/transform/add_empty_entry_point.h

@@ -0,0 +1,44 @@
+// Copyright 2021 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_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_
+#define SRC_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_
+
+#include "src/transform/transform.h"
+
+namespace tint {
+namespace transform {
+
+/// Add an empty entry point to the module, if no other entry points exist.
+class AddEmptyEntryPoint : public Castable<AddEmptyEntryPoint, Transform> {
+ public:
+  /// Constructor
+  AddEmptyEntryPoint();
+  /// Destructor
+  ~AddEmptyEntryPoint() override;
+
+ protected:
+  /// Runs the transform using the CloneContext built for transforming a
+  /// program. Run() is responsible for calling Clone() on the CloneContext.
+  /// @param ctx the CloneContext primed with the input program and
+  /// ProgramBuilder
+  /// @param inputs optional extra transform-specific input data
+  /// @param outputs optional extra transform-specific output data
+  void Run(CloneContext& ctx, const DataMap& inputs, DataMap& outputs) override;
+};
+
+}  // namespace transform
+}  // namespace tint
+
+#endif  // SRC_TRANSFORM_ADD_EMPTY_ENTRY_POINT_H_

src/transform/add_empty_entry_point_test.cc

@@ -0,0 +1,73 @@
+// Copyright 2021 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/transform/add_empty_entry_point.h"
+
+#include <utility>
+
+#include "src/transform/test_helper.h"
+
+namespace tint {
+namespace transform {
+namespace {
+
+using AddEmptyEntryPointTest = TransformTest;
+
+TEST_F(AddEmptyEntryPointTest, EmptyModule) {
+  auto* src = R"()";
+
+  auto* expect = R"(
+[[stage(compute), workgroup_size(1)]]
+fn unused_entry_point() {
+}
+)";
+
+  auto got = Run<AddEmptyEntryPoint>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddEmptyEntryPointTest, ExistingEntryPoint) {
+  auto* src = R"(
+[[stage(fragment)]]
+fn main() {
+}
+)";
+
+  auto* expect = src;
+
+  auto got = Run<AddEmptyEntryPoint>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddEmptyEntryPointTest, NameClash) {
+  auto* src = R"(var<private> unused_entry_point : f32;)";
+
+  auto* expect = R"(
+[[stage(compute), workgroup_size(1)]]
+fn unused_entry_point_1() {
+}
+
+var<private> unused_entry_point : f32;
+)";
+
+  auto got = Run<AddEmptyEntryPoint>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+}  // namespace
+}  // namespace transform
+}  // namespace tint

src/transform/hlsl.cc

@@ -17,6 +17,7 @@
 #include <utility>
 
 #include "src/program_builder.h"
+#include "src/transform/add_empty_entry_point.h"
 #include "src/transform/calculate_array_length.h"
 #include "src/transform/canonicalize_entry_point_io.h"
 #include "src/transform/decompose_memory_access.h"
@@ -71,6 +72,7 @@ Output Hlsl::Run(const Program* in, const DataMap& inputs) {
   manager.Add<ExternalTextureTransform>();
   manager.Add<PromoteInitializersToConstVar>();
   manager.Add<PadArrayElements>();
+  manager.Add<AddEmptyEntryPoint>();
 
   data.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
@@ -81,24 +83,12 @@ Output Hlsl::Run(const Program* in, const DataMap& inputs) {
 
   ProgramBuilder builder;
   CloneContext ctx(&builder, &out.program);
-  AddEmptyEntryPoint(ctx);
+  // TODO(jrprice): Move the sanitizer into the backend.
   ctx.Clone();
   builder.SetTransformApplied(this);
   return Output{Program(std::move(builder))};
 }
 
-void Hlsl::AddEmptyEntryPoint(CloneContext& ctx) const {
-  for (auto* func : ctx.src->AST().Functions()) {
-    if (func->IsEntryPoint()) {
-      return;
-    }
-  }
-  ctx.dst->Func(ctx.dst->Symbols().New("unused_entry_point"), {},
-                ctx.dst->ty.void_(), {},
-                {ctx.dst->Stage(ast::PipelineStage::kCompute),
-                 ctx.dst->WorkgroupSize(1)});
-}
-
 Hlsl::Config::Config(bool disable_wi) : disable_workgroup_init(disable_wi) {}
 Hlsl::Config::Config(const Config&) = default;
 Hlsl::Config::~Config() = default;

src/transform/hlsl.h

@@ -55,10 +55,6 @@ class Hlsl : public Castable<Hlsl, Transform> {
   /// @param data optional extra transform-specific data
   /// @returns the transformation result
   Output Run(const Program* program, const DataMap& data = {}) override;
-
- private:
-  /// Add an empty shader entry point if none exist in the module.
-  void AddEmptyEntryPoint(CloneContext& ctx) const;
 };
 
 }  // namespace transform

src/transform/hlsl_test.cc

@@ -22,19 +22,7 @@ namespace {
 
 using HlslTest = TransformTest;
 
-TEST_F(HlslTest, AddEmptyEntryPoint) {
-  auto* src = R"()";
-
-  auto* expect = R"(
-[[stage(compute), workgroup_size(1)]]
-fn unused_entry_point() {
-}
-)";
-
-  auto got = Run<Hlsl>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
+// TODO(jrprice): Remove this file when we remove the sanitizer transforms.
 
 }  // namespace
 }  // namespace transform

src/transform/spirv.cc

@@ -20,6 +20,7 @@
 #include "src/ast/stage_decoration.h"
 #include "src/program_builder.h"
 #include "src/sem/variable.h"
+#include "src/transform/add_empty_entry_point.h"
 #include "src/transform/canonicalize_entry_point_io.h"
 #include "src/transform/external_texture_transform.h"
 #include "src/transform/fold_constants.h"
@@ -52,6 +53,7 @@ Output Spirv::Run(const Program* in, const DataMap& data) {
   manager.Add<ExternalTextureTransform>();
   manager.Add<ForLoopToLoop>();  // Must come after ZeroInitWorkgroupMemory
   manager.Add<CanonicalizeEntryPointIO>();
+  manager.Add<AddEmptyEntryPoint>();
 
   internal_inputs.Add<CanonicalizeEntryPointIO::Config>(
       CanonicalizeEntryPointIO::Config(
@@ -67,7 +69,6 @@ Output Spirv::Run(const Program* in, const DataMap& data) {
   ProgramBuilder builder;
   CloneContext ctx(&builder, &transformedInput.program);
   HandleSampleMaskBuiltins(ctx);
-  AddEmptyEntryPoint(ctx);
   ctx.Clone();
 
   builder.SetTransformApplied(this);
@@ -122,17 +123,6 @@ void Spirv::HandleSampleMaskBuiltins(CloneContext& ctx) const {
   }
 }
 
-void Spirv::AddEmptyEntryPoint(CloneContext& ctx) const {
-  for (auto* func : ctx.src->AST().Functions()) {
-    if (func->IsEntryPoint()) {
-      return;
-    }
-  }
-  ctx.dst->Func(ctx.dst->Sym("unused_entry_point"), {}, ctx.dst->ty.void_(), {},
-                {ctx.dst->Stage(ast::PipelineStage::kCompute),
-                 ctx.dst->WorkgroupSize(1)});
-}
-
 Spirv::Config::Config(bool emit_vps, bool disable_wi)
     : emit_vertex_point_size(emit_vps), disable_workgroup_init(disable_wi) {}
 

src/transform/spirv.h

@@ -71,8 +71,6 @@ class Spirv : public Castable<Spirv, Transform> {
  private:
   /// Change type of sample mask builtin variables to single element arrays.
   void HandleSampleMaskBuiltins(CloneContext& ctx) const;
-  /// Add an empty shader entry point if none exist in the module.
-  void AddEmptyEntryPoint(CloneContext& ctx) const;
 };
 
 }  // namespace transform

src/transform/spirv_test.cc

@@ -101,20 +101,6 @@ fn main() {
   EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(SpirvTest, AddEmptyEntryPoint) {
-  auto* src = R"()";
-
-  auto* expect = R"(
-[[stage(compute), workgroup_size(1)]]
-fn unused_entry_point() {
-}
-)";
-
-  auto got = Run<Spirv>(src);
-
-  EXPECT_EQ(expect, str(got));
-}
-
 // Test that different transforms within the sanitizer interact correctly.
 TEST_F(SpirvTest, MultipleTransforms) {
   auto* src = R"(

test/BUILD.gn

@@ -287,6 +287,7 @@ tint_unittests_source_set("tint_unittests_core_src") {
     "../src/symbol_table_test.cc",
     "../src/symbol_test.cc",
     "../src/traits_test.cc",
+    "../src/transform/add_empty_entry_point_test.cc",
     "../src/transform/array_length_from_uniform_test.cc",
     "../src/transform/binding_remapper_test.cc",
     "../src/transform/calculate_array_length_test.cc",