transform: Handle arrayLength for non-struct buffers

These two transforms previously assumed that the argument to
arrayLength had the form `&struct_var.array_member`. We now also need
to handle the case where it is just `&array_var`.

Bug: tint:1372
Change-Id: I173a84bd32c324445573a295b281a51e291c2ae2
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/76163
Reviewed-by: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>

src/transform/array_length_from_uniform.cc

@@ -59,30 +59,26 @@ static void IterateArrayLengthOnStorageVar(CloneContext& ctx, F&& functor) {
     }
 
     // Get the storage buffer that contains the runtime array.
-    // We assume that the argument to `arrayLength` has the form
+    // Since we require SimplifyPointers, we can assume that the arrayLength()
-    // `&resource.array`, which requires that `SimplifyPointers` have been run
+    // call has one of two forms:
-    // before this transform.
+    //   arrayLength(&struct_var.array_member)
+    //   arrayLength(&array_var)
     auto* param = call_expr->args[0]->As<ast::UnaryOpExpression>();
     if (!param || param->op != ast::UnaryOp::kAddressOf) {
       TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "expected form of arrayLength argument to be "
+          << "expected form of arrayLength argument to be &array_var or "
-             "&resource.array";
+             "&struct_var.array_member";
       break;
     }
-    auto* accessor = param->expr->As<ast::MemberAccessorExpression>();
+    auto* storage_buffer_expr = param->expr;
-    if (!accessor) {
+    if (auto* accessor = param->expr->As<ast::MemberAccessorExpression>()) {
-      TINT_ICE(Transform, ctx.dst->Diagnostics())
+      storage_buffer_expr = accessor->structure;
-          << "expected form of arrayLength argument to be "
-             "&resource.array";
-      break;
     }
-    auto* storage_buffer_expr = accessor->structure;
+    auto* storage_buffer_sem = sem.Get<sem::VariableUser>(storage_buffer_expr);
-    auto* storage_buffer_sem =
-        sem.Get(storage_buffer_expr)->As<sem::VariableUser>();
     if (!storage_buffer_sem) {
       TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "expected form of arrayLength argument to be "
+          << "expected form of arrayLength argument to be &array_var or "
-             "&resource.array";
+             "&struct_var.array_member";
       break;
     }
 
@@ -183,14 +179,25 @@ void ArrayLengthFromUniform::Run(CloneContext& ctx,
         //                total_storage_buffer_size - array_offset
         // array_length = ----------------------------------------
         //                             array_stride
-        auto* storage_buffer_type =
+        const ast::Expression* total_size = total_storage_buffer_size;
-            storage_buffer_sem->Type()->UnwrapRef()->As<sem::Struct>();
+        auto* storage_buffer_type = storage_buffer_sem->Type()->UnwrapRef();
-        auto* array_member_sem = storage_buffer_type->Members().back();
+        const sem::Array* array_type = nullptr;
-        uint32_t array_offset = array_member_sem->Offset();
+        if (auto* str = storage_buffer_type->As<sem::Struct>()) {
-        uint32_t array_stride = array_member_sem->Size();
+          // The variable is a struct, so subtract the byte offset of the array
-        auto* array_length =
+          // member.
-            ctx.dst->Div(ctx.dst->Sub(total_storage_buffer_size, array_offset),
+          auto* array_member_sem = str->Members().back();
-                         array_stride);
+          array_type = array_member_sem->Type()->As<sem::Array>();
+          total_size = ctx.dst->Sub(total_storage_buffer_size,
+                                    array_member_sem->Offset());
+        } else if (auto* arr = storage_buffer_type->As<sem::Array>()) {
+          array_type = arr;
+        } else {
+          TINT_ICE(Transform, ctx.dst->Diagnostics())
+              << "expected form of arrayLength argument to be &array_var or "
+                 "&struct_var.array_member";
+          return;
+        }
+        auto* array_length = ctx.dst->Div(total_size, array_type->Stride());
 
         ctx.Replace(call_expr, array_length);
       });

src/transform/array_length_from_uniform_test.cc

