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>
diff --git a/src/tint/cmd/main.cc b/src/tint/cmd/main.cc
index f27e643..bf5a993 100644
--- a/src/tint/cmd/main.cc
+++ b/src/tint/cmd/main.cc
@@ -85,6 +85,8 @@
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 @@
--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 @@
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 @@
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 @@
}
#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));
diff --git a/test/tint/bug/chromium/1236161.wgsl.expected.glsl b/test/tint/bug/chromium/1236161.wgsl.expected.glsl
index a7666fb..04734ec 100644
--- a/test/tint/bug/chromium/1236161.wgsl.expected.glsl
+++ b/test/tint/bug/chromium/1236161.wgsl.expected.glsl
@@ -16,7 +16,7 @@
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;
}
diff --git a/test/tint/bug/tint/1725.wgsl b/test/tint/bug/tint/1725.wgsl
new file mode 100644
index 0000000..3c398b7
--- /dev/null
+++ 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];
+}
diff --git a/test/tint/bug/tint/1725.wgsl.expected.dxc.hlsl b/test/tint/bug/tint/1725.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..cea26ec
--- /dev/null
+++ 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;
+}
diff --git a/test/tint/bug/tint/1725.wgsl.expected.fxc.hlsl b/test/tint/bug/tint/1725.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..cea26ec
--- /dev/null
+++ 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;
+}
diff --git a/test/tint/bug/tint/1725.wgsl.expected.glsl b/test/tint/bug/tint/1725.wgsl.expected.glsl
new file mode 100644
index 0000000..3d8ff77
--- /dev/null
+++ 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;
+}
diff --git a/test/tint/bug/tint/1725.wgsl.expected.msl b/test/tint/bug/tint/1725.wgsl.expected.msl
new file mode 100644
index 0000000..a5c54fe
--- /dev/null
+++ 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;
+}
+
diff --git a/test/tint/bug/tint/1725.wgsl.expected.spvasm b/test/tint/bug/tint/1725.wgsl.expected.spvasm
new file mode 100644
index 0000000..7b1e735
--- /dev/null
+++ 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
diff --git a/test/tint/bug/tint/1725.wgsl.expected.wgsl b/test/tint/bug/tint/1725.wgsl.expected.wgsl
new file mode 100644
index 0000000..eefb582
--- /dev/null
+++ 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))];
+}
