From 791b4351d1860e1d447f182619e8bd1023dc47eb Mon Sep 17 00:00:00 2001 From: James Price Date: Wed, 11 May 2022 13:50:33 +0000 Subject: [PATCH] tint: Add transform to disable uniformity analysis This is done via a new extension, which in the future could also be used by shader authors as an escape hatch while we are still refining the analysis. The transform is run by the sanitizers for all of the non-WGSL backends. Bug: tint:880 Change-Id: Ibe90d7437d34c741a91eda65dff6d21d8469b9c7 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/88464 Kokoro: Kokoro Reviewed-by: Ben Clayton --- src/tint/BUILD.gn | 2 + src/tint/CMakeLists.txt | 3 + src/tint/ast/enable.cc | 6 ++ src/tint/ast/enable.h | 2 + .../transform/disable_uniformity_analysis.cc | 40 ++++++++++ .../transform/disable_uniformity_analysis.h | 47 ++++++++++++ .../disable_uniformity_analysis_test.cc | 73 +++++++++++++++++++ src/tint/writer/glsl/generator_impl.cc | 9 ++- src/tint/writer/hlsl/generator_impl.cc | 12 +-- src/tint/writer/msl/generator_impl.cc | 3 + src/tint/writer/spirv/generator_impl.cc | 3 + test/tint/BUILD.gn | 1 + 12 files changed, 192 insertions(+), 9 deletions(-) create mode 100644 src/tint/transform/disable_uniformity_analysis.cc create mode 100644 src/tint/transform/disable_uniformity_analysis.h create mode 100644 src/tint/transform/disable_uniformity_analysis_test.cc diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn index 259a26b51f..843f7fdd45 100644 --- a/src/tint/BUILD.gn +++ b/src/tint/BUILD.gn @@ -451,6 +451,8 @@ libtint_source_set("libtint_core_all_src") { "transform/decompose_strided_array.h", "transform/decompose_strided_matrix.cc", "transform/decompose_strided_matrix.h", + "transform/disable_uniformity_analysis.cc", + "transform/disable_uniformity_analysis.h", "transform/expand_compound_assignment.cc", "transform/expand_compound_assignment.h", "transform/first_index_offset.cc", diff --git a/src/tint/CMakeLists.txt b/src/tint/CMakeLists.txt index 682817f7bd..9635bef3f4 100644 --- a/src/tint/CMakeLists.txt +++ b/src/tint/CMakeLists.txt @@ -333,6 +333,8 @@ set(TINT_LIB_SRCS transform/decompose_strided_array.h transform/decompose_strided_matrix.cc transform/decompose_strided_matrix.h + transform/disable_uniformity_analysis.cc + transform/disable_uniformity_analysis.h transform/first_index_offset.cc transform/first_index_offset.h transform/fold_constants.cc @@ -1044,6 +1046,7 @@ if(TINT_BUILD_TESTS) transform/decompose_memory_access_test.cc transform/decompose_strided_array_test.cc transform/decompose_strided_matrix_test.cc + transform/disable_uniformity_analysis_test.cc transform/expand_compound_assignment_test.cc transform/first_index_offset_test.cc transform/fold_constants_test.cc diff --git a/src/tint/ast/enable.cc b/src/tint/ast/enable.cc index 857e1108c8..acb669f8b2 100644 --- a/src/tint/ast/enable.cc +++ b/src/tint/ast/enable.cc @@ -25,6 +25,9 @@ Enable::ExtensionKind Enable::NameToKind(const std::string& name) { if (name == "chromium_experimental_dp4a") { return Enable::ExtensionKind::kChromiumExperimentalDP4a; } + if (name == "chromium_disable_uniformity_analysis") { + return Enable::ExtensionKind::kChromiumDisableUniformityAnalysis; + } // The reserved internal extension name for testing if (name == "InternalExtensionForTesting") { @@ -38,6 +41,9 @@ std::string Enable::KindToName(ExtensionKind kind) { switch (kind) { case ExtensionKind::kChromiumExperimentalDP4a: return "chromium_experimental_dp4a"; + case ExtensionKind::kChromiumDisableUniformityAnalysis: + return "chromium_disable_uniformity_analysis"; + // The reserved internal extension for testing case ExtensionKind::kInternalExtensionForTesting: return "InternalExtensionForTesting"; diff --git a/src/tint/ast/enable.h b/src/tint/ast/enable.h index 7bcd20e4df..c5d1bdd653 100644 --- a/src/tint/ast/enable.h +++ b/src/tint/ast/enable.h @@ -36,6 +36,8 @@ class Enable : public Castable { /// "chromium_experimental_dp4a". /// See crbug.com/tint/1497 for more details kChromiumExperimentalDP4a, + /// A Chromium-specific extension for disabling uniformity analysis. + kChromiumDisableUniformityAnalysis, /// An internal reserved extension for test, named /// "InternalExtensionForTesting" diff --git a/src/tint/transform/disable_uniformity_analysis.cc b/src/tint/transform/disable_uniformity_analysis.cc new file mode 100644 index 0000000000..c02503128c --- /dev/null +++ b/src/tint/transform/disable_uniformity_analysis.cc @@ -0,0 +1,40 @@ +// 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/transform/disable_uniformity_analysis.h" + +#include + +#include "src/tint/program_builder.h" + +TINT_INSTANTIATE_TYPEINFO(tint::transform::DisableUniformityAnalysis); + +namespace tint::transform { + +DisableUniformityAnalysis::DisableUniformityAnalysis() = default; + +DisableUniformityAnalysis::~DisableUniformityAnalysis() = default; + +bool DisableUniformityAnalysis::ShouldRun(const Program* program, const DataMap&) const { + return !program->AST().Extensions().count( + ast::Enable::ExtensionKind::kChromiumDisableUniformityAnalysis); +} + +void DisableUniformityAnalysis::Run(CloneContext& ctx, const DataMap&, DataMap&) const { + ctx.dst->AST().AddEnable(ctx.dst->create( + ast::Enable::KindToName(ast::Enable::ExtensionKind::kChromiumDisableUniformityAnalysis))); + ctx.Clone(); +} + +} // namespace tint::transform diff --git a/src/tint/transform/disable_uniformity_analysis.h b/src/tint/transform/disable_uniformity_analysis.h new file mode 100644 index 0000000000..3c9fb53743 --- /dev/null +++ b/src/tint/transform/disable_uniformity_analysis.h @@ -0,0 +1,47 @@ +// 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_TRANSFORM_DISABLE_UNIFORMITY_ANALYSIS_H_ +#define SRC_TINT_TRANSFORM_DISABLE_UNIFORMITY_ANALYSIS_H_ + +#include "src/tint/transform/transform.h" + +namespace tint::transform { + +/// Disable uniformity analysis for the program. +class DisableUniformityAnalysis final : public Castable { + public: + /// Constructor + DisableUniformityAnalysis(); + /// Destructor + ~DisableUniformityAnalysis() override; + + /// @param program the program to inspect + /// @param data optional extra transform-specific input data + /// @returns true if this transform should be run for the given program + bool ShouldRun(const Program* program, const DataMap& data = {}) const 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) const override; +}; + +} // namespace tint::transform + +#endif // SRC_TINT_TRANSFORM_DISABLE_UNIFORMITY_ANALYSIS_H_ diff --git a/src/tint/transform/disable_uniformity_analysis_test.cc b/src/tint/transform/disable_uniformity_analysis_test.cc new file mode 100644 index 0000000000..a502442f5c --- /dev/null +++ b/src/tint/transform/disable_uniformity_analysis_test.cc @@ -0,0 +1,73 @@ +// 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/transform/disable_uniformity_analysis.h" + +#include +#include + +#include "src/tint/transform/test_helper.h" + +namespace tint::transform { +namespace { + +using DisableUniformityAnalysisTest = TransformTest; + +TEST_F(DisableUniformityAnalysisTest, ShouldRunEmptyModule) { + auto* src = R"()"; + + EXPECT_TRUE(ShouldRun(src)); +} + +TEST_F(DisableUniformityAnalysisTest, ShouldRunExtensionAlreadyPresent) { + auto* src = R"( +enable chromium_disable_uniformity_analysis; +)"; + + EXPECT_FALSE(ShouldRun(src)); +} + +TEST_F(DisableUniformityAnalysisTest, EmptyModule) { + auto* src = R"()"; + + auto* expect = R"( +enable chromium_disable_uniformity_analysis; +)"; + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +TEST_F(DisableUniformityAnalysisTest, NonEmptyModule) { + auto* src = R"( +@group(0) @binding(0) var global : i32; + +@stage(compute) @workgroup_size(64) +fn main() { + if ((global == 42)) { + workgroupBarrier(); + } +} +)"; + + auto expect = "\nenable chromium_disable_uniformity_analysis;\n" + std::string(src); + + auto got = Run(src); + + EXPECT_EQ(expect, str(got)); +} + +} // namespace +} // namespace tint::transform diff --git a/src/tint/writer/glsl/generator_impl.cc b/src/tint/writer/glsl/generator_impl.cc index f52862b2c6..98066d2efc 100644 --- a/src/tint/writer/glsl/generator_impl.cc +++ b/src/tint/writer/glsl/generator_impl.cc @@ -52,6 +52,7 @@ #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/disable_uniformity_analysis.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" @@ -157,6 +158,8 @@ SanitizedResult Sanitize(const Program* in, transform::Manager manager; transform::DataMap data; + manager.Add(); + { // Builtin polyfills transform::BuiltinPolyfill::Builtins polyfills; polyfills.count_leading_zeros = true; @@ -2595,10 +2598,8 @@ bool GeneratorImpl::EmitType(std::ostream& out, if (storage && storage->access() != ast::Access::kRead) { out << "writeonly "; } - auto* subtype = sampled ? sampled->type() - : storage ? storage->type() - : ms ? ms->type() - : nullptr; + auto* subtype = + sampled ? sampled->type() : storage ? storage->type() : ms ? ms->type() : nullptr; if (!subtype || subtype->Is()) { } else if (subtype->Is()) { out << "i"; diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc index 7ee5337f9f..814d6a0da6 100644 --- a/src/tint/writer/hlsl/generator_impl.cc +++ b/src/tint/writer/hlsl/generator_impl.cc @@ -52,6 +52,7 @@ #include "src/tint/transform/calculate_array_length.h" #include "src/tint/transform/canonicalize_entry_point_io.h" #include "src/tint/transform/decompose_memory_access.h" +#include "src/tint/transform/disable_uniformity_analysis.h" #include "src/tint/transform/expand_compound_assignment.h" #include "src/tint/transform/fold_trivial_single_use_lets.h" #include "src/tint/transform/localize_struct_array_assignment.h" @@ -139,6 +140,8 @@ SanitizedResult Sanitize(const Program* in, const Options& options) { transform::Manager manager; transform::DataMap data; + manager.Add(); + { // Builtin polyfills transform::BuiltinPolyfill::Builtins polyfills; // TODO(crbug.com/tint/1449): Some of these can map to HLSL's `firstbitlow` @@ -246,6 +249,10 @@ bool GeneratorImpl::Generate() { if (decl->Is()) { continue; // Ignore aliases. } + if (decl->Is()) { + // Currently we don't have to do anything for using a extension in HLSL. + continue; + } // Emit a new line between declarations if the type of declaration has // changed, or we're about to emit a function @@ -285,11 +292,6 @@ bool GeneratorImpl::Generate() { } return EmitFunction(func); }, - [&](const ast::Enable*) { - // Currently we don't have to do anything for using a extension in - // HLSL - return true; - }, [&](Default) { TINT_ICE(Writer, diagnostics_) << "unhandled module-scope declaration: " << decl->TypeInfo().name; diff --git a/src/tint/writer/msl/generator_impl.cc b/src/tint/writer/msl/generator_impl.cc index 8f5800c565..8e8dc7aaa3 100644 --- a/src/tint/writer/msl/generator_impl.cc +++ b/src/tint/writer/msl/generator_impl.cc @@ -59,6 +59,7 @@ #include "src/tint/transform/array_length_from_uniform.h" #include "src/tint/transform/builtin_polyfill.h" #include "src/tint/transform/canonicalize_entry_point_io.h" +#include "src/tint/transform/disable_uniformity_analysis.h" #include "src/tint/transform/expand_compound_assignment.h" #include "src/tint/transform/manager.h" #include "src/tint/transform/module_scope_var_to_entry_point_param.h" @@ -121,6 +122,8 @@ SanitizedResult Sanitize(const Program* in, const Options& options) { transform::Manager manager; transform::DataMap data; + manager.Add(); + { // Builtin polyfills transform::BuiltinPolyfill::Builtins polyfills; polyfills.extract_bits = transform::BuiltinPolyfill::Level::kClampParameters; diff --git a/src/tint/writer/spirv/generator_impl.cc b/src/tint/writer/spirv/generator_impl.cc index e466a24241..21216bf467 100644 --- a/src/tint/writer/spirv/generator_impl.cc +++ b/src/tint/writer/spirv/generator_impl.cc @@ -21,6 +21,7 @@ #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/disable_uniformity_analysis.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" @@ -41,6 +42,8 @@ SanitizedResult Sanitize(const Program* in, const Options& options) { transform::Manager manager; transform::DataMap data; + manager.Add(); + { // Builtin polyfills transform::BuiltinPolyfill::Builtins polyfills; polyfills.count_leading_zeros = true; diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn index bed8729b43..784fedb98d 100644 --- a/test/tint/BUILD.gn +++ b/test/tint/BUILD.gn @@ -323,6 +323,7 @@ tint_unittests_source_set("tint_unittests_transform_src") { "../../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/disable_uniformity_analysis_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",