diff --git a/include/tint/tint.h b/include/tint/tint.h
index ac1c8b3..9d28687 100644
--- a/include/tint/tint.h
+++ b/include/tint/tint.h
@@ -59,4 +59,7 @@
 #include "src/writer/hlsl/generator.h"
 #endif  // TINT_BUILD_HLSL_WRITER
 
+// TODO(crbug/984): Remove once Dawn builds with this flag
+#define TINT_EXPECTS_UBOS_TO_BE_MULTIPLE_OF_16
+
 #endif  // INCLUDE_TINT_TINT_H_
diff --git a/src/transform/array_length_from_uniform.cc b/src/transform/array_length_from_uniform.cc
index bb182c9..35fbcb9 100644
--- a/src/transform/array_length_from_uniform.cc
+++ b/src/transform/array_length_from_uniform.cc
@@ -67,11 +67,16 @@
   ast::Variable* buffer_size_ubo = nullptr;
   auto get_ubo = [&]() {
     if (!buffer_size_ubo) {
+      // Emit an array<vec4<u32>, N>, where N is 1/4 number of elements.
+      // We do this because UBOs require an element stride that is 16-byte
+      // aligned.
       auto* buffer_size_struct = ctx.dst->Structure(
           ctx.dst->Sym(),
           {ctx.dst->Member(
               kBufferSizeMemberName,
-              ctx.dst->ty.array(ctx.dst->ty.u32(), max_buffer_size_index + 1))},
+              ctx.dst->ty.array(ctx.dst->ty.vec4(ctx.dst->ty.u32()),
+                                (max_buffer_size_index / 4) + 1))},
+
           ast::DecorationList{ctx.dst->create<ast::StructBlockDecoration>()});
       buffer_size_ubo = ctx.dst->Global(
           ctx.dst->Sym(), ctx.dst->ty.Of(buffer_size_struct),
@@ -99,18 +104,20 @@
 
     // Get the storage buffer that contains the runtime array.
     // We assume that the argument to `arrayLength` has the form
-    // `&resource.array`, which requires that `InlinePointerLets` and `Simplify`
-    // have been run before this transform.
+    // `&resource.array`, which requires that `InlinePointerLets` and
+    // `Simplify` have been run before this transform.
     auto* param = call_expr->params()[0]->As<ast::UnaryOpExpression>();
     if (!param || param->op() != ast::UnaryOp::kAddressOf) {
       TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "expected form of arrayLength argument to be &resource.array";
+          << "expected form of arrayLength argument to be "
+             "&resource.array";
       break;
     }
     auto* accessor = param->expr()->As<ast::MemberAccessorExpression>();
     if (!accessor) {
       TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "expected form of arrayLength argument to be &resource.array";
+          << "expected form of arrayLength argument to be "
+             "&resource.array";
       break;
     }
     auto* storage_buffer_expr = accessor->structure();
@@ -118,7 +125,8 @@
         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 &resource.array";
+          << "expected form of arrayLength argument to be "
+             "&resource.array";
       break;
     }
 
@@ -135,9 +143,13 @@
     }
 
     // Load the total storage buffer size from the UBO.
-    auto* total_storage_buffer_size = ctx.dst->IndexAccessor(
+    uint32_t array_index = idx_itr->second / 4;
+    auto* vec_expr = ctx.dst->IndexAccessor(
         ctx.dst->MemberAccessor(get_ubo()->symbol(), kBufferSizeMemberName),
-        idx_itr->second);
+        array_index);
+    uint32_t vec_index = idx_itr->second % 4;
+    auto* total_storage_buffer_size =
+        ctx.dst->IndexAccessor(vec_expr, vec_index);
 
     // Calculate actual array length
     //                total_storage_buffer_size - array_offset
diff --git a/src/transform/array_length_from_uniform_test.cc b/src/transform/array_length_from_uniform_test.cc
index a173de8..6ab39ee 100644
--- a/src/transform/array_length_from_uniform_test.cc
+++ b/src/transform/array_length_from_uniform_test.cc
@@ -81,7 +81,7 @@
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<u32, 1>;
+  buffer_size : array<vec4<u32>, 1>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -96,7 +96,7 @@
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var len : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u);
+  var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 4u) / 4u);
 }
 )";
 
@@ -134,7 +134,7 @@
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<u32, 1>;
+  buffer_size : array<vec4<u32>, 1>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -150,7 +150,7 @@
 
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var len : u32 = ((tint_symbol_1.buffer_size[0u] - 8u) / 64u);
+  var len : u32 = ((tint_symbol_1.buffer_size[0u][0u] - 8u) / 64u);
 }
 )";
 
@@ -175,29 +175,48 @@
   x : i32;
   arr1 : array<i32>;
 };
-
 [[block]]
 struct SB2 {
   x : i32;
   arr2 : array<vec4<f32>>;
 };
+[[block]]
+struct SB3 {
+  x : i32;
+  arr3 : array<vec4<f32>>;
+};
+[[block]]
+struct SB4 {
+  x : i32;
+  arr4 : array<vec4<f32>>;
+};
+[[block]]
+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(3), binding(2)]] var<storage, read> sb4 : SB4;
+[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
 
 [[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.arr3));
+  var len4 : u32 = arrayLength(&(sb4.arr4));
+  var len5 : u32 = arrayLength(&(sb5.arr5));
+  var x : u32 = (len1 + len2 + len3 + len4 + len5);
 }
 )";
 
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<u32, 2>;
+  buffer_size : array<vec4<u32>, 2>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -214,21 +233,51 @@
   arr2 : array<vec4<f32>>;
 };
 
