tint/cmd: Add `--rename-all` flag

Renames all symbols Add a test for tint:1725 Bug: tint:1725 Change-Id: Idac45c677d15361d76510068ad756e2f9bffacb0 Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/106880 Auto-Submit: Ben Clayton <bclayton@google.com> Kokoro: Ben Clayton <bclayton@google.com> Reviewed-by: Dan Sinclair <dsinclair@chromium.org> Commit-Queue: Ben Clayton <bclayton@google.com>
2022-11-01 16:12:23 +00:00 · 2022-11-01 16:12:23 +00:00 · c02feff258
parent d336733c83
commit c02feff258
9 changed files with 226 additions and 44 deletions
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@ -85,6 +85,8 @@ struct Options {
    bool emit_single_entry_point = false;
    std::string ep_name;

+    bool rename_all = false;
+
    std::vector<std::string> transforms;

    std::string fxc_path;
@ -131,6 +133,7 @@ ${transforms}
  --xcrun                   -- Path to xcrun executable, used to validate MSL output.
                               When specified, automatically enables MSL validation
  --overrides               -- Override values as IDENTIFIER=VALUE, comma-separated.
+  --rename-all              -- Renames all symbols.
 )";

 Format parse_format(const std::string& fmt) {
@ -451,6 +454,9 @@ bool ParseArgs(const std::vector<std::string>& args, Options* opts) {
                auto parts = split_on_equal(o);
                opts->overrides.insert({parts[0], std::stod(parts[1])});
            }
+        } else if (arg == "--rename-all") {
+            ++i;
+            opts->rename_all = true;
        } else if (arg == "--hlsl-root-constant-binding-point") {
            ++i;
            if (i >= args.size()) {
@ -1284,50 +1290,13 @@ int main(int argc, const char** argv) {
    tint::transform::Manager transform_manager;
    tint::transform::DataMap transform_inputs;

-    // If overrides are provided, add the SubstituteOverride transform.
-    if (!options.overrides.empty()) {
-        for (auto& t : transforms) {
-            if (t.name == std::string("substitute_override")) {
-                if (!t.make(inspector, transform_manager, transform_inputs)) {
-                    return 1;
-                }
-                break;
-            }
-        }
-    }
-
-    for (const auto& name : options.transforms) {
-        // TODO(dsinclair): The vertex pulling transform requires setup code to
-        // be run that needs user input. Should we find a way to support that here
-        // maybe through a provided file?
-
-        bool found = false;
-        for (auto& t : transforms) {
-            if (t.name == name) {
-                if (!t.make(inspector, transform_manager, transform_inputs)) {
-                    return 1;
-                }
-                found = true;
-                break;
-            }
-        }
-        if (!found) {
-            std::cerr << "Unknown transform: " << name << std::endl;
-            std::cerr << "Available transforms: " << std::endl << transform_names();
-            return 1;
-        }
-    }
-
-    if (options.emit_single_entry_point) {
-        transform_manager.append(std::make_unique<tint::transform::SingleEntryPoint>());
-        transform_inputs.Add<tint::transform::SingleEntryPoint::Config>(options.ep_name);
-    }
-
+    // Renaming must always come first
    switch (options.format) {
        case Format::kMsl: {
 #if TINT_BUILD_MSL_WRITER
            transform_inputs.Add<tint::transform::Renamer::Config>(
-                tint::transform::Renamer::Target::kMslKeywords,
+                options.rename_all ? tint::transform::Renamer::Target::kAll
+                                   : tint::transform::Renamer::Target::kMslKeywords,
                /* preserve_unicode */ false);
            transform_manager.Add<tint::transform::Renamer>();
 #endif  // TINT_BUILD_MSL_WRITER
@ -1335,20 +1304,63 @@ int main(int argc, const char** argv) {
        }
 #if TINT_BUILD_GLSL_WRITER
        case Format::kGlsl: {
+            transform_inputs.Add<tint::transform::Renamer::Config>(
+                options.rename_all ? tint::transform::Renamer::Target::kAll
+                                   : tint::transform::Renamer::Target::kGlslKeywords,
+                /* preserve_unicode */ false);
+            transform_manager.Add<tint::transform::Renamer>();
            break;
        }
 #endif  // TINT_BUILD_GLSL_WRITER
        case Format::kHlsl: {
 #if TINT_BUILD_HLSL_WRITER
            transform_inputs.Add<tint::transform::Renamer::Config>(
-                tint::transform::Renamer::Target::kHlslKeywords,
+                options.rename_all ? tint::transform::Renamer::Target::kAll
+                                   : tint::transform::Renamer::Target::kHlslKeywords,
                /* preserve_unicode */ false);
            transform_manager.Add<tint::transform::Renamer>();
 #endif  // TINT_BUILD_HLSL_WRITER
            break;
        }
-        default:
+        default: {
+            if (options.rename_all) {
+                transform_manager.Add<tint::transform::Renamer>();
+            }
            break;
+        }
+    }
+
+    auto enable_transform = [&](std::string_view name) {
+        for (auto& t : transforms) {
+            if (t.name == name) {
+                return t.make(inspector, transform_manager, transform_inputs);
+            }
+        }
+
+        std::cerr << "Unknown transform: " << name << std::endl;
+        std::cerr << "Available transforms: " << std::endl << transform_names();
+        return false;
+    };
+
+    // If overrides are provided, add the SubstituteOverride transform.
+    if (!options.overrides.empty()) {
+        if (!enable_transform("substitute_override")) {
+            return 1;
+        }
+    }
+
+    for (const auto& name : options.transforms) {
+        // TODO(dsinclair): The vertex pulling transform requires setup code to
+        // be run that needs user input. Should we find a way to support that here
+        // maybe through a provided file?
+        if (!enable_transform(name)) {
+            return 1;
+        }
+    }
+
+    if (options.emit_single_entry_point) {
+        transform_manager.append(std::make_unique<tint::transform::SingleEntryPoint>());
+        transform_inputs.Add<tint::transform::SingleEntryPoint::Config>(options.ep_name);
    }

    auto out = transform_manager.Run(program.get(), std::move(transform_inputs));
--- a/test/tint/bug/chromium/1236161.wgsl.expected.glsl
+++ b/test/tint/bug/chromium/1236161.wgsl.expected.glsl
@ -16,7 +16,7 @@ layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
 void unused_entry_point() {
  return;
 }
-void tint_symbol() {
-  float tint_symbol_1 = tint_modf(1.0f).whole;
+void i() {
+  float s = tint_modf(1.0f).whole;
 }

--- a/test/tint/bug/tint/1725.wgsl
+++ b/test/tint/bug/tint/1725.wgsl
@ -0,0 +1,10 @@
+// flags: --transform robustness --rename-all
+@group(0) @binding(0) var<storage> data : array<u32>;
+
+@compute @workgroup_size(1)
+fn main(@builtin(local_invocation_index) local_invocation_index : u32) {
+    let min = 0;
+    let max = 0;
+    let arrayLength = 0;
+    let x = data[local_invocation_index];
+}
--- a/test/tint/bug/tint/1725.wgsl.expected.dxc.hlsl
+++ b/test/tint/bug/tint/1725.wgsl.expected.dxc.hlsl
@ -0,0 +1,21 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+
+struct tint_symbol_8 {
+  uint tint_symbol_2 : SV_GroupIndex;
+};
+
+void tint_symbol_1_inner(uint tint_symbol_2) {
+  uint tint_symbol_11 = 0u;
+  tint_symbol.GetDimensions(tint_symbol_11);
+  const uint tint_symbol_12 = (tint_symbol_11 / 4u);
+  const int tint_symbol_3 = 0;
+  const int tint_symbol_4 = 0;
+  const int tint_symbol_5 = 0;
+  const uint tint_symbol_6 = tint_symbol.Load((4u * min(tint_symbol_2, (tint_symbol_12 - 1u))));
+}
+
+[numthreads(1, 1, 1)]
+void tint_symbol_1(tint_symbol_8 tint_symbol_7) {
+  tint_symbol_1_inner(tint_symbol_7.tint_symbol_2);
+  return;
+}
--- a/test/tint/bug/tint/1725.wgsl.expected.fxc.hlsl
+++ b/test/tint/bug/tint/1725.wgsl.expected.fxc.hlsl
@ -0,0 +1,21 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+
+struct tint_symbol_8 {
+  uint tint_symbol_2 : SV_GroupIndex;
+};
+
+void tint_symbol_1_inner(uint tint_symbol_2) {
+  uint tint_symbol_11 = 0u;
+  tint_symbol.GetDimensions(tint_symbol_11);
+  const uint tint_symbol_12 = (tint_symbol_11 / 4u);
+  const int tint_symbol_3 = 0;
+  const int tint_symbol_4 = 0;
+  const int tint_symbol_5 = 0;
+  const uint tint_symbol_6 = tint_symbol.Load((4u * min(tint_symbol_2, (tint_symbol_12 - 1u))));
+}
+
+[numthreads(1, 1, 1)]
+void tint_symbol_1(tint_symbol_8 tint_symbol_7) {
+  tint_symbol_1_inner(tint_symbol_7.tint_symbol_2);
+  return;
+}
--- a/test/tint/bug/tint/1725.wgsl.expected.glsl
+++ b/test/tint/bug/tint/1725.wgsl.expected.glsl
@ -0,0 +1,18 @@
+#version 310 es
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+  uint inner[];
+} tint_symbol;
+
+void tint_symbol_1(uint tint_symbol_2) {
+  int tint_symbol_3 = 0;
+  int tint_symbol_4 = 0;
+  int tint_symbol_5 = 0;
+  uint tint_symbol_6 = tint_symbol.inner[min(tint_symbol_2, (uint(tint_symbol.inner.length()) - 1u))];
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+  tint_symbol_1(gl_LocalInvocationIndex);
+  return;
+}
--- a/test/tint/bug/tint/1725.wgsl.expected.msl
+++ b/test/tint/bug/tint/1725.wgsl.expected.msl
@ -0,0 +1,36 @@
+#include <metal_stdlib>
+
+using namespace metal;
+
+template<typename T, size_t N>
+struct tint_array {
+    const constant T& operator[](size_t i) const constant { return elements[i]; }
+    device T& operator[](size_t i) device { return elements[i]; }
+    const device T& operator[](size_t i) const device { return elements[i]; }
+    thread T& operator[](size_t i) thread { return elements[i]; }
+    const thread T& operator[](size_t i) const thread { return elements[i]; }
+    threadgroup T& operator[](size_t i) threadgroup { return elements[i]; }
+    const threadgroup T& operator[](size_t i) const threadgroup { return elements[i]; }
+    T elements[N];
+};
+
+struct tint_symbol_12 {
+  /* 0x0000 */ tint_array<uint, 1> arr;
+};
+
+struct tint_symbol_7 {
+  /* 0x0000 */ tint_array<uint4, 1> buffer_size;
+};
+
+void tint_symbol_1_inner(uint tint_symbol_2, const device tint_array<uint, 1>* const tint_symbol_9, const constant tint_symbol_7* const tint_symbol_10) {
+  int const tint_symbol_3 = 0;
+  int const tint_symbol_4 = 0;
+  int const tint_symbol_5 = 0;
+  uint const tint_symbol_6 = (*(tint_symbol_9))[min(tint_symbol_2, (((*(tint_symbol_10)).buffer_size[0u][0u] / 4u) - 1u))];
+}
+
+kernel void tint_symbol_1(const device tint_symbol_12* tint_symbol_11 [[buffer(0)]], const constant tint_symbol_7* tint_symbol_13 [[buffer(30)]], uint tint_symbol_2 [[thread_index_in_threadgroup]]) {
+  tint_symbol_1_inner(tint_symbol_2, &((*(tint_symbol_11)).arr), tint_symbol_13);
+  return;
+}
+
--- a/test/tint/bug/tint/1725.wgsl.expected.spvasm
+++ b/test/tint/bug/tint/1725.wgsl.expected.spvasm
@ -0,0 +1,55 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 29
+; Schema: 0
+               OpCapability Shader
+         %17 = OpExtInstImport "GLSL.std.450"
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %tint_symbol_1 "tint_symbol_1" %tint_symbol_2_1
+               OpExecutionMode %tint_symbol_1 LocalSize 1 1 1
+               OpName %tint_symbol_2_1 "tint_symbol_2_1"
+               OpName %tint_symbol_block "tint_symbol_block"
+               OpMemberName %tint_symbol_block 0 "inner"
+               OpName %tint_symbol "tint_symbol"
+               OpName %tint_symbol_1_inner "tint_symbol_1_inner"
+               OpName %tint_symbol_2 "tint_symbol_2"
+               OpName %tint_symbol_1 "tint_symbol_1"
+               OpDecorate %tint_symbol_2_1 BuiltIn LocalInvocationIndex
+               OpDecorate %tint_symbol_block Block
+               OpMemberDecorate %tint_symbol_block 0 Offset 0
+               OpDecorate %_runtimearr_uint ArrayStride 4
+               OpDecorate %tint_symbol NonWritable
+               OpDecorate %tint_symbol DescriptorSet 0
+               OpDecorate %tint_symbol Binding 0
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol_2_1 = OpVariable %_ptr_Input_uint Input
+%_runtimearr_uint = OpTypeRuntimeArray %uint
+%tint_symbol_block = OpTypeStruct %_runtimearr_uint
+%_ptr_StorageBuffer_tint_symbol_block = OpTypePointer StorageBuffer %tint_symbol_block
+%tint_symbol = OpVariable %_ptr_StorageBuffer_tint_symbol_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void %uint
+        %int = OpTypeInt 32 1
+         %14 = OpConstantNull %int
+     %uint_0 = OpConstant %uint 0
+     %uint_1 = OpConstant %uint 1
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+         %24 = OpTypeFunction %void
+%tint_symbol_1_inner = OpFunction %void None %8
+%tint_symbol_2 = OpFunctionParameter %uint
+         %12 = OpLabel
+         %18 = OpArrayLength %uint %tint_symbol 0
+         %20 = OpISub %uint %18 %uint_1
+         %16 = OpExtInst %uint %17 UMin %tint_symbol_2 %20
+         %22 = OpAccessChain %_ptr_StorageBuffer_uint %tint_symbol %uint_0 %16
+         %23 = OpLoad %uint %22
+               OpReturn
+               OpFunctionEnd
+%tint_symbol_1 = OpFunction %void None %24
+         %26 = OpLabel
+         %28 = OpLoad %uint %tint_symbol_2_1
+         %27 = OpFunctionCall %void %tint_symbol_1_inner %28
+               OpReturn
+               OpFunctionEnd
--- a/test/tint/bug/tint/1725.wgsl.expected.wgsl
+++ b/test/tint/bug/tint/1725.wgsl.expected.wgsl
@ -0,0 +1,9 @@
+@group(0) @binding(0) var<storage> tint_symbol : array<u32>;
+
+@compute @workgroup_size(1)
+fn tint_symbol_1(@builtin(local_invocation_index) tint_symbol_2 : u32) {
+  let tint_symbol_3 = 0;
+  let tint_symbol_4 = 0;
+  let tint_symbol_5 = 0;
+  let tint_symbol_6 = tint_symbol[min(tint_symbol_2, (arrayLength(&(tint_symbol)) - 1u))];
+}