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>

src/transform/calculate_array_length.cc

@@ -119,7 +119,14 @@ void CalculateArrayLength::Run(CloneContext& ctx,
               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 {

src/transform/calculate_array_length_test.cc

@@ -65,7 +65,7 @@ fn main() {
   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 @@ fn main() {
   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>;;

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);
+}

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.
+
+
+

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;
+}

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;
+}
+

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

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));
+}