Import Tint changes from Dawn
Changes:
- 2517d99f351f307f0f339dc2c4173ed2e265f17d Fixup kokoro lint issue by dan sinclair <dsinclair@chromium.org>
- 4d5723b5e57070448cbc1c832e3bcec0fa741972 Remove fuzzer CPPLINT file. by dan sinclair <dsinclair@chromium.org>
- a5f5f5a1aafa781d40fd00ad31f608738e8d68b7 use stamp file for tint_generate_wgsl_corpus by Fumitoshi Ukai <ukai@google.com>
- c85f8d828e717076d750327797f51cd8c5f456ec tint: Fix doxygen failure by Ben Clayton <bclayton@google.com>
- f625a6d57e462946f2c38a4a4c611ee1939f4b53 tint: Make GLSL backend consistent with the others by Antonio Maiorano <amaiorano@google.com>
- b5c46c30ec440899910fdabcafa3b01027bb5e9e tint: spirv writer: add a GeneralImpl to be consistent wi... by Antonio Maiorano <amaiorano@google.com>
GitOrigin-RevId: 2517d99f351f307f0f339dc2c4173ed2e265f17d
Change-Id: I2abe6ee2d1a03d4e28275187dac5c0f88ee0323c
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/86641
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
Commit-Queue: Ben Clayton <bclayton@google.com>
diff --git a/include/tint/.clang-format b/include/tint/.clang-format
new file mode 100644
index 0000000..2fb833a
--- /dev/null
+++ b/include/tint/.clang-format
@@ -0,0 +1,2 @@
+# http://clang.llvm.org/docs/ClangFormatStyleOptions.html
+BasedOnStyle: Chromium
diff --git a/include/tint/tint.h b/include/tint/tint.h
index 1a04196..2b8430e 100644
--- a/include/tint/tint.h
+++ b/include/tint/tint.h
@@ -61,7 +61,6 @@
#endif // TINT_BUILD_HLSL_WRITER
#if TINT_BUILD_GLSL_WRITER
-#include "src/tint/transform/glsl.h"
#include "src/tint/writer/glsl/generator.h"
#endif // TINT_BUILD_GLSL_WRITER
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 79beb3d..5594bf8 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -688,6 +688,8 @@
"writer/spirv/function.h",
"writer/spirv/generator.cc",
"writer/spirv/generator.h",
+ "writer/spirv/generator_impl.cc",
+ "writer/spirv/generator_impl.h",
"writer/spirv/instruction.cc",
"writer/spirv/instruction.h",
"writer/spirv/operand.cc",
@@ -749,8 +751,6 @@
libtint_source_set("libtint_glsl_writer_src") {
sources = [
- "transform/glsl.cc",
- "transform/glsl.h",
"writer/glsl/generator.cc",
"writer/glsl/generator.h",
"writer/glsl/generator_impl.cc",
diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt
index 046b814..527eeaf 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -339,8 +339,6 @@
transform/for_loop_to_loop.h
transform/expand_compound_assignment.cc
transform/expand_compound_assignment.h
- transform/glsl.cc
- transform/glsl.h
transform/loop_to_for_loop.cc
transform/loop_to_for_loop.h
transform/manager.cc
@@ -519,6 +517,8 @@
writer/spirv/function.h
writer/spirv/generator.cc
writer/spirv/generator.h
+ writer/spirv/generator_impl.cc
+ writer/spirv/generator_impl.h
writer/spirv/instruction.cc
writer/spirv/instruction.h
writer/spirv/operand.cc
diff --git a/src/tint/fuzzers/BUILD.gn b/src/tint/fuzzers/BUILD.gn
index 3253e94..afd1e3a 100644
--- a/src/tint/fuzzers/BUILD.gn
+++ b/src/tint/fuzzers/BUILD.gn
@@ -22,14 +22,16 @@
import("//testing/libfuzzer/fuzzer_test.gni")
fuzzer_corpus_wgsl_dir = "${target_gen_dir}/fuzzer_corpus_wgsl"
+ fuzzer_corpus_wgsl_stamp = "${fuzzer_corpus_wgsl_dir}.stamp"
action("tint_generate_wgsl_corpus") {
script = "generate_wgsl_corpus.py"
sources = [ "generate_wgsl_corpus.py" ]
args = [
+ "--stamp=${fuzzer_corpus_wgsl_stamp}",
rebase_path("${tint_root_dir}/test", root_build_dir),
rebase_path(fuzzer_corpus_wgsl_dir, root_build_dir),
]
- outputs = [ fuzzer_corpus_wgsl_dir ]
+ outputs = [ fuzzer_corpus_wgsl_stamp ]
}
tint_fuzzer_common_libfuzzer_options = [
diff --git a/src/tint/fuzzers/generate_wgsl_corpus.py b/src/tint/fuzzers/generate_wgsl_corpus.py
index 65c564f..983acd9 100644
--- a/src/tint/fuzzers/generate_wgsl_corpus.py
+++ b/src/tint/fuzzers/generate_wgsl_corpus.py
@@ -26,6 +26,7 @@
# Usage:
# generate_wgsl_corpus.py <input_dir> <corpus_dir>
+import optparse
import os
import pathlib
import shutil
@@ -40,11 +41,14 @@
def main():
- if len(sys.argv) != 3:
- print("Usage: " + sys.argv[0] + " <input dir> <output dir>")
- return 1
- input_dir: str = os.path.abspath(sys.argv[1].rstrip(os.sep))
- corpus_dir: str = os.path.abspath(sys.argv[2])
+ parser = optparse.OptionParser(
+ usage="usage: %prog [option] input-dir output-dir")
+ parser.add_option('--stamp', dest='stamp', help='stamp file')
+ options, args = parser.parse_args(sys.argv[1:])
+ if len(args) != 2:
+ parser.error("incorrect number of arguments")
+ input_dir: str = os.path.abspath(args[0].rstrip(os.sep))
+ corpus_dir: str = os.path.abspath(args[1])
if os.path.exists(corpus_dir):
shutil.rmtree(corpus_dir)
os.makedirs(corpus_dir)
@@ -53,6 +57,8 @@
continue
out_file = in_file[len(input_dir) + 1:].replace(os.sep, '_')
shutil.copy(in_file, corpus_dir + os.sep + out_file)
+ if options.stamp:
+ pathlib.Path(options.stamp).touch(mode=0o644, exist_ok=True)
if __name__ == "__main__":
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/CPPLINT.cfg b/src/tint/fuzzers/tint_regex_fuzzer/CPPLINT.cfg
deleted file mode 100644
index 96988ad..0000000
--- a/src/tint/fuzzers/tint_regex_fuzzer/CPPLINT.cfg
+++ /dev/null
@@ -1 +0,0 @@
-filter=-build/c++11
\ No newline at end of file
diff --git a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
index 2397e7f..d507944 100644
--- a/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
+++ b/src/tint/fuzzers/tint_regex_fuzzer/wgsl_mutator.cc
@@ -17,7 +17,7 @@
#include <cassert>
#include <cstring>
#include <map>
-#include <regex>
+#include <regex> // NOLINT(build/c++11)
#include <string>
#include <utility>
#include <vector>
diff --git a/src/tint/transform/glsl.cc b/src/tint/transform/glsl.cc
deleted file mode 100644
index f0120e8..0000000
--- a/src/tint/transform/glsl.cc
+++ /dev/null
@@ -1,145 +0,0 @@
-// 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/tint/transform/glsl.h"
-
-#include <utility>
-
-#include "src/tint/program_builder.h"
-#include "src/tint/transform/add_empty_entry_point.h"
-#include "src/tint/transform/add_spirv_block_attribute.h"
-#include "src/tint/transform/binding_remapper.h"
-#include "src/tint/transform/builtin_polyfill.h"
-#include "src/tint/transform/canonicalize_entry_point_io.h"
-#include "src/tint/transform/combine_samplers.h"
-#include "src/tint/transform/decompose_memory_access.h"
-#include "src/tint/transform/expand_compound_assignment.h"
-#include "src/tint/transform/fold_trivial_single_use_lets.h"
-#include "src/tint/transform/loop_to_for_loop.h"
-#include "src/tint/transform/manager.h"
-#include "src/tint/transform/promote_initializers_to_const_var.h"
-#include "src/tint/transform/promote_side_effects_to_decl.h"
-#include "src/tint/transform/remove_phonies.h"
-#include "src/tint/transform/renamer.h"
-#include "src/tint/transform/simplify_pointers.h"
-#include "src/tint/transform/single_entry_point.h"
-#include "src/tint/transform/unshadow.h"
-#include "src/tint/transform/unwind_discard_functions.h"
-#include "src/tint/transform/zero_init_workgroup_memory.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
-
-TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl);
-TINT_INSTANTIATE_TYPEINFO(tint::transform::Glsl::Config);
-
-namespace tint::transform {
-
-Glsl::Glsl() = default;
-Glsl::~Glsl() = default;
-
-Output Glsl::Run(const Program* in, const DataMap& inputs) const {
- Manager manager;
- DataMap data;
-
- auto* cfg = inputs.Get<Config>();
-
- { // Builtin polyfills
- BuiltinPolyfill::Builtins polyfills;
- polyfills.count_leading_zeros = true;
- polyfills.count_trailing_zeros = true;
- polyfills.extract_bits = BuiltinPolyfill::Level::kClampParameters;
- polyfills.first_leading_bit = true;
- polyfills.first_trailing_bit = true;
- polyfills.insert_bits = BuiltinPolyfill::Level::kClampParameters;
- data.Add<BuiltinPolyfill::Config>(polyfills);
- manager.Add<BuiltinPolyfill>();
- }
-
- if (cfg && !cfg->entry_point.empty()) {
- manager.Add<SingleEntryPoint>();
- data.Add<SingleEntryPoint::Config>(cfg->entry_point);
- }
- manager.Add<Renamer>();
- data.Add<Renamer::Config>(Renamer::Target::kGlslKeywords,
- /* preserve_unicode */ false);
- manager.Add<Unshadow>();
-
- // Attempt to convert `loop`s into for-loops. This is to try and massage the
- // output into something that will not cause FXC to choke or misbehave.
- manager.Add<FoldTrivialSingleUseLets>();
- manager.Add<LoopToForLoop>();
-
- if (!cfg || !cfg->disable_workgroup_init) {
- // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
- // ZeroInitWorkgroupMemory may inject new builtin parameters.
- manager.Add<ZeroInitWorkgroupMemory>();
- }
- manager.Add<CanonicalizeEntryPointIO>();
- manager.Add<ExpandCompoundAssignment>();
- manager.Add<PromoteSideEffectsToDecl>();
- manager.Add<UnwindDiscardFunctions>();
- manager.Add<SimplifyPointers>();
-
- manager.Add<RemovePhonies>();
-
- if (cfg && cfg->generate_external_texture_bindings) {
- auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
- data.Add<MultiplanarExternalTexture::NewBindingPoints>(new_bindings_map);
- }
- manager.Add<MultiplanarExternalTexture>();
-
- manager.Add<CombineSamplers>();
- if (auto* binding_info = inputs.Get<CombineSamplers::BindingInfo>()) {
- data.Add<CombineSamplers::BindingInfo>(*binding_info);
- } else {
- data.Add<CombineSamplers::BindingInfo>(CombineSamplers::BindingMap(),
- sem::BindingPoint());
- }
- manager.Add<BindingRemapper>();
- if (auto* remappings = inputs.Get<BindingRemapper::Remappings>()) {
- data.Add<BindingRemapper::Remappings>(*remappings);
- } else {
- BindingRemapper::BindingPoints bp;
- BindingRemapper::AccessControls ac;
- data.Add<BindingRemapper::Remappings>(bp, ac, /* mayCollide */ true);
- }
-
- manager.Add<PromoteInitializersToConstVar>();
-
- manager.Add<AddEmptyEntryPoint>();
- manager.Add<AddSpirvBlockAttribute>();
-
- data.Add<CanonicalizeEntryPointIO::Config>(
- CanonicalizeEntryPointIO::ShaderStyle::kGlsl);
- auto out = manager.Run(in, data);
- if (!out.program.IsValid()) {
- return out;
- }
-
- ProgramBuilder builder;
- CloneContext ctx(&builder, &out.program);
- ctx.Clone();
- return Output{Program(std::move(builder))};
-}
-
-Glsl::Config::Config(const std::string& entry_point_in,
- bool disable_workgroup_init_in,
- bool generate_external_texture_bindings_in)
- : entry_point(entry_point_in),
- disable_workgroup_init(disable_workgroup_init_in),
- generate_external_texture_bindings(
- generate_external_texture_bindings_in) {}
-Glsl::Config::Config(const Config&) = default;
-Glsl::Config::~Config() = default;
-
-} // namespace tint::transform
diff --git a/src/tint/transform/glsl.h b/src/tint/transform/glsl.h
deleted file mode 100644
index c84ff29..0000000
--- a/src/tint/transform/glsl.h
+++ /dev/null
@@ -1,75 +0,0 @@
-// 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_TINT_TRANSFORM_GLSL_H_
-#define SRC_TINT_TRANSFORM_GLSL_H_
-
-#include <string>
-
-#include "src/tint/transform/transform.h"
-
-// Forward declarations
-namespace tint {
-class CloneContext;
-} // namespace tint
-
-namespace tint::transform {
-
-/// Glsl is a transform used to sanitize a Program for use with the Glsl writer.
-/// Passing a non-sanitized Program to the Glsl writer will result in undefined
-/// behavior.
-class Glsl final : public Castable<Glsl, Transform> {
- public:
- /// Configuration options for the Glsl sanitizer transform.
- struct Config final : public Castable<Data, transform::Data> {
- /// Constructor
- /// @param entry_point the root entry point function to generate
- /// @param disable_workgroup_init `true` to disable workgroup memory zero
- /// initialization
- /// @param generate_external_texture_bindings 'true' to generates binding
- /// mappings for external textures
- explicit Config(const std::string& entry_point,
- bool disable_workgroup_init,
- bool generate_external_texture_bindings);
-
- /// Copy constructor
- Config(const Config&);
-
- /// Destructor
- ~Config() override;
-
- /// GLSL generator wraps a single entry point in a main() function.
- std::string entry_point;
-
- /// Set to `true` to disable workgroup memory zero initialization
- bool disable_workgroup_init = false;
-
- /// Set to 'true' to generates binding mappings for external textures
- bool generate_external_texture_bindings = false;
- };
-
- /// Constructor
- Glsl();
- ~Glsl() override;
-
- /// Runs the transform on `program`, returning the transformation result.
- /// @param program the source program to transform
- /// @param data optional extra transform-specific data
- /// @returns the transformation result
- Output Run(const Program* program, const DataMap& data = {}) const override;
-};
-
-} // namespace tint::transform
-
-#endif // SRC_TINT_TRANSFORM_GLSL_H_
diff --git a/src/tint/transform/glsl_test.cc b/src/tint/transform/glsl_test.cc
deleted file mode 100644
index 555701b..0000000
--- a/src/tint/transform/glsl_test.cc
+++ /dev/null
@@ -1,39 +0,0 @@
-// 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/tint/transform/glsl.h"
-
-#include "src/tint/transform/test_helper.h"
-
-namespace tint::transform {
-namespace {
-
-using GlslTest = TransformTest;
-
-TEST_F(GlslTest, AddEmptyEntryPoint) {
- auto* src = R"()";
-
- auto* expect = R"(
-@stage(compute) @workgroup_size(1)
-fn unused_entry_point() {
-}
-)";
-
- auto got = Run<Glsl>(src);
-
- EXPECT_EQ(expect, str(got));
-}
-
-} // namespace
-} // namespace tint::transform
diff --git a/src/tint/writer/glsl/generator.cc b/src/tint/writer/glsl/generator.cc
index ad06ac4..83ae8c4 100644
--- a/src/tint/writer/glsl/generator.cc
+++ b/src/tint/writer/glsl/generator.cc
@@ -16,7 +16,6 @@
#include "src/tint/transform/binding_remapper.h"
#include "src/tint/transform/combine_samplers.h"
-#include "src/tint/transform/glsl.h"
#include "src/tint/writer/glsl/generator_impl.h"
namespace tint::writer::glsl {
@@ -34,34 +33,25 @@
const std::string& entry_point) {
Result result;
- // Run the GLSL sanitizer.
- transform::DataMap data;
- data.Add<transform::BindingRemapper::Remappings>(options.binding_points,
- options.access_controls,
- options.allow_collisions);
- data.Add<transform::CombineSamplers::BindingInfo>(
- options.binding_map, options.placeholder_binding_point);
- data.Add<transform::Glsl::Config>(entry_point,
- /* disable_workgroup_init */ false,
- options.generate_external_texture_bindings);
- transform::Glsl sanitizer;
- auto output = sanitizer.Run(program, data);
- if (!output.program.IsValid()) {
+ // Sanitize the program.
+ auto sanitized_result = Sanitize(program, options, entry_point);
+ if (!sanitized_result.program.IsValid()) {
result.success = false;
- result.error = output.program.Diagnostics().str();
+ result.error = sanitized_result.program.Diagnostics().str();
return result;
}
// Generate the GLSL code.
- auto impl = std::make_unique<GeneratorImpl>(&output.program, options.version);
+ auto impl = std::make_unique<GeneratorImpl>(&sanitized_result.program,
+ options.version);
result.success = impl->Generate();
result.error = impl->error();
result.glsl = impl->result();
// Collect the list of entry points in the sanitized program.
- for (auto* func : output.program.AST().Functions()) {
+ for (auto* func : sanitized_result.program.AST().Functions()) {
if (func->IsEntryPoint()) {
- auto name = output.program.Symbols().NameFor(func->symbol);
+ auto name = sanitized_result.program.Symbols().NameFor(func->symbol);
result.entry_points.push_back({name, func->PipelineStage()});
}
}
diff --git a/src/tint/writer/glsl/generator.h b/src/tint/writer/glsl/generator.h
index ab5c9bb..11206cf 100644
--- a/src/tint/writer/glsl/generator.h
+++ b/src/tint/writer/glsl/generator.h
@@ -70,6 +70,9 @@
/// generated by the BindingRemapper transform
bool allow_collisions = false;
+ /// Set to `true` to disable workgroup memory zero initialization
+ bool disable_workgroup_init = false;
+
/// Set to 'true' to generates binding mappings for external textures
bool generate_external_texture_bindings = false;
@@ -106,6 +109,7 @@
/// information.
/// @param program the program to translate to GLSL
/// @param options the configuration options to use when generating GLSL
+/// @param entry_point the entry point to generate GLSL for
/// @returns the resulting GLSL and supplementary information
Result Generate(const Program* program,
const Options& options,
diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc
index 729a8c0..ef7467d 100644
--- a/src/tint/writer/glsl/generator_impl.cc
+++ b/src/tint/writer/glsl/generator_impl.cc
@@ -45,12 +45,32 @@
#include "src/tint/sem/type_constructor.h"
#include "src/tint/sem/type_conversion.h"
#include "src/tint/sem/variable.h"
-#include "src/tint/transform/glsl.h"
+#include "src/tint/transform/add_empty_entry_point.h"
+#include "src/tint/transform/add_spirv_block_attribute.h"
+#include "src/tint/transform/binding_remapper.h"
+#include "src/tint/transform/builtin_polyfill.h"
+#include "src/tint/transform/canonicalize_entry_point_io.h"
+#include "src/tint/transform/combine_samplers.h"
+#include "src/tint/transform/decompose_memory_access.h"
+#include "src/tint/transform/expand_compound_assignment.h"
+#include "src/tint/transform/fold_trivial_single_use_lets.h"
+#include "src/tint/transform/loop_to_for_loop.h"
+#include "src/tint/transform/manager.h"
+#include "src/tint/transform/promote_initializers_to_const_var.h"
+#include "src/tint/transform/promote_side_effects_to_decl.h"
+#include "src/tint/transform/remove_phonies.h"
+#include "src/tint/transform/renamer.h"
+#include "src/tint/transform/simplify_pointers.h"
+#include "src/tint/transform/single_entry_point.h"
+#include "src/tint/transform/unshadow.h"
+#include "src/tint/transform/unwind_discard_functions.h"
+#include "src/tint/transform/zero_init_workgroup_memory.h"
#include "src/tint/utils/defer.h"
#include "src/tint/utils/map.h"
#include "src/tint/utils/scoped_assignment.h"
#include "src/tint/writer/append_vector.h"
#include "src/tint/writer/float_to_string.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
namespace {
@@ -127,6 +147,86 @@
} // namespace
+SanitizedResult::SanitizedResult() = default;
+SanitizedResult::~SanitizedResult() = default;
+SanitizedResult::SanitizedResult(SanitizedResult&&) = default;
+
+SanitizedResult Sanitize(const Program* in,
+ const Options& options,
+ const std::string& entry_point) {
+ transform::Manager manager;
+ transform::DataMap data;
+
+ { // Builtin polyfills
+ transform::BuiltinPolyfill::Builtins polyfills;
+ polyfills.count_leading_zeros = true;
+ polyfills.count_trailing_zeros = true;
+ polyfills.extract_bits =
+ transform::BuiltinPolyfill::Level::kClampParameters;
+ polyfills.first_leading_bit = true;
+ polyfills.first_trailing_bit = true;
+ polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
+ data.Add<transform::BuiltinPolyfill::Config>(polyfills);
+ manager.Add<transform::BuiltinPolyfill>();
+ }
+
+ if (!entry_point.empty()) {
+ manager.Add<transform::SingleEntryPoint>();
+ data.Add<transform::SingleEntryPoint::Config>(entry_point);
+ }
+ manager.Add<transform::Renamer>();
+ data.Add<transform::Renamer::Config>(
+ transform::Renamer::Target::kGlslKeywords,
+ /* preserve_unicode */ false);
+ manager.Add<transform::Unshadow>();
+
+ // Attempt to convert `loop`s into for-loops. This is to try and massage the
+ // output into something that will not cause FXC to choke or misbehave.
+ manager.Add<transform::FoldTrivialSingleUseLets>();
+ manager.Add<transform::LoopToForLoop>();
+
+ if (!options.disable_workgroup_init) {
+ // ZeroInitWorkgroupMemory must come before CanonicalizeEntryPointIO as
+ // ZeroInitWorkgroupMemory may inject new builtin parameters.
+ manager.Add<transform::ZeroInitWorkgroupMemory>();
+ }
+ manager.Add<transform::CanonicalizeEntryPointIO>();
+ manager.Add<transform::ExpandCompoundAssignment>();
+ manager.Add<transform::PromoteSideEffectsToDecl>();
+ manager.Add<transform::UnwindDiscardFunctions>();
+ manager.Add<transform::SimplifyPointers>();
+
+ manager.Add<transform::RemovePhonies>();
+
+ if (options.generate_external_texture_bindings) {
+ auto new_bindings_map = writer::GenerateExternalTextureBindings(in);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ new_bindings_map);
+ }
+ manager.Add<transform::MultiplanarExternalTexture>();
+
+ data.Add<transform::CombineSamplers::BindingInfo>(
+ options.binding_map, options.placeholder_binding_point);
+ manager.Add<transform::CombineSamplers>();
+
+ data.Add<transform::BindingRemapper::Remappings>(options.binding_points,
+ options.access_controls,
+ options.allow_collisions);
+ manager.Add<transform::BindingRemapper>();
+
+ manager.Add<transform::PromoteInitializersToConstVar>();
+ manager.Add<transform::AddEmptyEntryPoint>();
+ manager.Add<transform::AddSpirvBlockAttribute>();
+ data.Add<transform::CanonicalizeEntryPointIO::Config>(
+ transform::CanonicalizeEntryPointIO::ShaderStyle::kGlsl);
+
+ auto out = manager.Run(in, data);
+
+ SanitizedResult result;
+ result.program = std::move(out.program);
+ return result;
+}
+
GeneratorImpl::GeneratorImpl(const Program* program, const Version& version)
: TextGenerator(program), version_(version) {}
diff --git a/src/tint/writer/glsl/generator_impl.h b/src/tint/writer/glsl/generator_impl.h
index 37be658..9104182 100644
--- a/src/tint/writer/glsl/generator_impl.h
+++ b/src/tint/writer/glsl/generator_impl.h
@@ -35,6 +35,7 @@
#include "src/tint/scope_stack.h"
#include "src/tint/transform/decompose_memory_access.h"
#include "src/tint/utils/hash.h"
+#include "src/tint/writer/glsl/generator.h"
#include "src/tint/writer/glsl/version.h"
#include "src/tint/writer/text_generator.h"
@@ -48,6 +49,28 @@
namespace tint::writer::glsl {
+/// The result of sanitizing a program for generation.
+struct SanitizedResult {
+ /// Constructor
+ SanitizedResult();
+ /// Destructor
+ ~SanitizedResult();
+ /// Move constructor
+ SanitizedResult(SanitizedResult&&);
+
+ /// The sanitized program.
+ Program program;
+};
+
+/// Sanitize a program in preparation for generating GLSL.
+/// @program The program to sanitize
+/// @param options The HLSL generator options.
+/// @param entry_point the entry point to generate GLSL for
+/// @returns the sanitized program and any supplementary information
+SanitizedResult Sanitize(const Program* program,
+ const Options& options,
+ const std::string& entry_point);
+
/// Implementation class for GLSL generator
class GeneratorImpl : public TextGenerator {
public:
diff --git a/src/tint/writer/glsl/test_helper.h b/src/tint/writer/glsl/test_helper.h
index 969490a..7266d0a 100644
--- a/src/tint/writer/glsl/test_helper.h
+++ b/src/tint/writer/glsl/test_helper.h
@@ -20,7 +20,6 @@
#include <utility>
#include "gtest/gtest.h"
-#include "src/tint/transform/glsl.h"
#include "src/tint/transform/manager.h"
#include "src/tint/writer/glsl/generator_impl.h"
@@ -60,8 +59,10 @@
/// @note The generator is only built once. Multiple calls to Build() will
/// return the same GeneratorImpl without rebuilding.
/// @param version the GLSL version
+ /// @param options the GLSL backend options
/// @return the built generator
- GeneratorImpl& SanitizeAndBuild(Version version = Version()) {
+ GeneratorImpl& SanitizeAndBuild(Version version = Version(),
+ const Options& options = {}) {
if (gen_) {
return *gen_;
}
@@ -76,15 +77,14 @@
<< formatter.format(program->Diagnostics());
}();
- transform::Manager transform_manager;
- transform::DataMap transform_data;
- transform_manager.Add<tint::transform::Glsl>();
- auto result = transform_manager.Run(program.get(), transform_data);
+ auto sanitized_result =
+ Sanitize(program.get(), options, /* entry_point */ "");
[&]() {
- ASSERT_TRUE(result.program.IsValid())
- << formatter.format(result.program.Diagnostics());
+ ASSERT_TRUE(sanitized_result.program.IsValid())
+ << formatter.format(sanitized_result.program.Diagnostics());
}();
- *program = std::move(result.program);
+
+ *program = std::move(sanitized_result.program);
gen_ = std::make_unique<GeneratorImpl>(program.get(), version);
return *gen_;
}
diff --git a/src/tint/writer/spirv/binary_writer.h b/src/tint/writer/spirv/binary_writer.h
index 89789fd..a071d7c 100644
--- a/src/tint/writer/spirv/binary_writer.h
+++ b/src/tint/writer/spirv/binary_writer.h
@@ -45,6 +45,9 @@
/// @returns the assembled SPIR-V
const std::vector<uint32_t>& result() const { return out_; }
+ /// @returns the assembled SPIR-V
+ std::vector<uint32_t>& result() { return out_; }
+
private:
void process_instruction(const Instruction& inst);
void process_op(const Operand& op);
diff --git a/src/tint/writer/spirv/builder.cc b/src/tint/writer/spirv/builder.cc
index 436f977..c033344 100644
--- a/src/tint/writer/spirv/builder.cc
+++ b/src/tint/writer/spirv/builder.cc
@@ -41,26 +41,10 @@
#include "src/tint/sem/type_conversion.h"
#include "src/tint/sem/variable.h"
#include "src/tint/sem/vector_type.h"
-#include "src/tint/transform/add_empty_entry_point.h"
#include "src/tint/transform/add_spirv_block_attribute.h"
-#include "src/tint/transform/builtin_polyfill.h"
-#include "src/tint/transform/canonicalize_entry_point_io.h"
-#include "src/tint/transform/expand_compound_assignment.h"
-#include "src/tint/transform/fold_constants.h"
-#include "src/tint/transform/for_loop_to_loop.h"
-#include "src/tint/transform/manager.h"
-#include "src/tint/transform/promote_side_effects_to_decl.h"
-#include "src/tint/transform/remove_unreachable_statements.h"
-#include "src/tint/transform/simplify_pointers.h"
-#include "src/tint/transform/unshadow.h"
-#include "src/tint/transform/unwind_discard_functions.h"
-#include "src/tint/transform/var_for_dynamic_index.h"
-#include "src/tint/transform/vectorize_scalar_matrix_constructors.h"
-#include "src/tint/transform/zero_init_workgroup_memory.h"
#include "src/tint/utils/defer.h"
#include "src/tint/utils/map.h"
#include "src/tint/writer/append_vector.h"
-#include "src/tint/writer/generate_external_texture_bindings.h"
namespace tint::writer::spirv {
namespace {
@@ -255,61 +239,6 @@
} // namespace
-SanitizedResult Sanitize(const Program* in, const Options& options) {
- transform::Manager manager;
- transform::DataMap data;
-
- { // Builtin polyfills
- transform::BuiltinPolyfill::Builtins polyfills;
- polyfills.count_leading_zeros = true;
- polyfills.count_trailing_zeros = true;
- polyfills.extract_bits =
- transform::BuiltinPolyfill::Level::kClampParameters;
- polyfills.first_leading_bit = true;
- polyfills.first_trailing_bit = true;
- polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
- data.Add<transform::BuiltinPolyfill::Config>(polyfills);
- manager.Add<transform::BuiltinPolyfill>();
- }
-
- if (options.generate_external_texture_bindings) {
- auto new_bindings_map = GenerateExternalTextureBindings(in);
- data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
- new_bindings_map);
- }
- manager.Add<transform::MultiplanarExternalTexture>();
-
- manager.Add<transform::Unshadow>();
- bool disable_workgroup_init_in_sanitizer =
- options.disable_workgroup_init ||
- options.use_zero_initialize_workgroup_memory_extension;
- if (!disable_workgroup_init_in_sanitizer) {
- manager.Add<transform::ZeroInitWorkgroupMemory>();
- }
- manager.Add<transform::RemoveUnreachableStatements>();
- manager.Add<transform::ExpandCompoundAssignment>();
- manager.Add<transform::PromoteSideEffectsToDecl>();
- manager.Add<transform::UnwindDiscardFunctions>();
- manager.Add<transform::SimplifyPointers>(); // Required for arrayLength()
- manager.Add<transform::FoldConstants>();
- manager.Add<transform::VectorizeScalarMatrixConstructors>();
- manager.Add<transform::ForLoopToLoop>(); // Must come after
- // ZeroInitWorkgroupMemory
- manager.Add<transform::CanonicalizeEntryPointIO>();
- manager.Add<transform::AddEmptyEntryPoint>();
- manager.Add<transform::AddSpirvBlockAttribute>();
- manager.Add<transform::VarForDynamicIndex>();
-
- data.Add<transform::CanonicalizeEntryPointIO::Config>(
- transform::CanonicalizeEntryPointIO::Config(
- transform::CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF,
- options.emit_vertex_point_size));
-
- SanitizedResult result;
- result.program = std::move(manager.Run(in, data).program);
- return result;
-}
-
Builder::AccessorInfo::AccessorInfo() : source_id(0), source_type(nullptr) {}
Builder::AccessorInfo::~AccessorInfo() {}
diff --git a/src/tint/writer/spirv/builder.h b/src/tint/writer/spirv/builder.h
index c36efcc..7ad9cf2 100644
--- a/src/tint/writer/spirv/builder.h
+++ b/src/tint/writer/spirv/builder.h
@@ -38,7 +38,6 @@
#include "src/tint/sem/builtin.h"
#include "src/tint/sem/storage_texture_type.h"
#include "src/tint/writer/spirv/function.h"
-#include "src/tint/writer/spirv/generator.h"
#include "src/tint/writer/spirv/scalar_constant.h"
// Forward declarations
@@ -51,17 +50,6 @@
namespace tint::writer::spirv {
-/// The result of sanitizing a program for generation.
-struct SanitizedResult {
- /// The sanitized program.
- Program program;
-};
-
-/// Sanitize a program in preparation for generating SPIR-V.
-/// @program The program to sanitize
-/// @param options The SPIR-V generator options.
-SanitizedResult Sanitize(const Program* program, const Options& options);
-
/// Builder class to create SPIR-V instructions from a module.
class Builder {
public:
diff --git a/src/tint/writer/spirv/generator.cc b/src/tint/writer/spirv/generator.cc
index 938f8f7..04f6e15 100644
--- a/src/tint/writer/spirv/generator.cc
+++ b/src/tint/writer/spirv/generator.cc
@@ -14,7 +14,9 @@
#include "src/tint/writer/spirv/generator.h"
-#include "src/tint/writer/spirv/binary_writer.h"
+#include <utility>
+
+#include "src/tint/writer/spirv/generator_impl.h"
namespace tint::writer::spirv {
@@ -37,20 +39,12 @@
bool zero_initialize_workgroup_memory =
!options.disable_workgroup_init &&
options.use_zero_initialize_workgroup_memory_extension;
- auto builder = std::make_unique<Builder>(&sanitized_result.program,
- zero_initialize_workgroup_memory);
- auto writer = std::make_unique<BinaryWriter>();
- if (!builder->Build()) {
- result.success = false;
- result.error = builder->error();
- return result;
- }
- writer->WriteHeader(builder->id_bound());
- writer->WriteBuilder(builder.get());
-
- result.success = true;
- result.spirv = writer->result();
+ auto impl = std::make_unique<GeneratorImpl>(&sanitized_result.program,
+ zero_initialize_workgroup_memory);
+ result.success = impl->Generate();
+ result.error = impl->error();
+ result.spirv = std::move(impl->result());
return result;
}
diff --git a/src/tint/writer/spirv/generator_impl.cc b/src/tint/writer/spirv/generator_impl.cc
new file mode 100644
index 0000000..3318808
--- /dev/null
+++ b/src/tint/writer/spirv/generator_impl.cc
@@ -0,0 +1,120 @@
+// Copyright 2022 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/writer/spirv/generator_impl.h"
+
+#include <utility>
+#include <vector>
+
+#include "src/tint/transform/add_empty_entry_point.h"
+#include "src/tint/transform/add_spirv_block_attribute.h"
+#include "src/tint/transform/builtin_polyfill.h"
+#include "src/tint/transform/canonicalize_entry_point_io.h"
+#include "src/tint/transform/expand_compound_assignment.h"
+#include "src/tint/transform/fold_constants.h"
+#include "src/tint/transform/for_loop_to_loop.h"
+#include "src/tint/transform/manager.h"
+#include "src/tint/transform/promote_side_effects_to_decl.h"
+#include "src/tint/transform/remove_unreachable_statements.h"
+#include "src/tint/transform/simplify_pointers.h"
+#include "src/tint/transform/unshadow.h"
+#include "src/tint/transform/unwind_discard_functions.h"
+#include "src/tint/transform/var_for_dynamic_index.h"
+#include "src/tint/transform/vectorize_scalar_matrix_constructors.h"
+#include "src/tint/transform/zero_init_workgroup_memory.h"
+#include "src/tint/writer/generate_external_texture_bindings.h"
+
+namespace tint::writer::spirv {
+
+SanitizedResult Sanitize(const Program* in, const Options& options) {
+ transform::Manager manager;
+ transform::DataMap data;
+
+ { // Builtin polyfills
+ transform::BuiltinPolyfill::Builtins polyfills;
+ polyfills.count_leading_zeros = true;
+ polyfills.count_trailing_zeros = true;
+ polyfills.extract_bits =
+ transform::BuiltinPolyfill::Level::kClampParameters;
+ polyfills.first_leading_bit = true;
+ polyfills.first_trailing_bit = true;
+ polyfills.insert_bits = transform::BuiltinPolyfill::Level::kClampParameters;
+ data.Add<transform::BuiltinPolyfill::Config>(polyfills);
+ manager.Add<transform::BuiltinPolyfill>();
+ }
+
+ if (options.generate_external_texture_bindings) {
+ auto new_bindings_map = GenerateExternalTextureBindings(in);
+ data.Add<transform::MultiplanarExternalTexture::NewBindingPoints>(
+ new_bindings_map);
+ }
+ manager.Add<transform::MultiplanarExternalTexture>();
+
+ manager.Add<transform::Unshadow>();
+ bool disable_workgroup_init_in_sanitizer =
+ options.disable_workgroup_init ||
+ options.use_zero_initialize_workgroup_memory_extension;
+ if (!disable_workgroup_init_in_sanitizer) {
+ manager.Add<transform::ZeroInitWorkgroupMemory>();
+ }
+ manager.Add<transform::RemoveUnreachableStatements>();
+ manager.Add<transform::ExpandCompoundAssignment>();
+ manager.Add<transform::PromoteSideEffectsToDecl>();
+ manager.Add<transform::UnwindDiscardFunctions>();
+ manager.Add<transform::SimplifyPointers>(); // Required for arrayLength()
+ manager.Add<transform::FoldConstants>();
+ manager.Add<transform::VectorizeScalarMatrixConstructors>();
+ manager.Add<transform::ForLoopToLoop>(); // Must come after
+ // ZeroInitWorkgroupMemory
+ manager.Add<transform::CanonicalizeEntryPointIO>();
+ manager.Add<transform::AddEmptyEntryPoint>();
+ manager.Add<transform::AddSpirvBlockAttribute>();
+ manager.Add<transform::VarForDynamicIndex>();
+
+ data.Add<transform::CanonicalizeEntryPointIO::Config>(
+ transform::CanonicalizeEntryPointIO::Config(
+ transform::CanonicalizeEntryPointIO::ShaderStyle::kSpirv, 0xFFFFFFFF,
+ options.emit_vertex_point_size));
+
+ SanitizedResult result;
+ result.program = std::move(manager.Run(in, data).program);
+ return result;
+}
+
+GeneratorImpl::GeneratorImpl(const Program* program,
+ bool zero_initialize_workgroup_memory)
+ : builder_(program, zero_initialize_workgroup_memory) {}
+
+bool GeneratorImpl::Generate() {
+ if (builder_.Build()) {
+ writer_.WriteHeader(builder_.id_bound());
+ writer_.WriteBuilder(&builder_);
+ return true;
+ }
+ return false;
+}
+
+const std::vector<uint32_t>& GeneratorImpl::result() const {
+ return writer_.result();
+}
+
+std::vector<uint32_t>& GeneratorImpl::result() {
+ return writer_.result();
+}
+
+std::string GeneratorImpl::error() const {
+ return builder_.error();
+}
+
+} // namespace tint::writer::spirv
diff --git a/src/tint/writer/spirv/generator_impl.h b/src/tint/writer/spirv/generator_impl.h
new file mode 100644
index 0000000..2e02af1
--- /dev/null
+++ b/src/tint/writer/spirv/generator_impl.h
@@ -0,0 +1,67 @@
+// Copyright 2022 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_WRITER_SPIRV_GENERATOR_IMPL_H_
+#define SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_H_
+
+#include <string>
+#include <vector>
+
+#include "src/tint/program.h"
+#include "src/tint/writer/spirv/binary_writer.h"
+#include "src/tint/writer/spirv/builder.h"
+#include "src/tint/writer/spirv/generator.h"
+
+namespace tint::writer::spirv {
+
+/// The result of sanitizing a program for generation.
+struct SanitizedResult {
+ /// The sanitized program.
+ Program program;
+};
+
+/// Sanitize a program in preparation for generating SPIR-V.
+/// @program The program to sanitize
+/// @param options The SPIR-V generator options.
+SanitizedResult Sanitize(const Program* program, const Options& options);
+
+/// Implementation class for SPIR-V generator
+class GeneratorImpl {
+ public:
+ /// Constructor
+ /// @param program the program to generate
+ /// @param zero_initialize_workgroup_memory `true` to initialize all the
+ /// variables in the Workgroup storage class with OpConstantNull
+ GeneratorImpl(const Program* program, bool zero_initialize_workgroup_memory);
+
+ /// @returns true on successful generation; false otherwise
+ bool Generate();
+
+ /// @returns the result data
+ const std::vector<uint32_t>& result() const;
+
+ /// @returns the result data
+ std::vector<uint32_t>& result();
+
+ /// @returns the error
+ std::string error() const;
+
+ private:
+ Builder builder_;
+ BinaryWriter writer_;
+};
+
+} // namespace tint::writer::spirv
+
+#endif // SRC_TINT_WRITER_SPIRV_GENERATOR_IMPL_H_
diff --git a/src/tint/writer/spirv/test_helper.h b/src/tint/writer/spirv/test_helper.h
index ecf2a21..c8052b2 100644
--- a/src/tint/writer/spirv/test_helper.h
+++ b/src/tint/writer/spirv/test_helper.h
@@ -22,6 +22,7 @@
#include "gtest/gtest.h"
#include "spirv-tools/libspirv.hpp"
#include "src/tint/writer/spirv/binary_writer.h"
+#include "src/tint/writer/spirv/generator_impl.h"
namespace tint::writer::spirv {
diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn
index d2ca288..bb217df 100644
--- a/test/tint/BUILD.gn
+++ b/test/tint/BUILD.gn
@@ -1,4 +1,4 @@
-# Copyright 2021 The Tint Authors
+# Copyright 2022 The Dawn & Tint Authors
#
# Licensed under the Apache License, Version 2.0 (the "License");
# you may not use this file except in compliance with the License.
@@ -175,6 +175,7 @@
"../../src/tint/ast/id_attribute_test.cc",
"../../src/tint/ast/identifier_expression_test.cc",
"../../src/tint/ast/if_statement_test.cc",
+ "../../src/tint/ast/increment_decrement_statement_test.cc",
"../../src/tint/ast/index_accessor_expression_test.cc",
"../../src/tint/ast/int_literal_expression_test.cc",
"../../src/tint/ast/interpolate_attribute_test.cc",
@@ -252,6 +253,7 @@
"../../src/tint/resolver/entry_point_validation_test.cc",
"../../src/tint/resolver/function_validation_test.cc",
"../../src/tint/resolver/host_shareable_validation_test.cc",
+ "../../src/tint/resolver/increment_decrement_validation_test.cc",
"../../src/tint/resolver/is_host_shareable_test.cc",
"../../src/tint/resolver/is_storeable_test.cc",
"../../src/tint/resolver/pipeline_overridable_constant_test.cc",
@@ -320,11 +322,11 @@
"../../src/tint/transform/decompose_memory_access_test.cc",
"../../src/tint/transform/decompose_strided_array_test.cc",
"../../src/tint/transform/decompose_strided_matrix_test.cc",
+ "../../src/tint/transform/expand_compound_assignment_test.cc",
"../../src/tint/transform/first_index_offset_test.cc",
"../../src/tint/transform/fold_constants_test.cc",
"../../src/tint/transform/fold_trivial_single_use_lets_test.cc",
"../../src/tint/transform/for_loop_to_loop_test.cc",
- "../../src/tint/transform/expand_compound_assignment_test.cc",
"../../src/tint/transform/localize_struct_array_assignment_test.cc",
"../../src/tint/transform/loop_to_for_loop_test.cc",
"../../src/tint/transform/module_scope_var_to_entry_point_param_test.cc",
@@ -497,6 +499,7 @@
"../../src/tint/reader/wgsl/parser_impl_global_variable_decl_test.cc",
"../../src/tint/reader/wgsl/parser_impl_if_stmt_test.cc",
"../../src/tint/reader/wgsl/parser_impl_inclusive_or_expression_test.cc",
+ "../../src/tint/reader/wgsl/parser_impl_increment_decrement_stmt_test.cc",
"../../src/tint/reader/wgsl/parser_impl_logical_and_expression_test.cc",
"../../src/tint/reader/wgsl/parser_impl_logical_or_expression_test.cc",
"../../src/tint/reader/wgsl/parser_impl_loop_stmt_test.cc",
@@ -664,7 +667,6 @@
tint_unittests_source_set("tint_unittests_glsl_writer_src") {
sources = [
- "../../src/tint/transform/glsl_test.cc",
"../../src/tint/writer/glsl/generator_impl_array_accessor_test.cc",
"../../src/tint/writer/glsl/generator_impl_assign_test.cc",
"../../src/tint/writer/glsl/generator_impl_binary_test.cc",
