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>

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

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

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

src/tint/ast/enable.h

@@ -36,6 +36,8 @@ class Enable : public Castable<Enable, Node> {
         /// "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"

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

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_

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

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<transform::DisableUniformityAnalysis>();
+
     {  // 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()
+        auto* subtype =
-                        : storage ? storage->type()
+            sampled ? sampled->type() : storage ? storage->type() : ms ? ms->type() : nullptr;
-                        : ms      ? ms->type()
-                                  : nullptr;
         if (!subtype || subtype->Is<sem::F32>()) {
         } else if (subtype->Is<sem::I32>()) {
             out << "i";

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<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 @@ bool GeneratorImpl::Generate() {
         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 @@ 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;

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<transform::DisableUniformityAnalysis>();
+
     {  // Builtin polyfills
         transform::BuiltinPolyfill::Builtins polyfills;
         polyfills.extract_bits = transform::BuiltinPolyfill::Level::kClampParameters;

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<transform::DisableUniformityAnalysis>();
+
     {  // Builtin polyfills
         transform::BuiltinPolyfill::Builtins polyfills;
         polyfills.count_leading_zeros = true;

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