transform: Fix CalculateArrayLength for arrays
The transform was not correctly inserting the intrinsic call after array element types.
Fixed: chromium:1290107
Change-Id: I7199d1846cb98305d789cf0bc362eb5872d9b917
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/78542
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: James Price <jrprice@google.com>
diff --git a/src/transform/calculate_array_length.cc b/src/transform/calculate_array_length.cc
index c54c3e8..b365f06 100644
--- a/src/transform/calculate_array_length.cc
+++ b/src/transform/calculate_array_length.cc
@@ -119,7 +119,14 @@
ctx.dst->ASTNodes().Create<BufferSizeIntrinsic>(ctx.dst->ID()),
},
ast::DecorationList{});
- if (auto* str = buffer_type->As<sem::Struct>()) {
+ // Insert the intrinsic function after the structure or array structure
+ // element type. TODO(crbug.com/tint/1266): Once we allow out-of-order
+ // declarations, this can be simplified.
+ auto* insert_after = buffer_type;
+ while (auto* arr = insert_after->As<sem::Array>()) {
+ insert_after = arr->ElemType();
+ }
+ if (auto* str = insert_after->As<sem::Struct>()) {
ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(),
func);
} else {
diff --git a/src/transform/calculate_array_length_test.cc b/src/transform/calculate_array_length_test.cc
index cebb276..fcf7001 100644
--- a/src/transform/calculate_array_length_test.cc
+++ b/src/transform/calculate_array_length_test.cc
@@ -65,7 +65,7 @@
EXPECT_TRUE(ShouldRun<CalculateArrayLength>(src));
}
-TEST_F(CalculateArrayLengthTest, Basic) {
+TEST_F(CalculateArrayLengthTest, BasicArray) {
auto* src = R"(
@group(0) @binding(0) var<storage, read> sb : array<i32>;
@@ -135,6 +135,80 @@
EXPECT_EQ(expect, str(got));
}
+TEST_F(CalculateArrayLengthTest, ArrayOfStruct) {
+ auto* src = R"(
+struct S {
+ f : f32;
+}
+
+@group(0) @binding(0) var<storage, read> arr : array<S>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+ let len = arrayLength(&arr);
+}
+)";
+ auto* expect = R"(
+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)
+fn main() {
+ var tint_symbol_1 : u32 = 0u;
+ tint_symbol(arr, &(tint_symbol_1));
+ let tint_symbol_2 : u32 = (tint_symbol_1 / 4u);
+ let len = tint_symbol_2;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, ArrayOfArrayOfStruct) {
+ auto* src = R"(
+struct S {
+ f : f32;
+}
+
+@group(0) @binding(0) var<storage, read> arr : array<array<S, 4>>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+ let len = arrayLength(&arr);
+}
+)";
+ auto* expect = R"(
+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)
+fn main() {
+ var tint_symbol_1 : u32 = 0u;
+ tint_symbol(arr, &(tint_symbol_1));
+ let tint_symbol_2 : u32 = (tint_symbol_1 / 16u);
+ let len = tint_symbol_2;
+}
+)";
+
+ auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+ EXPECT_EQ(expect, str(got));
+}
+
TEST_F(CalculateArrayLengthTest, InSameBlock) {
auto* src = R"(
@group(0) @binding(0) var<storage, read> sb : array<i32>;;
diff --git a/test/bug/chromium/1290107.wgsl b/test/bug/chromium/1290107.wgsl
new file mode 100644
index 0000000..c2e4778
--- /dev/null
+++ b/test/bug/chromium/1290107.wgsl
@@ -0,0 +1,10 @@
+struct S {
+ f : f32;
+}
+
+@group(0) @binding(0) var<storage, read> arr : array<S>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+ let len = arrayLength(&arr);
+}
diff --git a/test/bug/chromium/1290107.wgsl.expected.glsl b/test/bug/chromium/1290107.wgsl.expected.glsl
new file mode 100644
index 0000000..debb813
--- /dev/null
+++ b/test/bug/chromium/1290107.wgsl.expected.glsl
@@ -0,0 +1,32 @@
+SKIP: FAILED
+
+#version 310 es
+precision mediump float;
+
+struct S {
+ float f;
+};
+
+layout(binding = 0) buffer arr_block_1 {
+ S inner[];
+} arr;
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void tint_symbol() {
+ uint tint_symbol_2 = 0u;
+ arr.inner.GetDimensions(tint_symbol_2);
+ uint tint_symbol_3 = (tint_symbol_2 / 4u);
+ uint len = tint_symbol_3;
+ return;
+}
+
+void main() {
+ tint_symbol();
+}
+
+Error parsing GLSL shader:
+ERROR: 0:14: '.' : cannot apply to an array: GetDimensions
+ERROR: 0:14: '' : compilation terminated
+ERROR: 2 compilation errors. No code generated.
+
+
+
diff --git a/test/bug/chromium/1290107.wgsl.expected.hlsl b/test/bug/chromium/1290107.wgsl.expected.hlsl
new file mode 100644
index 0000000..95e8cb7
--- /dev/null
+++ b/test/bug/chromium/1290107.wgsl.expected.hlsl
@@ -0,0 +1,10 @@
+ByteAddressBuffer arr : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ uint tint_symbol_1 = 0u;
+ arr.GetDimensions(tint_symbol_1);
+ const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+ const uint len = tint_symbol_2;
+ return;
+}
diff --git a/test/bug/chromium/1290107.wgsl.expected.msl b/test/bug/chromium/1290107.wgsl.expected.msl
new file mode 100644
index 0000000..4255195
--- /dev/null
+++ b/test/bug/chromium/1290107.wgsl.expected.msl
@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+ /* 0x0000 */ uint4 buffer_size[1];
+};
+struct S {
+ float f;
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+ uint const len = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+ return;
+}
+
diff --git a/test/bug/chromium/1290107.wgsl.expected.spvasm b/test/bug/chromium/1290107.wgsl.expected.spvasm
new file mode 100644
index 0000000..1c835ab
--- /dev/null
+++ b/test/bug/chromium/1290107.wgsl.expected.spvasm
@@ -0,0 +1,36 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %arr_block "arr_block"
+ OpMemberName %arr_block 0 "inner"
+ OpName %S "S"
+ OpMemberName %S 0 "f"
+ OpName %arr "arr"
+ OpName %main "main"
+ OpDecorate %arr_block Block
+ OpMemberDecorate %arr_block 0 Offset 0
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %_runtimearr_S ArrayStride 4
+ OpDecorate %arr NonWritable
+ OpDecorate %arr DescriptorSet 0
+ OpDecorate %arr Binding 0
+ %float = OpTypeFloat 32
+ %S = OpTypeStruct %float
+%_runtimearr_S = OpTypeRuntimeArray %S
+ %arr_block = OpTypeStruct %_runtimearr_S
+%_ptr_StorageBuffer_arr_block = OpTypePointer StorageBuffer %arr_block
+ %arr = OpVariable %_ptr_StorageBuffer_arr_block StorageBuffer
+ %void = OpTypeVoid
+ %7 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %main = OpFunction %void None %7
+ %10 = OpLabel
+ %11 = OpArrayLength %uint %arr 0
+ OpReturn
+ OpFunctionEnd
diff --git a/test/bug/chromium/1290107.wgsl.expected.wgsl b/test/bug/chromium/1290107.wgsl.expected.wgsl
new file mode 100644
index 0000000..4041cef
--- /dev/null
+++ b/test/bug/chromium/1290107.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+struct S {
+ f : f32;
+}
+
+@group(0) @binding(0) var<storage, read> arr : array<S>;
+
+@stage(compute) @workgroup_size(1)
+fn main() {
+ let len = arrayLength(&(arr));
+}
