resolver: Enable support of out of order declarations
Allow module-scope declarations to be made in any order.
Bug: tint:1266
Change-Id: Ib2607b6c33fad7c83e2c36f85b0a965eac922ec5
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/79769
Reviewed-by: David Neto <dneto@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
diff --git a/src/transform/calculate_array_length_test.cc b/src/transform/calculate_array_length_test.cc
index fcf7001..fa1f274 100644
--- a/src/transform/calculate_array_length_test.cc
+++ b/src/transform/calculate_array_length_test.cc
@@ -111,14 +111,14 @@
)";
auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
struct SB {
x : i32;
arr : array<i32>;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@@ -149,13 +149,13 @@
}
)";
auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<S>, result : ptr<function, u32>)
+
struct S {
f : f32;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<S>, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> arr : array<S>;
@stage(compute) @workgroup_size(1)
@@ -186,13 +186,13 @@
}
)";
auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<array<S, 4u>>, result : ptr<function, u32>)
+
struct S {
f : f32;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<array<S, 4u>>, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> arr : array<array<S, 4>>;
@stage(compute) @workgroup_size(1)
@@ -261,14 +261,14 @@
)";
auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
struct SB {
x : i32;
arr : array<i32>;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@@ -334,15 +334,15 @@
)";
auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
struct SB {
x : i32;
y : f32;
arr : @stride(64) array<i32>;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@@ -381,14 +381,14 @@
)";
auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
struct SB {
x : i32;
arr : array<i32>;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> sb : SB;
@stage(compute) @workgroup_size(1)
@@ -443,6 +443,12 @@
auto* expect = R"(
@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
+
+@internal(intrinsic_buffer_size)
+fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
+
+@internal(intrinsic_buffer_size)
fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
struct SB1 {
@@ -450,17 +456,11 @@
arr1 : array<i32>;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
-
struct SB2 {
x : i32;
arr2 : array<vec4<f32>>;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> sb1 : SB1;
@group(0) @binding(1) var<storage, read> sb2 : SB2;
@@ -512,14 +512,14 @@
auto* expect =
R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
+
struct SB {
x : i32;
arr : array<i32>;
}
-@internal(intrinsic_buffer_size)
-fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB, result : ptr<function, u32>)
-
@group(0) @binding(0) var<storage, read> a : SB;
@group(0) @binding(1) var<storage, read> b : SB;
@@ -544,6 +544,82 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(CalculateArrayLengthTest, OutOfOrder) {
+ auto* src = R"(
+@stage(compute) @workgroup_size(1)
+fn main() {
+ var len1 : u32 = arrayLength(&(sb1.arr1));
+ var len2 : u32 = arrayLength(&(sb2.arr2));
+ var len3 : u32 = arrayLength(&sb3);
+ var x : u32 = (len1 + len2 + len3);
+}
+
+@group(0) @binding(0) var<storage, read> sb1 : SB1;
+
+struct SB1 {
+ x : i32;
+ arr1 : array<i32>;
+};
+
+@group(0) @binding(1) var<storage, read> sb2 : SB2;
+
+struct SB2 {
+ x : i32;
+ arr2 : array<vec4<f32>>;
+};
+
+@group(0) @binding(2) var<storage, read> sb3 : array<i32>;
+)";
+
+ auto* expect = R"(
+@internal(intrinsic_buffer_size)
+fn tint_symbol(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB1, result : ptr<function, u32>)
+
+@internal(intrinsic_buffer_size)
+fn tint_symbol_3(@internal(disable_validation__ignore_constructible_function_parameter) buffer : SB2, result : ptr<function, u32>)
+
+@internal(intrinsic_buffer_size)
+fn tint_symbol_6(@internal(disable_validation__ignore_constructible_function_parameter) buffer : array<i32>, result : ptr<function, u32>)
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+ var tint_symbol_1 : u32 = 0u;
+ tint_symbol(sb1, &(tint_symbol_1));
+ let tint_symbol_2 : u32 = ((tint_symbol_1 - 4u) / 4u);
+ var tint_symbol_4 : u32 = 0u;
+ tint_symbol_3(sb2, &(tint_symbol_4));
+ let tint_symbol_5 : u32 = ((tint_symbol_4 - 16u) / 16u);
+ var tint_symbol_7 : u32 = 0u;
+ tint_symbol_6(sb3, &(tint_symbol_7));
+ let tint_symbol_8 : u32 = (tint_symbol_7 / 4u);
+ var len1 : u32 = tint_symbol_2;
+ var len2 : u32 = tint_symbol_5;
+ var len3 : u32 = tint_symbol_8;
+ var x : u32 = ((len1 + len2) + len3);
+}
+
+@group(0) @binding(0) var<storage, read> sb1 : SB1;
+
+struct SB1 {
+ x : i32;
+ arr1 : array<i32>;
+}
+
+@group(0) @binding(1) var<storage, read> sb2 : SB2;
+
+struct SB2 {
+ x : i32;
+ arr2 : array<vec4<f32>>;
+}
+
+@group(0) @binding(2) var<storage, read> sb3 : array<i32>;
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
} // namespace
} // namespace transform
} // namespace tint