@@ -52,6 +52,44 @@ TEST_F(ArrayLengthFromUniformTest, Error_MissingSimplifyPointers) {
 
 TEST_F(ArrayLengthFromUniformTest, Basic) {
   auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var len : u32 = arrayLength(&sb);
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  buffer_size : array<vec4<u32>, 1u>;
+};
+
+[[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
+
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 4u);
+}
+)";
+
+  ArrayLengthFromUniform::Config cfg({0, 30u});
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
+
+  DataMap data;
+  data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+  EXPECT_EQ(std::unordered_set<uint32_t>({0}),
+            got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
+}
+
+TEST_F(ArrayLengthFromUniformTest, BasicInStruct) {
+  auto* src = R"(
 struct SB {
   x : i32;
   arr : array<i32>;
@@ -100,6 +138,44 @@ fn main() {
 
 TEST_F(ArrayLengthFromUniformTest, WithStride) {
   auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var len : u32 = arrayLength(&sb);
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol {
+  buffer_size : array<vec4<u32>, 1u>;
+};
+
+[[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
+
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var len : u32 = (tint_symbol_1.buffer_size[0u][0u] / 64u);
+}
+)";
+
+  ArrayLengthFromUniform::Config cfg({0, 30u});
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 0}, 0);
+
+  DataMap data;
+  data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
+
+  auto got = Run<Unshadow, SimplifyPointers, ArrayLengthFromUniform>(src, data);
+
+  EXPECT_EQ(expect, str(got));
+  EXPECT_EQ(std::unordered_set<uint32_t>({0}),
+            got.data.Get<ArrayLengthFromUniform::Result>()->used_size_indices);
+}
+
+TEST_F(ArrayLengthFromUniformTest, WithStride_InStruct) {
+  auto* src = R"(
 struct SB {
   x : i32;
   y : f32;
@@ -158,32 +234,24 @@ struct SB2 {
   x : i32;
   arr2 : array<vec4<f32>>;
 };
-struct SB3 {
-  x : i32;
-  arr3 : array<vec4<f32>>;
-};
 struct SB4 {
   x : i32;
   arr4 : array<vec4<f32>>;
 };
-struct SB5 {
-  x : i32;
-  arr5 : array<vec4<f32>>;
-};
 
 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
-[[group(2), binding(2)]] var<storage, read> sb3 : SB3;
+[[group(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = arrayLength(&(sb1.arr1));
   var len2 : u32 = arrayLength(&(sb2.arr2));
-  var len3 : u32 = arrayLength(&(sb3.arr3));
+  var len3 : u32 = arrayLength(&sb3);
   var len4 : u32 = arrayLength(&(sb4.arr4));
-  var len5 : u32 = arrayLength(&(sb5.arr5));
+  var len5 : u32 = arrayLength(&sb5);
   var x : u32 = (len1 + len2 + len3 + len4 + len5);
 }
 )";
@@ -205,38 +273,28 @@ struct SB2 {
   arr2 : array<vec4<f32>>;
 };
 
-struct SB3 {
-  x : i32;
-  arr3 : array<vec4<f32>>;
-};
-
 struct SB4 {
   x : i32;
   arr4 : array<vec4<f32>>;
 };
 
-struct SB5 {
-  x : i32;
-  arr5 : array<vec4<f32>>;
-};
-
 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
 
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
 
-[[group(2), binding(2)]] var<storage, read> sb3 : SB3;
+[[group(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
 
 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
 
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
   var len2 : u32 = ((tint_symbol_1.buffer_size[0u][1u] - 16u) / 16u);
-  var len3 : u32 = ((tint_symbol_1.buffer_size[0u][2u] - 16u) / 16u);
+  var len3 : u32 = (tint_symbol_1.buffer_size[0u][2u] / 16u);
   var len4 : u32 = ((tint_symbol_1.buffer_size[0u][3u] - 16u) / 16u);
-  var len5 : u32 = ((tint_symbol_1.buffer_size[1u][0u] - 16u) / 16u);
+  var len5 : u32 = (tint_symbol_1.buffer_size[1u][0u] / 16u);
   var x : u32 = ((((len1 + len2) + len3) + len4) + len5);
 }
 )";
@@ -268,29 +326,21 @@ struct SB2 {
   x : i32;
   arr2 : array<vec4<f32>>;
 };
-struct SB3 {
-  x : i32;
-  arr3 : array<vec4<f32>>;
-};
 struct SB4 {
   x : i32;
   arr4 : array<vec4<f32>>;
 };
-struct SB5 {
-  x : i32;
-  arr5 : array<vec4<f32>>;
-};
 
 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
-[[group(2), binding(2)]] var<storage, read> sb3 : SB3;
+[[group(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = arrayLength(&(sb1.arr1));
-  var len3 : u32 = arrayLength(&(sb3.arr3));
+  var len3 : u32 = arrayLength(&sb3);
   var x : u32 = (len1 + len3);
 }
 )";
@@ -312,35 +362,25 @@ struct SB2 {
   arr2 : array<vec4<f32>>;
 };
 
-struct SB3 {
-  x : i32;
-  arr3 : array<vec4<f32>>;
-};
-
 struct SB4 {
   x : i32;
   arr4 : array<vec4<f32>>;
 };
 
-struct SB5 {
-  x : i32;
-  arr5 : array<vec4<f32>>;
-};
-
 [[group(0), binding(2)]] var<storage, read> sb1 : SB1;
 
 [[group(1), binding(2)]] var<storage, read> sb2 : SB2;
 
-[[group(2), binding(2)]] var<storage, read> sb3 : SB3;
+[[group(2), binding(2)]] var<storage, read> sb3 : array<vec4<f32>>;
 
 [[group(3), binding(2)]] var<storage, read> sb4 : SB4;
 
-[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+[[group(4), binding(2)]] var<storage, read> sb5 : array<vec4<f32>>;
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
-  var len3 : u32 = ((tint_symbol_1.buffer_size[0u][2u] - 16u) / 16u);
+  var len3 : u32 = (tint_symbol_1.buffer_size[0u][2u] / 16u);
   var x : u32 = (len1 + len3);
 }
 )";

src/transform/calculate_array_length.cc

@@ -42,7 +42,7 @@ namespace {
 /// It is used as a key by the array_length_by_usage map.
 struct ArrayUsage {
   ast::BlockStatement const* const block;
-  sem::Node const* const buffer;
+  sem::Variable const* const buffer;
   bool operator==(const ArrayUsage& rhs) const {
     return block == rhs.block && buffer == rhs.buffer;
   }
@@ -80,12 +80,11 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) {
   // get_buffer_size_intrinsic() emits the function decorated with
   // BufferSizeIntrinsic that is transformed by the HLSL writer into a call to
   // [RW]ByteAddressBuffer.GetDimensions().
-  std::unordered_map<const sem::Struct*, Symbol> buffer_size_intrinsics;
+  std::unordered_map<const sem::Type*, Symbol> buffer_size_intrinsics;
-  auto get_buffer_size_intrinsic = [&](const sem::Struct* buffer_type) {
+  auto get_buffer_size_intrinsic = [&](const sem::Type* buffer_type) {
     return utils::GetOrCreate(buffer_size_intrinsics, buffer_type, [&] {
       auto name = ctx.dst->Sym();
-      auto* buffer_typename =
+      auto* type = CreateASTTypeFor(ctx, buffer_type);
-          ctx.dst->ty.type_name(ctx.Clone(buffer_type->Declaration()->name));
       auto* disable_validation = ctx.dst->Disable(
           ast::DisabledValidation::kIgnoreConstructibleFunctionParameter);
       auto* func = ctx.dst->create<ast::Function>(
@@ -95,7 +94,7 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) {
               // in order for HLSL to emit this as a ByteAddressBuffer.
               ctx.dst->create<ast::Variable>(
                   ctx.dst->Sym("buffer"), ast::StorageClass::kStorage,
-                  ast::Access::kUndefined, buffer_typename, true, nullptr,
+                  ast::Access::kUndefined, type, true, nullptr,
                   ast::DecorationList{disable_validation}),
               ctx.dst->Param("result",
                              ctx.dst->ty.pointer(ctx.dst->ty.u32(),
@@ -106,8 +105,12 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) {
               ctx.dst->ASTNodes().Create<BufferSizeIntrinsic>(ctx.dst->ID()),
           },
           ast::DecorationList{});
-      ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(),
+      if (auto* str = buffer_type->As<sem::Struct>()) {
-                      buffer_type->Declaration(), func);
+        ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(),
+                        func);
+      } else {
+        ctx.InsertFront(ctx.src->AST().GlobalDeclarations(), func);
+      }
       return name;
     });
   };
@@ -123,71 +126,47 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) {
         if (intrinsic->Type() == sem::IntrinsicType::kArrayLength) {
           // We're dealing with an arrayLength() call
 
-          // https://gpuweb.github.io/gpuweb/wgsl/#array-types states:
+          // A runtime-sized array can only appear as the store type of a
-          //
+          // variable, or the last element of a structure (which cannot itself
-          // * The last member of the structure type defining the store type for
+          // be nested). Given that we require SimplifyPointers, we can assume
-          //   a variable in the storage storage class may be a runtime-sized
+          // that the arrayLength() call has one of two forms:
-          //   array.
+          //   arrayLength(&struct_var.array_member)
-          // * A runtime-sized array must not be used as the store type or
+          //   arrayLength(&array_var)
-          //   contained within a store type in any other cases.
-          // * An expression must not evaluate to a runtime-sized array type.
-          //
-          // We can assume that the arrayLength() call has a single argument of
-          // the form: arrayLength(&X.Y) where X is an expression that resolves
-          // to the storage buffer structure, and Y is the runtime sized array.
           auto* arg = call_expr->args[0];
           auto* address_of = arg->As<ast::UnaryOpExpression>();
           if (!address_of || address_of->op != ast::UnaryOp::kAddressOf) {
             TINT_ICE(Transform, ctx.dst->Diagnostics())
-                << "arrayLength() expected pointer to member access, got "
+                << "arrayLength() expected address-of, got "
-                << address_of->TypeInfo().name;
+                << arg->TypeInfo().name;
           }
-          auto* array_expr = address_of->expr;
+          auto* storage_buffer_expr = address_of->expr;
-
+          if (auto* accessor =
-          auto* accessor = array_expr->As<ast::MemberAccessorExpression>();
+                  storage_buffer_expr->As<ast::MemberAccessorExpression>()) {
-          if (!accessor) {
+            storage_buffer_expr = accessor->structure;
+          }
+          auto* storage_buffer_sem =
+              sem.Get<sem::VariableUser>(storage_buffer_expr);
+          if (!storage_buffer_sem) {
             TINT_ICE(Transform, ctx.dst->Diagnostics())
-                << "arrayLength() expected pointer to member access, got "
+                << "expected form of arrayLength argument to be &array_var or "
-                   "pointer to "
+                   "&struct_var.array_member";
-                << array_expr->TypeInfo().name;
             break;
           }
-          auto* storage_buffer_expr = accessor->structure;
+          auto* storage_buffer_var = storage_buffer_sem->Variable();
-          auto* storage_buffer_sem = sem.Get(storage_buffer_expr);
+          auto* storage_buffer_type = storage_buffer_sem->Type()->UnwrapRef();
-          auto* storage_buffer_type =
-              storage_buffer_sem->Type()->UnwrapRef()->As<sem::Struct>();
 
           // Generate BufferSizeIntrinsic for this storage type if we haven't
           // already
           auto buffer_size = get_buffer_size_intrinsic(storage_buffer_type);
 
-          if (!storage_buffer_type) {
-            TINT_ICE(Transform, ctx.dst->Diagnostics())
-                << "arrayLength(X.Y) expected X to be sem::Struct, got "
-                << storage_buffer_type->FriendlyName(ctx.src->Symbols());
-            break;
-          }
-
           // Find the current statement block
           auto* block = call->Stmt()->Block()->Declaration();
 
-          // If the storage_buffer_expr is resolves to a variable (typically
-          // true) then key the array_length from the variable. If not, key off
-          // the expression semantic node, which will be unique per call to
-          // arrayLength().
-          const sem::Node* storage_buffer_usage = storage_buffer_sem;
-          if (auto* user = storage_buffer_sem->As<sem::VariableUser>()) {
-            storage_buffer_usage = user->Variable();
-          }
-
           auto array_length = utils::GetOrCreate(
-              array_length_by_usage, {block, storage_buffer_usage}, [&] {
+              array_length_by_usage, {block, storage_buffer_var}, [&] {
                 // First time this array length is used for this block.
                 // Let's calculate it.
 
-                // Semantic info for the runtime array structure member
-                auto* array_member_sem = storage_buffer_type->Members().back();
-
                 // Construct the variable that'll hold the result of
                 // RWByteAddressBuffer.GetDimensions()
                 auto* buffer_size_result = ctx.dst->Decl(
@@ -208,14 +187,28 @@ void CalculateArrayLength::Run(CloneContext& ctx, const DataMap&, DataMap&) {
                 // array_length = ----------------------------------------
                 //                             array_stride
                 auto name = ctx.dst->Sym();
-                uint32_t array_offset = array_member_sem->Offset();
+                const ast::Expression* total_size =
-                uint32_t array_stride = array_member_sem->Size();
+                    ctx.dst->Expr(buffer_size_result->variable);
-                auto* array_length_var = ctx.dst->Decl(ctx.dst->Const(
+                const sem::Array* array_type = nullptr;
-                    name, ctx.dst->ty.u32(),
+                if (auto* str = storage_buffer_type->As<sem::Struct>()) {
-                    ctx.dst->Div(
+                  // The variable is a struct, so subtract the byte offset of
-                        ctx.dst->Sub(buffer_size_result->variable->symbol,
+                  // the array member.
-                                     array_offset),
+                  auto* array_member_sem = str->Members().back();
-                        array_stride)));
+                  array_type = array_member_sem->Type()->As<sem::Array>();
+                  total_size =
+                      ctx.dst->Sub(total_size, array_member_sem->Offset());
+                } else if (auto* arr = storage_buffer_type->As<sem::Array>()) {
+                  array_type = arr;
+                } else {
+                  TINT_ICE(Transform, ctx.dst->Diagnostics())
+                      << "expected form of arrayLength argument to be "
+                         "&array_var or &struct_var.array_member";
+                  return name;
+                }
+                uint32_t array_stride = array_type->Size();
+                auto* array_length_var = ctx.dst->Decl(
+                    ctx.dst->Const(name, ctx.dst->ty.u32(),
+                                   ctx.dst->Div(total_size, array_stride)));
 
                 // Insert the array length calculations at the top of the block
                 ctx.InsertBefore(block->statements, block->statements[0],

src/transform/calculate_array_length_test.cc

@@ -38,6 +38,36 @@ TEST_F(CalculateArrayLengthTest, Error_MissingCalculateArrayLength) {
 
 TEST_F(CalculateArrayLengthTest, Basic) {
   auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var len : u32 = arrayLength(&sb);
+}
+)";
+
+  auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : array<i32>, result : ptr<function, u32>)
+
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var tint_symbol_1 : u32 = 0u;
+  tint_symbol(sb, &(tint_symbol_1));
+  let tint_symbol_2 : u32 = (tint_symbol_1 / 4u);
+  var len : u32 = tint_symbol_2;
+}
+)";
+
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, BasicInStruct) {
+  auto* src = R"(
 struct SB {
   x : i32;
   arr : array<i32>;
@@ -78,6 +108,40 @@ fn main() {
 
 TEST_F(CalculateArrayLengthTest, InSameBlock) {
   auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var a : u32 = arrayLength(&sb);
+  var b : u32 = arrayLength(&sb);
+  var c : u32 = arrayLength(&sb);
+}
+)";
+
+  auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : array<i32>, result : ptr<function, u32>)
+
+[[group(0), binding(0)]] var<storage, read> sb : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var tint_symbol_1 : u32 = 0u;
+  tint_symbol(sb, &(tint_symbol_1));
+  let tint_symbol_2 : u32 = (tint_symbol_1 / 4u);
+  var a : u32 = tint_symbol_2;
+  var b : u32 = tint_symbol_2;
+  var c : u32 = tint_symbol_2;
+}
+)";
+
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, InSameBlock_Struct) {
+  auto* src = R"(
 struct SB {
   x : i32;
   arr : array<i32>;
@@ -122,6 +186,36 @@ fn main() {
 
 TEST_F(CalculateArrayLengthTest, WithStride) {
   auto* src = R"(
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var len : u32 = arrayLength(&sb);
+}
+)";
+
+  auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : [[stride(64)]] array<i32>, result : ptr<function, u32>)
+
+[[group(0), binding(0)]] var<storage, read> sb : [[stride(64)]] array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var tint_symbol_1 : u32 = 0u;
+  tint_symbol(sb, &(tint_symbol_1));
+  let tint_symbol_2 : u32 = (tint_symbol_1 / 64u);
+  var len : u32 = tint_symbol_2;
+}
+)";
+
+  auto got = Run<Unshadow, SimplifyPointers, CalculateArrayLength>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(CalculateArrayLengthTest, WithStride_InStruct) {
+  auto* src = R"(
 struct SB {
   x : i32;
   y : f32;
@@ -233,15 +327,21 @@ struct SB2 {
 
 [[group(0), binding(1)]] var<storage, read> sb2 : SB2;
 
+[[group(0), binding(2)]] var<storage, read> sb3 : array<i32>;
+
 [[stage(compute), workgroup_size(1)]]
 fn main() {
   var len1 : u32 = arrayLength(&(sb1.arr1));
   var len2 : u32 = arrayLength(&(sb2.arr2));
-  var x : u32 = (len1 + len2);
+  var len3 : u32 = arrayLength(&sb3);
+  var x : u32 = (len1 + len2 + len3);
 }
 )";
 
   auto* expect = R"(
+[[internal(intrinsic_buffer_size)]]
+fn tint_symbol_6([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : array<i32>, result : ptr<function, u32>)
+
 struct SB1 {
   x : i32;
   arr1 : array<i32>;
@@ -262,6 +362,8 @@ fn tint_symbol_3([[internal(disable_validation__ignore_constructible_function_pa
 
 [[group(0), binding(1)]] var<storage, read> sb2 : SB2;
 
+[[group(0), binding(2)]] var<storage, read> sb3 : array<i32>;
+
 [[stage(compute), workgroup_size(1)]]
 fn main() {
   var tint_symbol_1 : u32 = 0u;
@@ -270,9 +372,13 @@ fn main() {
   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 x : u32 = (len1 + len2);
+  var len3 : u32 = tint_symbol_8;
+  var x : u32 = ((len1 + len2) + len3);
 }
 )";
 

test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    let p = &G;
+    let p2 = &(*p);
+    let l1 : u32 = arrayLength(p2);
+}

test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/intrinsics/arrayLength/complex_via_let_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let p = &(G);
+  let p2 = &(*(p));
+  let l1 : u32 = arrayLength(p2);
+}

test/intrinsics/arrayLength/simple_no_struct.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    let l1 : u32 = arrayLength(&G);
+}

test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/intrinsics/arrayLength/simple_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let l1 : u32 = arrayLength(&(G));
+}

test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl

@@ -0,0 +1,9 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let p = &*&G;
+  let p2 = &*p;
+  let p3 = &(*p);
+  let l1 : u32 = arrayLength(&*p3);
+}

test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/intrinsics/arrayLength/via_let_complex_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,9 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let p = &(*(&(G)));
+  let p2 = &(*(p));
+  let p3 = &(*(p));
+  let l1 : u32 = arrayLength(&(*(p3)));
+}

test/intrinsics/arrayLength/via_let_no_struct.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    let p = &G;
+    let p2 = p;
+    let l1 : u32 = arrayLength(p2);
+}

test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.hlsl

@@ -0,0 +1,10 @@
+ByteAddressBuffer G : register(t0, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  uint tint_symbol_1 = 0u;
+  G.GetDimensions(tint_symbol_1);
+  const uint tint_symbol_2 = (tint_symbol_1 / 4u);
+  const uint l1 = tint_symbol_2;
+  return;
+}

test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.msl

@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_symbol_1 {
+  /* 0x0000 */ uint4 buffer_size[1];
+};
+
+kernel void tint_symbol(const constant tint_symbol_1* tint_symbol_3 [[buffer(30)]]) {
+  uint const l1 = ((*(tint_symbol_3)).buffer_size[0u][0u] / 4u);
+  return;
+}
+

test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.spvasm

@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %G_block "G_block"
+               OpMemberName %G_block 0 "inner"
+               OpName %G "G"
+               OpName %main "main"
+               OpDecorate %G_block Block
+               OpMemberDecorate %G_block 0 Offset 0
+               OpDecorate %_runtimearr_int ArrayStride 4
+               OpDecorate %G NonWritable
+               OpDecorate %G DescriptorSet 0
+               OpDecorate %G Binding 0
+        %int = OpTypeInt 32 1
+%_runtimearr_int = OpTypeRuntimeArray %int
+    %G_block = OpTypeStruct %_runtimearr_int
+%_ptr_StorageBuffer_G_block = OpTypePointer StorageBuffer %G_block
+          %G = OpVariable %_ptr_StorageBuffer_G_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpArrayLength %uint %G 0
+               OpReturn
+               OpFunctionEnd

test/intrinsics/arrayLength/via_let_no_struct.wgsl.expected.wgsl

@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> G : array<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let p = &(G);
+  let p2 = p;
+  let l1 : u32 = arrayLength(p2);
+}