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>
diff --git a/src/BUILD.gn b/src/BUILD.gn
index 32a2b15..04afd87 100644
--- a/src/BUILD.gn
+++ b/src/BUILD.gn
@@ -416,6 +416,8 @@
"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",
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 0cf9989..9226231 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -286,6 +286,8 @@
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_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
diff --git a/src/transform/add_empty_entry_point.cc b/src/transform/add_empty_entry_point.cc
new file mode 100644
index 0000000..a9f057e
--- /dev/null
+++ b/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
diff --git a/src/transform/add_empty_entry_point.h b/src/transform/add_empty_entry_point.h
new file mode 100644
index 0000000..dc65a0b
--- /dev/null
+++ b/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_
diff --git a/src/transform/add_empty_entry_point_test.cc b/src/transform/add_empty_entry_point_test.cc
new file mode 100644
index 0000000..722af1a
--- /dev/null
+++ b/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
diff --git a/src/transform/hlsl.cc b/src/transform/hlsl.cc
index 224d416..0fe35b5 100644
--- a/src/transform/hlsl.cc
+++ b/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 @@
manager.Add<ExternalTextureTransform>();
manager.Add<PromoteInitializersToConstVar>();
manager.Add<PadArrayElements>();
+ manager.Add<AddEmptyEntryPoint>();
data.Add<CanonicalizeEntryPointIO::Config>(
CanonicalizeEntryPointIO::ShaderStyle::kHlsl);
@@ -81,24 +83,12 @@
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;
diff --git a/src/transform/hlsl.h b/src/transform/hlsl.h
index 01aa2d5..ea8a27e 100644
--- a/src/transform/hlsl.h
+++ b/src/transform/hlsl.h
@@ -55,10 +55,6 @@
/// @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
diff --git a/src/transform/hlsl_test.cc b/src/transform/hlsl_test.cc
index b7a8975..4389d7e 100644
--- a/src/transform/hlsl_test.cc
+++ b/src/transform/hlsl_test.cc
@@ -22,19 +22,7 @@
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
diff --git a/src/transform/spirv.cc b/src/transform/spirv.cc
index ea2a70e..6bd220a 100644
--- a/src/transform/spirv.cc
+++ b/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 @@
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 @@
ProgramBuilder builder;
CloneContext ctx(&builder, &transformedInput.program);
HandleSampleMaskBuiltins(ctx);
- AddEmptyEntryPoint(ctx);
ctx.Clone();
builder.SetTransformApplied(this);
@@ -122,17 +123,6 @@
}
}
-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) {}
diff --git a/src/transform/spirv.h b/src/transform/spirv.h
index 00e8d2d..faf5b96 100644
--- a/src/transform/spirv.h
+++ b/src/transform/spirv.h
@@ -71,8 +71,6 @@
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
diff --git a/src/transform/spirv_test.cc b/src/transform/spirv_test.cc
index 6f68d49..26ed757 100644
--- a/src/transform/spirv_test.cc
+++ b/src/transform/spirv_test.cc
@@ -101,20 +101,6 @@
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"(
diff --git a/test/BUILD.gn b/test/BUILD.gn
index 8f5fc33..56483e5 100644
--- a/test/BUILD.gn
+++ b/test/BUILD.gn
@@ -287,6 +287,7 @@
"../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",