+[[block]]
+struct SB3 {
+  x : i32;
+  arr3 : array<vec4<f32>>;
+};
+
+[[block]]
+struct SB4 {
+  x : i32;
+  arr4 : array<vec4<f32>>;
+};
+
+[[block]]
+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(3), binding(2)]] var<storage, read> sb4 : SB4;
+
+[[group(4), binding(2)]] var<storage, read> sb5 : SB5;
+
 [[stage(compute), workgroup_size(1)]]
 fn main() {
-  var len1 : u32 = ((tint_symbol_1.buffer_size[0u] - 4u) / 4u);
-  var len2 : u32 = ((tint_symbol_1.buffer_size[1u] - 16u) / 16u);
-  var x : u32 = (len1 + len2);
+  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 len4 : u32 = ((tint_symbol_1.buffer_size[0u][3u] - 16u) / 16u);
+  var len5 : u32 = ((tint_symbol_1.buffer_size[1u][0u] - 16u) / 16u);
+  var x : u32 = ((((len1 + len2) + len3) + len4) + len5);
 }
 )";
 
   ArrayLengthFromUniform::Config cfg({0, 30u});
   cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{0, 2u}, 0);
   cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{1u, 2u}, 1);
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{2u, 2u}, 2);
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{3u, 2u}, 3);
+  cfg.bindpoint_to_size_index.emplace(sem::BindingPoint{4u, 2u}, 4);
 
   DataMap data;
   data.Add<ArrayLengthFromUniform::Config>(std::move(cfg));
diff --git a/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl b/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl
index d690d24..9c7e44c 100644
--- a/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl
+++ b/test/intrinsics/arrayLength/complex_via_let.wgsl.expected.msl
@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 
diff --git a/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl b/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl
index c86fc59..40f5a45 100644
--- a/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl
+++ b/test/intrinsics/arrayLength/deprecated.wgsl.expected.msl
@@ -2,15 +2,15 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
-  uint const l2 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
+  uint const l2 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 
diff --git a/test/intrinsics/arrayLength/simple.wgsl.expected.msl b/test/intrinsics/arrayLength/simple.wgsl.expected.msl
index d690d24..9c7e44c 100644
--- a/test/intrinsics/arrayLength/simple.wgsl.expected.msl
+++ b/test/intrinsics/arrayLength/simple.wgsl.expected.msl
@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 
diff --git a/test/intrinsics/arrayLength/via_let.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let.wgsl.expected.msl
index d690d24..9c7e44c 100644
--- a/test/intrinsics/arrayLength/via_let.wgsl.expected.msl
+++ b/test/intrinsics/arrayLength/via_let.wgsl.expected.msl
@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 
diff --git a/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl b/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl
index d690d24..9c7e44c 100644
--- a/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl
+++ b/test/intrinsics/arrayLength/via_let_complex.wgsl.expected.msl
@@ -2,14 +2,14 @@
 
 using namespace metal;
 struct tint_symbol_1 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct S {
   /* 0x0000 */ int a[1];
 };
 
 kernel void tint_symbol(constant tint_symbol_1& tint_symbol_2 [[buffer(30)]]) {
-  uint const l1 = ((tint_symbol_2.buffer_size[0u] - 0u) / 4u);
+  uint const l1 = ((tint_symbol_2.buffer_size[0u][0u] - 0u) / 4u);
   return;
 }
 
diff --git a/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl
index 4f53a60..3553097 100644
--- a/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl
+++ b/test/intrinsics/gen/arrayLength/1588cd.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[2];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RO {
   /* 0x0000 */ int arg_0[1];
@@ -12,7 +12,7 @@
 };
 
 void arrayLength_1588cd(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
diff --git a/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl
index 71c7a83..97d2857 100644
--- a/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl
+++ b/test/intrinsics/gen/arrayLength/61b1c7.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RW {
   /* 0x0000 */ int arg_0[1];
@@ -12,7 +12,7 @@
 };
 
 void arrayLength_61b1c7(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
diff --git a/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl
index ad28d92..ed94999 100644
--- a/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl
+++ b/test/intrinsics/gen/arrayLength/a0f5ca.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[2];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RO {
   /* 0x0000 */ float arg_0[1];
@@ -12,7 +12,7 @@
 };
 
 void arrayLength_a0f5ca(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
diff --git a/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl
index fdbb89d..525920c 100644
--- a/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl
+++ b/test/intrinsics/gen/arrayLength/cdd123.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RW {
   /* 0x0000 */ float arg_0[1];
@@ -12,7 +12,7 @@
 };
 
 void arrayLength_cdd123(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
diff --git a/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl
index 14068e5..97cbb4f 100644
--- a/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl
+++ b/test/intrinsics/gen/arrayLength/cfca0a.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[2];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RO {
   /* 0x0000 */ uint arg_0[1];
@@ -12,7 +12,7 @@
 };
 
 void arrayLength_cfca0a(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[1u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][1u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {
diff --git a/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl b/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl
index 04f8508..d345929 100644
--- a/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl
+++ b/test/intrinsics/gen/arrayLength/eb510f.wgsl.expected.msl
@@ -2,7 +2,7 @@
 
 using namespace metal;
 struct tint_symbol_2 {
-  /* 0x0000 */ uint buffer_size[1];
+  /* 0x0000 */ uint4 buffer_size[1];
 };
 struct SB_RW {
   /* 0x0000 */ uint arg_0[1];
@@ -12,7 +12,7 @@
 };
 
 void arrayLength_eb510f(constant tint_symbol_2& tint_symbol_3) {
-  uint res = ((tint_symbol_3.buffer_size[0u] - 0u) / 4u);
+  uint res = ((tint_symbol_3.buffer_size[0u][0u] - 0u) / 4u);
 }
 
 vertex tint_symbol vertex_main(constant tint_symbol_2& tint_symbol_3 [[buffer(30)]]) {