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 <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@google.com>
diff --git a/src/tint/BUILD.gn b/src/tint/BUILD.gn
index 259a26b..843f7fd 100644
--- a/src/tint/BUILD.gn
+++ b/src/tint/BUILD.gn
@@ -451,6 +451,8 @@
"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 682817f..9635bef 100644
--- a/src/tint/CMakeLists.txt
+++ b/src/tint/CMakeLists.txt
@@ -333,6 +333,8 @@
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 @@
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 857e110..acb669f 100644
--- a/src/tint/ast/enable.cc
+++ b/src/tint/ast/enable.cc
@@ -25,6 +25,9 @@
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 @@
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 7bcd20e..c5d1bdd 100644
--- a/src/tint/ast/enable.h
+++ b/src/tint/ast/enable.h
@@ -36,6 +36,8 @@
/// "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 0000000..c025031
--- /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 <utility>
+
+#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>(
+ 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 0000000..3c9fb53
--- /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<DisableUniformityAnalysis, Transform> {
+ 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 0000000..a502442
--- /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 <string>
+#include <utility>
+
+#include "src/tint/transform/test_helper.h"
+
+namespace tint::transform {
+namespace {
+
+using DisableUniformityAnalysisTest = TransformTest;
+
+TEST_F(DisableUniformityAnalysisTest, ShouldRunEmptyModule) {
+ auto* src = R"()";
+
+ EXPECT_TRUE(ShouldRun<DisableUniformityAnalysis>(src));
+}
+
+TEST_F(DisableUniformityAnalysisTest, ShouldRunExtensionAlreadyPresent) {
+ auto* src = R"(
+enable chromium_disable_uniformity_analysis;
+)";
+
+ EXPECT_FALSE(ShouldRun<DisableUniformityAnalysis>(src));
+}
+
+TEST_F(DisableUniformityAnalysisTest, EmptyModule) {
+ auto* src = R"()";
+
+ auto* expect = R"(
+enable chromium_disable_uniformity_analysis;
+)";
+
+ auto got = Run<DisableUniformityAnalysis>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(DisableUniformityAnalysisTest, NonEmptyModule) {
+ auto* src = R"(
+@group(0) @binding(0) var<storage, read> 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<DisableUniformityAnalysis>(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 f52862b..98066d2 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 @@
transform::Manager manager;
transform::DataMap data;
+ manager.Add<transform::DisableUniformityAnalysis>();
+
{ // Builtin polyfills
transform::BuiltinPolyfill::Builtins polyfills;
polyfills.count_leading_zeros = true;
@@ -2595,10 +2598,8 @@
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<sem::F32>()) {
} else if (subtype->Is<sem::I32>()) {
out << "i";
diff --git a/src/tint/writer/hlsl/generator_impl.cc b/src/tint/writer/hlsl/generator_impl.cc
index 7ee5337..814d6a0 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 @@
transform::Manager manager;
transform::DataMap data;
+ manager.Add<transform::DisableUniformityAnalysis>();
+
{ // Builtin polyfills
transform::BuiltinPolyfill::Builtins polyfills;
// TODO(crbug.com/tint/1449): Some of these can map to HLSL's `firstbitlow`
@@ -246,6 +249,10 @@
if (decl->Is<ast::Alias>()) {
continue; // Ignore aliases.
}
+ if (decl->Is<ast::Enable>()) {
+ // 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 @@
}
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 8f5800c..8e8dc7a 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 @@
transform::Manager manager;
transform::DataMap data;
+ manager.Add<transform::DisableUniformityAnalysis>();
+
{ // 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 e466a24..21216bf 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 @@
transform::Manager manager;
transform::DataMap data;
+ manager.Add<transform::DisableUniformityAnalysis>();
+
{ // Builtin polyfills
transform::BuiltinPolyfill::Builtins polyfills;
polyfills.count_leading_zeros = true;
diff --git a/test/tint/BUILD.gn b/test/tint/BUILD.gn
index bed8729..784fedb 100644
--- a/test/tint/BUILD.gn
+++ b/test/tint/BUILD.gn
@@ -323,6 +323,7 @@
"../../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",
