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",