diff --git a/docs/origin-trial-changes.md b/docs/origin-trial-changes.md
new file mode 100644
index 0000000000..2742515025
--- /dev/null
+++ b/docs/origin-trial-changes.md
@@ -0,0 +1,6 @@
+# Tint changes during Origin Trial
+
+## Changes for M95
+
+### New Features
+* The size of an array can now be defined using a non-overridable module-scope constant
diff --git a/src/ast/array.cc b/src/ast/array.cc
index 2da027fef8..955089707a 100644
--- a/src/ast/array.cc
+++ b/src/ast/array.cc
@@ -23,10 +23,32 @@ TINT_INSTANTIATE_TYPEINFO(tint::ast::Array);
 namespace tint {
 namespace ast {
 
+namespace {
+// Returns the string representation of an array size expression.
+std::string SizeExprToString(const ast::Expression* size,
+                             const SymbolTable* symbols = nullptr) {
+  if (auto* ident = size->As<ast::IdentifierExpression>()) {
+    if (symbols) {
+      return symbols->NameFor(ident->symbol());
+    } else {
+      return ident->symbol().to_str();
+    }
+  } else if (auto* scalar = size->As<ast::ScalarConstructorExpression>()) {
+    auto* literal = scalar->literal()->As<ast::IntLiteral>();
+    if (literal) {
+      return std::to_string(literal->value_as_u32());
+    }
+  }
+  // This will never be exposed to the user as the Resolver will reject this
+  // expression for array size.
+  return "<invalid>";
+}
+}  // namespace
+
 Array::Array(ProgramID program_id,
              const Source& source,
              Type* subtype,
-             uint32_t size,
+             ast::Expression* size,
              ast::DecorationList decorations)
     : Base(program_id, source),
       subtype_(subtype),
@@ -42,7 +64,7 @@ std::string Array::type_name() const {
 
   std::string type_name = "__array" + subtype_->type_name();
   if (!IsRuntimeArray()) {
-    type_name += "_" + std::to_string(size_);
+    type_name += "_" + SizeExprToString(size_);
   }
   for (auto* deco : decos_) {
     if (auto* stride = deco->As<ast::StrideDecoration>()) {
@@ -62,7 +84,7 @@ std::string Array::FriendlyName(const SymbolTable& symbols) const {
   }
   out << "array<" << subtype_->FriendlyName(symbols);
   if (!IsRuntimeArray()) {
-    out << ", " << size_;
+    out << ", " << SizeExprToString(size_, &symbols);
   }
   out << ">";
   return out.str();
@@ -72,8 +94,9 @@ Array* Array::Clone(CloneContext* ctx) const {
   // Clone arguments outside of create() call to have deterministic ordering
   auto src = ctx->Clone(source());
   auto* ty = ctx->Clone(type());
+  auto* size = ctx->Clone(Size());
   auto decos = ctx->Clone(decorations());
-  return ctx->dst->create<Array>(src, ty, size_, decos);
+  return ctx->dst->create<Array>(src, ty, size, decos);
 }
 
 }  // namespace ast
diff --git a/src/ast/array.h b/src/ast/array.h
index c5ba1b8f09..7e4b004dad 100644
--- a/src/ast/array.h
+++ b/src/ast/array.h
@@ -23,6 +23,9 @@
 namespace tint {
 namespace ast {
 
+// Forward declarations.
+class Expression;
+
 /// An array type. If size is zero then it is a runtime array.
 class Array : public Castable<Array, Type> {
  public:
@@ -30,13 +33,13 @@ class Array : public Castable<Array, Type> {
   /// @param program_id the identifier of the program that owns this node
   /// @param source the source of this node
   /// @param subtype the type of the array elements
-  /// @param size the number of elements in the array. `0` represents a
+  /// @param size the number of elements in the array. nullptr represents a
   /// runtime-sized array.
   /// @param decorations the array decorations
   Array(ProgramID program_id,
         const Source& source,
         Type* subtype,
-        uint32_t size,
+        ast::Expression* size,
         ast::DecorationList decorations);
   /// Move constructor
   Array(Array&&);
@@ -44,15 +47,16 @@ class Array : public Castable<Array, Type> {
 
   /// @returns true if this is a runtime array.
   /// i.e. the size is determined at runtime
-  bool IsRuntimeArray() const { return size_ == 0; }
+  bool IsRuntimeArray() const { return size_ == nullptr; }
 
   /// @returns the array decorations
   const ast::DecorationList& decorations() const { return decos_; }
 
   /// @returns the array type
   Type* type() const { return subtype_; }
-  /// @returns the array size. Size is 0 for a runtime array
-  uint32_t size() const { return size_; }
+
+  /// @returns the array size, or nullptr for a runtime array
+  ast::Expression* Size() const { return size_; }
 
   /// @returns the name for the type
   std::string type_name() const override;
@@ -69,7 +73,7 @@ class Array : public Castable<Array, Type> {
 
  private:
   Type* const subtype_;
-  uint32_t const size_;
+  ast::Expression* const size_;
   ast::DecorationList const decos_;
 };
 
diff --git a/src/ast/array_test.cc b/src/ast/array_test.cc
index 5e8b9c81d6..f6413d9d70 100644
--- a/src/ast/array_test.cc
+++ b/src/ast/array_test.cc
@@ -24,57 +24,76 @@ using AstArrayTest = TestHelper;
 
 TEST_F(AstArrayTest, CreateSizedArray) {
   auto* u32 = create<U32>();
-  auto* arr = create<Array>(u32, 3, DecorationList{});
+  auto* size = Expr(3);
+  auto* arr = create<Array>(u32, size, DecorationList{});
   EXPECT_EQ(arr->type(), u32);
-  EXPECT_EQ(arr->size(), 3u);
+  EXPECT_EQ(arr->Size(), size);
   EXPECT_TRUE(arr->Is<Array>());
   EXPECT_FALSE(arr->IsRuntimeArray());
 }
 
 TEST_F(AstArrayTest, CreateRuntimeArray) {
   auto* u32 = create<U32>();
-  auto* arr = create<Array>(u32, 0, DecorationList{});
+  auto* arr = create<Array>(u32, nullptr, DecorationList{});
   EXPECT_EQ(arr->type(), u32);
-  EXPECT_EQ(arr->size(), 0u);
+  EXPECT_EQ(arr->Size(), nullptr);
   EXPECT_TRUE(arr->Is<Array>());
   EXPECT_TRUE(arr->IsRuntimeArray());
 }
 
 TEST_F(AstArrayTest, TypeName) {
   auto* i32 = create<I32>();
-  auto* arr = create<Array>(i32, 0, DecorationList{});
+  auto* arr = create<Array>(i32, nullptr, DecorationList{});
   EXPECT_EQ(arr->type_name(), "__array__i32");
 }
 
-TEST_F(AstArrayTest, FriendlyNameRuntimeSized) {
+TEST_F(AstArrayTest, FriendlyName_RuntimeSized) {
   auto* i32 = create<I32>();
-  auto* arr = create<Array>(i32, 0, DecorationList{});
+  auto* arr = create<Array>(i32, nullptr, DecorationList{});
   EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32>");
 }
 
-TEST_F(AstArrayTest, FriendlyNameStaticSized) {
+TEST_F(AstArrayTest, FriendlyName_LiteralSized) {
   auto* i32 = create<I32>();
-  auto* arr = create<Array>(i32, 5, DecorationList{});
+  auto* arr = create<Array>(i32, Expr(5), DecorationList{});
   EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32, 5>");
 }
 
-TEST_F(AstArrayTest, FriendlyNameWithStride) {
+TEST_F(AstArrayTest, FriendlyName_ConstantSized) {
+  auto* i32 = create<I32>();
+  auto* arr = create<Array>(i32, Expr("size"), DecorationList{});
+  EXPECT_EQ(arr->FriendlyName(Symbols()), "array<i32, size>");
+}
+
+TEST_F(AstArrayTest, FriendlyName_WithStride) {
   auto* i32 = create<I32>();
   auto* arr =
-      create<Array>(i32, 5, DecorationList{create<StrideDecoration>(32)});
+      create<Array>(i32, Expr(5), DecorationList{create<StrideDecoration>(32)});
   EXPECT_EQ(arr->FriendlyName(Symbols()), "[[stride(32)]] array<i32, 5>");
 }
 
-TEST_F(AstArrayTest, TypeName_RuntimeArray) {
+TEST_F(AstArrayTest, TypeName_RuntimeSized) {
   auto* i32 = create<I32>();
-  auto* arr = create<Array>(i32, 3, DecorationList{});
+  auto* arr = create<Array>(i32, nullptr, DecorationList{});
+  EXPECT_EQ(arr->type_name(), "__array__i32");
+}
+
+TEST_F(AstArrayTest, TypeName_LiteralSized) {
+  auto* i32 = create<I32>();
+  auto* arr = create<Array>(i32, Expr(3), DecorationList{});
   EXPECT_EQ(arr->type_name(), "__array__i32_3");
 }
 
+TEST_F(AstArrayTest, TypeName_ConstantSized) {
+  auto* i32 = create<I32>();
+  auto* arr = create<Array>(i32, Expr("size"), DecorationList{});
+  EXPECT_EQ(arr->type_name(), "__array__i32_$1");
+}
+
 TEST_F(AstArrayTest, TypeName_WithStride) {
   auto* i32 = create<I32>();
   auto* arr =
-      create<Array>(i32, 3, DecorationList{create<StrideDecoration>(16)});
+      create<Array>(i32, Expr(3), DecorationList{create<StrideDecoration>(16)});
   EXPECT_EQ(arr->type_name(), "__array__i32_3_stride_16");
 }
 
diff --git a/src/program_builder.h b/src/program_builder.h
index 41865e50fa..fc3b59e994 100644
--- a/src/program_builder.h
+++ b/src/program_builder.h
@@ -622,66 +622,85 @@ class ProgramBuilder {
     }
 
     /// @param subtype the array element type
-    /// @param n the array size. 0 represents a runtime-array
+    /// @param n the array size. nullptr represents a runtime-array
     /// @param decos the optional decorations for the array
     /// @return the tint AST type for a array of size `n` of type `T`
+    template <typename EXPR = ast::Expression*>
     ast::Array* array(ast::Type* subtype,
-                      uint32_t n = 0,
+                      EXPR&& n = nullptr,
                       ast::DecorationList decos = {}) const {
-      return builder->create<ast::Array>(subtype, n, decos);
+      return builder->create<ast::Array>(
+          subtype, builder->Expr(std::forward<EXPR>(n)), decos);
     }
 
     /// @param source the Source of the node
     /// @param subtype the array element type
-    /// @param n the array size. 0 represents a runtime-array
+    /// @param n the array size. nullptr represents a runtime-array
     /// @param decos the optional decorations for the array
     /// @return the tint AST type for a array of size `n` of type `T`
+    template <typename EXPR = ast::Expression*>
     ast::Array* array(const Source& source,
                       ast::Type* subtype,
-                      uint32_t n = 0,
+                      EXPR&& n = nullptr,
                       ast::DecorationList decos = {}) const {
-      return builder->create<ast::Array>(source, subtype, n, decos);
+      return builder->create<ast::Array>(
+          source, subtype, builder->Expr(std::forward<EXPR>(n)), decos);
     }
 
     /// @param subtype the array element type
-    /// @param n the array size. 0 represents a runtime-array
+    /// @param n the array size. nullptr represents a runtime-array
     /// @param stride the array stride. 0 represents implicit stride
     /// @return the tint AST type for a array of size `n` of type `T`
-    ast::Array* array(ast::Type* subtype, uint32_t n, uint32_t stride) const {
+    template <typename EXPR>
+    ast::Array* array(ast::Type* subtype, EXPR&& n, uint32_t stride) const {
       ast::DecorationList decos;
       if (stride) {
         decos.emplace_back(builder->create<ast::StrideDecoration>(stride));
       }
-      return array(subtype, n, std::move(decos));
+      return array(subtype, std::forward<EXPR>(n), std::move(decos));
     }
 
     /// @param source the Source of the node
     /// @param subtype the array element type
-    /// @param n the array size. 0 represents a runtime-array
+    /// @param n the array size. nullptr represents a runtime-array
     /// @param stride the array stride. 0 represents implicit stride
     /// @return the tint AST type for a array of size `n` of type `T`
+    template <typename EXPR>
     ast::Array* array(const Source& source,
                       ast::Type* subtype,
-                      uint32_t n,
+                      EXPR&& n,
                       uint32_t stride) const {
       ast::DecorationList decos;
       if (stride) {
         decos.emplace_back(builder->create<ast::StrideDecoration>(stride));
       }
-      return array(source, subtype, n, std::move(decos));
+      return array(source, subtype, std::forward<EXPR>(n), std::move(decos));
+    }
+
+    /// @return the tint AST type for a runtime-sized array of type `T`
+    template <typename T>
+    ast::Array* array() const {
+      return array(Of<T>(), nullptr);
     }
 
     /// @return the tint AST type for an array of size `N` of type `T`
-    template <typename T, int N = 0>
+    template <typename T, int N>
     ast::Array* array() const {
-      return array(Of<T>(), N);
+      return array(Of<T>(), builder->Expr(N));
+    }
+
+    /// @param stride the array stride
+    /// @return the tint AST type for a runtime-sized array of type `T`
+    template <typename T>
+    ast::Array* array(uint32_t stride) const {
+      return array(Of<T>(), nullptr, stride);
     }
 
     /// @param stride the array stride
     /// @return the tint AST type for an array of size `N` of type `T`
-    template <typename T, int N = 0>
+    template <typename T, int N>
     ast::Array* array(uint32_t stride) const {
-      return array(Of<T>(), N, stride);
+      return array(Of<T>(), builder->Expr(N), stride);
     }
 
     /// Creates a type name
@@ -1273,22 +1292,23 @@ class ProgramBuilder {
 
   /// @param args the arguments for the array constructor
   /// @return an `ast::TypeConstructorExpression` of an array with element type
-  /// `T`, constructed with the values `args`.
-  template <typename T, int N = 0, typename... ARGS>
+  /// `T` and size `N`, constructed with the values `args`.
+  template <typename T, int N, typename... ARGS>
   ast::TypeConstructorExpression* array(ARGS&&... args) {
     return Construct(ty.array<T, N>(), std::forward<ARGS>(args)...);
   }
 
   /// @param subtype the array element type
-  /// @param n the array size. 0 represents a runtime-array.
+  /// @param n the array size. nullptr represents a runtime-array.
   /// @param args the arguments for the array constructor
   /// @return an `ast::TypeConstructorExpression` of an array with element type
   /// `subtype`, constructed with the values `args`.
-  template <typename... ARGS>
+  template <typename EXPR, typename... ARGS>
   ast::TypeConstructorExpression* array(ast::Type* subtype,
-                                        uint32_t n,
+                                        EXPR&& n,
                                         ARGS&&... args) {
-    return Construct(ty.array(subtype, n), std::forward<ARGS>(args)...);
+    return Construct(ty.array(subtype, std::forward<EXPR>(n)),
+                     std::forward<ARGS>(args)...);
   }
 
   /// @param name the variable name
diff --git a/src/reader/spirv/parser_type.cc b/src/reader/spirv/parser_type.cc
index 7802bada83..7556659292 100644
--- a/src/reader/spirv/parser_type.cc
+++ b/src/reader/spirv/parser_type.cc
@@ -187,7 +187,11 @@ Array::Array(const Type* t, uint32_t sz, uint32_t st)
 Array::Array(const Array&) = default;
 
 ast::Type* Array::Build(ProgramBuilder& b) const {
-  return b.ty.array(type->Build(b), size, stride);
+  if (size > 0) {
+    return b.ty.array(type->Build(b), size, stride);
+  } else {
+    return b.ty.array(type->Build(b), nullptr, stride);
+  }
 }
 
 Sampler::Sampler(ast::SamplerKind k) : kind(k) {}
diff --git a/src/reader/wgsl/parser_impl.cc b/src/reader/wgsl/parser_impl.cc
index 637375d8f1..c86277a1ec 100644
--- a/src/reader/wgsl/parser_impl.cc
+++ b/src/reader/wgsl/parser_impl.cc
@@ -1187,7 +1187,7 @@ Expect<ast::Type*> ParserImpl::expect_type_decl_array(
     ast::DecorationList decos) {
   const char* use = "array declaration";
 
-  uint32_t size = 0;
+  ast::Expression* size = nullptr;
 
   auto subtype = expect_lt_gt_block(use, [&]() -> Expect<ast::Type*> {
     auto type = expect_type(use);
@@ -1195,10 +1195,14 @@ Expect<ast::Type*> ParserImpl::expect_type_decl_array(
       return Failure::kErrored;
 
     if (match(Token::Type::kComma)) {
-      auto val = expect_nonzero_positive_sint("array size");
-      if (val.errored)
+      auto expr = primary_expression();
+      if (expr.errored) {
         return Failure::kErrored;
-      size = val.value;
+      } else if (!expr.matched) {
+        return add_error(peek(), "expected array size expression");
+      }
+
+      size = std::move(expr.value);
     }
 
     return type.value;
diff --git a/src/reader/wgsl/parser_impl_error_msg_test.cc b/src/reader/wgsl/parser_impl_error_msg_test.cc
index 07007b4857..0211d214d2 100644
--- a/src/reader/wgsl/parser_impl_error_msg_test.cc
+++ b/src/reader/wgsl/parser_impl_error_msg_test.cc
@@ -862,19 +862,18 @@ TEST_F(ParserImplErrorTest, GlobalDeclVarArrayMissingType) {
          "              ^\n");
 }
 
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayInvalidSize) {
-  EXPECT(
-      "var i : array<u32, x>;",
-      "test.wgsl:1:20 error: expected signed integer literal for array size\n"
-      "var i : array<u32, x>;\n"
-      "                   ^\n");
+TEST_F(ParserImplErrorTest, GlobalDeclVarArrayMissingSize) {
+  EXPECT("var i : array<u32, >;",
+         "test.wgsl:1:20 error: expected array size expression\n"
+         "var i : array<u32, >;\n"
+         "                   ^\n");
 }
 
-TEST_F(ParserImplErrorTest, GlobalDeclVarArrayNegativeSize) {
-  EXPECT("var i : array<u32, -3>;",
-         "test.wgsl:1:20 error: array size must be greater than 0\n"
-         "var i : array<u32, -3>;\n"
-         "                   ^^\n");
+TEST_F(ParserImplErrorTest, GlobalDeclVarArrayInvalidSize) {
+  EXPECT("var i : array<u32, !>;",
+         "test.wgsl:1:20 error: expected array size expression\n"
+         "var i : array<u32, !>;\n"
+         "                   ^\n");
 }
 
 TEST_F(ParserImplErrorTest, GlobalDeclVarDecoListEmpty) {
diff --git a/src/reader/wgsl/parser_impl_type_decl_test.cc b/src/reader/wgsl/parser_impl_type_decl_test.cc
index 68812c3d5e..6b9f5dcacd 100644
--- a/src/reader/wgsl/parser_impl_type_decl_test.cc
+++ b/src/reader/wgsl/parser_impl_type_decl_test.cc
@@ -440,7 +440,7 @@ TEST_F(ParserImplTest, TypeDecl_Atomic_MissingType) {
   ASSERT_EQ(p->error(), "1:8: invalid type for atomic declaration");
 }
 
-TEST_F(ParserImplTest, TypeDecl_Array) {
+TEST_F(ParserImplTest, TypeDecl_Array_SintLiteralSize) {
   auto p = parser("array<f32, 5>");
   auto t = p->type_decl();
   EXPECT_TRUE(t.matched);
@@ -451,10 +451,57 @@ TEST_F(ParserImplTest, TypeDecl_Array) {
 
   auto* a = t.value->As<ast::Array>();
   ASSERT_FALSE(a->IsRuntimeArray());
-  ASSERT_EQ(a->size(), 5u);
   ASSERT_TRUE(a->type()->Is<ast::F32>());
   EXPECT_EQ(a->decorations().size(), 0u);
   EXPECT_EQ(t.value->source().range, (Source::Range{{1u, 1u}, {1u, 14u}}));
+
+  auto* size_expr = a->Size()->As<ast::ScalarConstructorExpression>();
+  ASSERT_NE(size_expr, nullptr);
+  auto* size = size_expr->literal()->As<ast::SintLiteral>();
+  ASSERT_NE(size, nullptr);
+  EXPECT_EQ(size->value_as_i32(), 5);
+}
+
+TEST_F(ParserImplTest, TypeDecl_Array_UintLiteralSize) {
+  auto p = parser("array<f32, 5u>");
+  auto t = p->type_decl();
+  EXPECT_TRUE(t.matched);
+  EXPECT_FALSE(t.errored);
+  ASSERT_NE(t.value, nullptr) << p->error();
+  ASSERT_FALSE(p->has_error());
+  ASSERT_TRUE(t.value->Is<ast::Array>());
+
+  auto* a = t.value->As<ast::Array>();
+  ASSERT_FALSE(a->IsRuntimeArray());
+  ASSERT_TRUE(a->type()->Is<ast::F32>());
+  EXPECT_EQ(a->decorations().size(), 0u);
+  EXPECT_EQ(t.value->source().range, (Source::Range{{1u, 1u}, {1u, 15u}}));
+
+  auto* size_expr = a->Size()->As<ast::ScalarConstructorExpression>();
+  ASSERT_NE(size_expr, nullptr);
+  auto* size = size_expr->literal()->As<ast::UintLiteral>();
+  ASSERT_NE(size, nullptr);
+  EXPECT_EQ(size->value_as_u32(), 5u);
+}
+
+TEST_F(ParserImplTest, TypeDecl_Array_ConstantSize) {
+  auto p = parser("array<f32, size>");
+  auto t = p->type_decl();
+  EXPECT_TRUE(t.matched);
+  EXPECT_FALSE(t.errored);
+  ASSERT_NE(t.value, nullptr) << p->error();
+  ASSERT_FALSE(p->has_error());
+  ASSERT_TRUE(t.value->Is<ast::Array>());
+
+  auto* a = t.value->As<ast::Array>();
+  ASSERT_FALSE(a->IsRuntimeArray());
+  ASSERT_TRUE(a->type()->Is<ast::F32>());
+  EXPECT_EQ(a->decorations().size(), 0u);
+  EXPECT_EQ(t.value->source().range, (Source::Range{{1u, 1u}, {1u, 17u}}));
+
+  auto* size_expr = a->Size()->As<ast::IdentifierExpression>();
+  ASSERT_NE(size_expr, nullptr);
+  EXPECT_EQ(p->builder().Symbols().NameFor(size_expr->symbol()), "size");
 }
 
 TEST_F(ParserImplTest, TypeDecl_Array_Stride) {
@@ -468,9 +515,14 @@ TEST_F(ParserImplTest, TypeDecl_Array_Stride) {
 
   auto* a = t.value->As<ast::Array>();
   ASSERT_FALSE(a->IsRuntimeArray());
-  ASSERT_EQ(a->size(), 5u);
   ASSERT_TRUE(a->type()->Is<ast::F32>());
 
+  auto* size_expr = a->Size()->As<ast::ScalarConstructorExpression>();
+  ASSERT_NE(size_expr, nullptr);
+  auto* size = size_expr->literal()->As<ast::SintLiteral>();
+  ASSERT_NE(size, nullptr);
+  EXPECT_EQ(size->value_as_i32(), 5);
+
   ASSERT_EQ(a->decorations().size(), 1u);
   auto* stride = a->decorations()[0];
   ASSERT_TRUE(stride->Is<ast::StrideDecoration>());
@@ -664,34 +716,24 @@ TEST_F(ParserImplTest, TypeDecl_Array_BadType) {
   ASSERT_EQ(p->error(), "1:7: unknown type 'unknown'");
 }
 
-TEST_F(ParserImplTest, TypeDecl_Array_ZeroSize) {
-  auto p = parser("array<f32, 0>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  ASSERT_EQ(p->error(), "1:12: array size must be greater than 0");
-}
-
-TEST_F(ParserImplTest, TypeDecl_Array_NegativeSize) {
-  auto p = parser("array<f32, -1>");
-  auto t = p->type_decl();
-  EXPECT_TRUE(t.errored);
-  EXPECT_FALSE(t.matched);
-  ASSERT_EQ(t.value, nullptr);
-  ASSERT_TRUE(p->has_error());
-  ASSERT_EQ(p->error(), "1:12: array size must be greater than 0");
-}
-
 TEST_F(ParserImplTest, TypeDecl_Array_BadSize) {
-  auto p = parser("array<f32, invalid>");
+  auto p = parser("array<f32, !>");
   auto t = p->type_decl();
   EXPECT_TRUE(t.errored);
   EXPECT_FALSE(t.matched);
   ASSERT_EQ(t.value, nullptr);
   ASSERT_TRUE(p->has_error());
-  ASSERT_EQ(p->error(), "1:12: expected signed integer literal for array size");
+  ASSERT_EQ(p->error(), "1:12: expected array size expression");
+}
+
+TEST_F(ParserImplTest, TypeDecl_Array_MissingSize) {
+  auto p = parser("array<f32,>");
+  auto t = p->type_decl();
+  EXPECT_TRUE(t.errored);
+  EXPECT_FALSE(t.matched);
+  ASSERT_EQ(t.value, nullptr);
+  ASSERT_TRUE(p->has_error());
+  ASSERT_EQ(p->error(), "1:11: expected array size expression");
 }
 
 TEST_F(ParserImplTest, TypeDecl_Array_MissingLessThan) {
diff --git a/src/resolver/assignment_validation_test.cc b/src/resolver/assignment_validation_test.cc
index 580e5521d0..0063a2ad50 100644
--- a/src/resolver/assignment_validation_test.cc
+++ b/src/resolver/assignment_validation_test.cc
@@ -62,6 +62,49 @@ TEST_F(ResolverAssignmentValidationTest, AssignIncompatibleTypes) {
   EXPECT_EQ(r()->error(), "12:34 error: cannot assign 'f32' to 'i32'");
 }
 
+TEST_F(ResolverAssignmentValidationTest,
+       AssignArraysWithDifferentSizeExpressions_Pass) {
+  // let len = 4u;
+  // {
+  //   var a : array<f32, 4>;
+  //   var b : array<f32, len>;
+  //   a = b;
+  // }
+
+  GlobalConst("len", nullptr, Expr(4u));
+
+  auto* a = Var("a", ty.array(ty.f32(), 4));
+  auto* b = Var("b", ty.array(ty.f32(), "len"));
+
+  auto* assign = Assign(Source{{12, 34}}, "a", "b");
+  WrapInFunction(a, b, assign);
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverAssignmentValidationTest,
+       AssignArraysWithDifferentSizeExpressions_Fail) {
+  // let len = 5u;
+  // {
+  //   var a : array<f32, 4>;
+  //   var b : array<f32, len>;
+  //   a = b;
+  // }
+
+  GlobalConst("len", nullptr, Expr(5u));
+
+  auto* a = Var("a", ty.array(ty.f32(), 4));
+  auto* b = Var("b", ty.array(ty.f32(), "len"));
+
+  auto* assign = Assign(Source{{12, 34}}, "a", "b");
+  WrapInFunction(a, b, assign);
+
+  ASSERT_FALSE(r()->Resolve());
+
+  EXPECT_EQ(r()->error(),
+            "12:34 error: cannot assign 'array<f32, len>' to 'array<f32, 4>'");
+}
+
 TEST_F(ResolverAssignmentValidationTest,
        AssignCompatibleTypesInBlockStatement_Pass) {
   // {
diff --git a/src/resolver/decoration_validation_test.cc b/src/resolver/decoration_validation_test.cc
index 2808ea063a..a1ceeea1e3 100644
--- a/src/resolver/decoration_validation_test.cc
+++ b/src/resolver/decoration_validation_test.cc
@@ -681,7 +681,7 @@ using ArrayDecorationTest = TestWithParams;
 TEST_P(ArrayDecorationTest, IsValid) {
   auto& params = GetParam();
 
-  auto* arr = ty.array(ty.f32(), 0,
+  auto* arr = ty.array(ty.f32(), nullptr,
                        createDecorations(Source{{12, 34}}, *this, params.kind));
   Structure("mystruct",
             {
diff --git a/src/resolver/function_validation_test.cc b/src/resolver/function_validation_test.cc
index 386c9dfdc8..4e87a8cc05 100644
--- a/src/resolver/function_validation_test.cc
+++ b/src/resolver/function_validation_test.cc
@@ -698,7 +698,7 @@ TEST_F(ResolverFunctionValidationTest, ReturnIsConstructible_StructOfAtomic) {
 }
 
 TEST_F(ResolverFunctionValidationTest, ReturnIsConstructible_RuntimeArray) {
-  auto* ret_type = ty.array(Source{{12, 34}}, ty.i32(), 0);
+  auto* ret_type = ty.array(Source{{12, 34}}, ty.i32());
   Func("f", {}, ret_type, {});
 
   EXPECT_FALSE(r()->Resolve());
diff --git a/src/resolver/resolver.cc b/src/resolver/resolver.cc
index 9c3058d6ac..20acfa86fe 100644
--- a/src/resolver/resolver.cc
+++ b/src/resolver/resolver.cc
@@ -3868,11 +3868,68 @@ sem::Array* Resolver::Array(const ast::Array* arr) {
 
   auto stride = explicit_stride ? explicit_stride : implicit_stride;
 
-  // WebGPU requires runtime arrays have at least one element, but the AST
-  // records an element count of 0 for it.
-  auto size = std::max<uint32_t>(arr->size(), 1) * stride;
-  auto* out = builder_->create<sem::Array>(elem_type, arr->size(), el_align,
-                                           size, stride, implicit_stride);
+  // Evaluate the constant array size expression.
+  // sem::Array uses a size of 0 for a runtime-sized array.
+  uint32_t count = 0;
+  if (auto* count_expr = arr->Size()) {
+    Mark(count_expr);
+    if (!Expression(count_expr)) {
+      return nullptr;
+    }
+
+    auto size_source = count_expr->source();
+
+    auto* ty = TypeOf(count_expr)->UnwrapRef();
+    if (!ty->is_integer_scalar()) {
+      AddError("array size must be integer scalar", size_source);
+      return nullptr;
+    }
+
+    if (auto* ident = count_expr->As<ast::IdentifierExpression>()) {
+      // Make sure the identifier is a non-overridable module-scope constant.
+      VariableInfo* var = nullptr;
+      bool is_global = false;
+      if (!variable_stack_.get(ident->symbol(), &var, &is_global) ||
+          !is_global || !var->declaration->is_const()) {
+        AddError("array size identifier must be a module-scope constant",
+                 size_source);
+        return nullptr;
+      }
+      if (ast::HasDecoration<ast::OverrideDecoration>(
+              var->declaration->decorations())) {
+        AddError("array size expression must not be pipeline-overridable",
+                 size_source);
+        return nullptr;
+      }
+
+      count_expr = var->declaration->constructor();
+    } else if (!count_expr->Is<ast::ScalarConstructorExpression>()) {
+      AddError(
+          "array size expression must be either a literal or a module-scope "
+          "constant",
+          size_source);
+      return nullptr;
+    }
+
+    auto count_val = ConstantValueOf(count_expr);
+    if (!count_val) {
+      TINT_ICE(Resolver, diagnostics_)
+          << "could not resolve array size expression";
+      return nullptr;
+    }
+
+    if (ty->is_signed_integer_scalar() ? count_val.Elements()[0].i32 < 1
+                                       : count_val.Elements()[0].u32 < 1u) {
+      AddError("array size must be at least 1", size_source);
+      return nullptr;
+    }
+
+    count = count_val.Elements()[0].u32;
+  }
+
+  auto size = std::max<uint32_t>(count, 1) * stride;
+  auto* out = builder_->create<sem::Array>(elem_type, count, el_align, size,
+                                           stride, implicit_stride);
 
   if (!ValidateArray(out, source)) {
     return nullptr;
diff --git a/src/resolver/resolver_test.cc b/src/resolver/resolver_test.cc
index d84dba8505..108a3bc419 100644
--- a/src/resolver/resolver_test.cc
+++ b/src/resolver/resolver_test.cc
@@ -431,6 +431,66 @@ TEST_F(ResolverTest, Stmt_VariableDecl_ModuleScopeAfterFunctionScope) {
   EXPECT_EQ(VarOf(fn_f32->constructor())->Declaration(), mod_f32);
 }
 
+TEST_F(ResolverTest, ArraySize_UnsignedLiteral) {
+  // var<private> a : array<f32, 10u>;
+  auto* a =
+      Global("a", ty.array(ty.f32(), Expr(10u)), ast::StorageClass::kPrivate);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(a), nullptr);
+  auto* ref = TypeOf(a)->As<sem::Reference>();
+  ASSERT_NE(ref, nullptr);
+  auto* ary = ref->StoreType()->As<sem::Array>();
+  EXPECT_EQ(ary->Count(), 10u);
+}
+
+TEST_F(ResolverTest, ArraySize_SignedLiteral) {
+  // var<private> a : array<f32, 10>;
+  auto* a =
+      Global("a", ty.array(ty.f32(), Expr(10)), ast::StorageClass::kPrivate);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(a), nullptr);
+  auto* ref = TypeOf(a)->As<sem::Reference>();
+  ASSERT_NE(ref, nullptr);
+  auto* ary = ref->StoreType()->As<sem::Array>();
+  EXPECT_EQ(ary->Count(), 10u);
+}
+
+TEST_F(ResolverTest, ArraySize_UnsignedConstant) {
+  // let size = 0u;
+  // var<private> a : array<f32, 10u>;
+  GlobalConst("size", nullptr, Expr(10u));
+  auto* a = Global("a", ty.array(ty.f32(), Expr("size")),
+                   ast::StorageClass::kPrivate);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(a), nullptr);
+  auto* ref = TypeOf(a)->As<sem::Reference>();
+  ASSERT_NE(ref, nullptr);
+  auto* ary = ref->StoreType()->As<sem::Array>();
+  EXPECT_EQ(ary->Count(), 10u);
+}
+
+TEST_F(ResolverTest, ArraySize_SignedConstant) {
+  // let size = 0;
+  // var<private> a : array<f32, 10>;
+  GlobalConst("size", nullptr, Expr(10));
+  auto* a = Global("a", ty.array(ty.f32(), Expr("size")),
+                   ast::StorageClass::kPrivate);
+
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+
+  ASSERT_NE(TypeOf(a), nullptr);
+  auto* ref = TypeOf(a)->As<sem::Reference>();
+  ASSERT_NE(ref, nullptr);
+  auto* ary = ref->StoreType()->As<sem::Array>();
+  EXPECT_EQ(ary->Count(), 10u);
+}
+
 TEST_F(ResolverTest, Expr_ArrayAccessor_Array) {
   auto* idx = Expr(2);
   Global("my_var", ty.array<f32, 3>(), ast::StorageClass::kPrivate);
diff --git a/src/resolver/type_constructor_validation_test.cc b/src/resolver/type_constructor_validation_test.cc
index 295b61f51e..0cd2e3f5d1 100644
--- a/src/resolver/type_constructor_validation_test.cc
+++ b/src/resolver/type_constructor_validation_test.cc
@@ -584,8 +584,9 @@ TEST_F(ResolverTypeConstructorValidationTest,
 }
 
 TEST_F(ResolverTypeConstructorValidationTest, Expr_Constructor_Array_Runtime) {
-  // array<f32>(1);
-  auto* tc = array<i32>(
+  // array<i32>(1);
+  auto* tc = array(
+      ty.i32(), nullptr,
       create<ast::ScalarConstructorExpression>(Source{{12, 34}}, Literal(1)));
   WrapInFunction(tc);
 
@@ -595,8 +596,8 @@ TEST_F(ResolverTypeConstructorValidationTest, Expr_Constructor_Array_Runtime) {
 
 TEST_F(ResolverTypeConstructorValidationTest,
        Expr_Constructor_Array_RuntimeZeroValue) {
-  // array<f32>();
-  auto* tc = array<i32>();
+  // array<i32>();
+  auto* tc = array(ty.i32(), nullptr);
   WrapInFunction(tc);
 
   EXPECT_FALSE(r()->Resolve());
diff --git a/src/resolver/type_validation_test.cc b/src/resolver/type_validation_test.cc
index 52925b4f54..549b814713 100644
--- a/src/resolver/type_validation_test.cc
+++ b/src/resolver/type_validation_test.cc
@@ -201,6 +201,180 @@ TEST_F(ResolverTypeValidationTest,
   EXPECT_TRUE(r()->Resolve()) << r()->error();
 }
 
+TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedLiteral_Pass) {
+  // var<private> a : array<f32, 4u>;
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, 4u)),
+         ast::StorageClass::kPrivate);
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedLiteral_Pass) {
+  // var<private> a : array<f32, 4>;
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, 4)),
+         ast::StorageClass::kPrivate);
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedConstant_Pass) {
+  // let size = 4u;
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Expr(4u));
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedConstant_Pass) {
+  // let size = 4;
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Expr(4));
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedLiteral_Zero) {
+  // var<private> a : array<f32, 0u>;
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, 0u)),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be at least 1");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedLiteral_Zero) {
+  // var<private> a : array<f32, 0>;
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, 0)),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be at least 1");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedLiteral_Negative) {
+  // var<private> a : array<f32, -10>;
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, -10)),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be at least 1");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_UnsignedConstant_Zero) {
+  // let size = 0u;
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Expr(0u));
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be at least 1");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedConstant_Zero) {
+  // let size = 0;
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Expr(0));
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be at least 1");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_SignedConstant_Negative) {
+  // let size = -10;
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Expr(-10));
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be at least 1");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_FloatLiteral) {
+  // var<private> a : array<f32, 10.0>;
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, 10.f)),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be integer scalar");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_IVecLiteral) {
+  // var<private> a : array<f32, vec2<i32>(10, 10)>;
+  Global(
+      "a",
+      ty.array(ty.f32(), Construct(Source{{12, 34}}, ty.vec2<i32>(), 10, 10)),
+      ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be integer scalar");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_FloatConstant) {
+  // let size = 10.0;
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Expr(10.f));
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be integer scalar");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_IVecConstant) {
+  // let size = vec2<i32>(100, 100);
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Construct(ty.vec2<i32>(), 100, 100));
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(), "12:34 error: array size must be integer scalar");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_OverridableConstant) {
+  // [[override]] let size = 10;
+  // var<private> a : array<f32, size>;
+  GlobalConst("size", nullptr, Expr(10), {Override()});
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      "12:34 error: array size expression must not be pipeline-overridable");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_ModuleVar) {
+  // var<private> size : i32 = 10;
+  // var<private> a : array<f32, size>;
+  Global("size", ty.i32(), Expr(10), ast::StorageClass::kPrivate);
+  Global("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")),
+         ast::StorageClass::kPrivate);
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      "12:34 error: array size identifier must be a module-scope constant");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_FunctionConstant) {
+  // {
+  //   let size = 10;
+  //   var a : array<f32, size>;
+  // }
+  auto* size = Const("size", nullptr, Expr(10));
+  auto* a = Var("a", ty.array(ty.f32(), Expr(Source{{12, 34}}, "size")));
+  WrapInFunction(Block(Decl(size), Decl(a)));
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(
+      r()->error(),
+      "12:34 error: array size identifier must be a module-scope constant");
+}
+
+TEST_F(ResolverTypeValidationTest, ArraySize_InvalidExpr) {
+  // var a : array<f32, i32(4)>;
+  auto* size = Const("size", nullptr, Expr(10));
+  auto* a =
+      Var("a", ty.array(ty.f32(), Construct(Source{{12, 34}}, ty.i32(), 4)));
+  WrapInFunction(Block(Decl(size), Decl(a)));
+  EXPECT_FALSE(r()->Resolve());
+  EXPECT_EQ(r()->error(),
+            "12:34 error: array size expression must be either a literal or a "
+            "module-scope constant");
+}
+
 TEST_F(ResolverTypeValidationTest, RuntimeArrayInFunction_Fail) {
   /// [[stage(vertex)]]
   // fn func() { var a : array<i32>; }
diff --git a/src/transform/array_length_from_uniform_test.cc b/src/transform/array_length_from_uniform_test.cc
index 6ab39ee055..c41ba2e932 100644
--- a/src/transform/array_length_from_uniform_test.cc
+++ b/src/transform/array_length_from_uniform_test.cc
@@ -81,7 +81,7 @@ fn main() {
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<vec4<u32>, 1>;
+  buffer_size : array<vec4<u32>, 1u>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -134,7 +134,7 @@ fn main() {
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<vec4<u32>, 1>;
+  buffer_size : array<vec4<u32>, 1u>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
@@ -216,7 +216,7 @@ fn main() {
   auto* expect = R"(
 [[block]]
 struct tint_symbol {
-  buffer_size : array<vec4<u32>, 2>;
+  buffer_size : array<vec4<u32>, 2u>;
 };
 
 [[group(0), binding(30)]] var<uniform> tint_symbol_1 : tint_symbol;
diff --git a/src/transform/decompose_memory_access_test.cc b/src/transform/decompose_memory_access_test.cc
index d4e54bdcf4..37fcc27485 100644
--- a/src/transform/decompose_memory_access_test.cc
+++ b/src/transform/decompose_memory_access_test.cc
@@ -180,8 +180,8 @@ fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_p
   return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
 }
 
-fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
-  var arr : array<vec3<f32>, 2>;
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
+  var arr : array<vec3<f32>, 2u>;
   for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
     arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
   }
@@ -378,8 +378,8 @@ fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_p
   return mat4x4<f32>(tint_symbol_11(buffer, (offset + 0u)), tint_symbol_11(buffer, (offset + 16u)), tint_symbol_11(buffer, (offset + 32u)), tint_symbol_11(buffer, (offset + 48u)));
 }
 
-fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> array<vec3<f32>, 2> {
-  var arr : array<vec3<f32>, 2>;
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : UB, offset : u32) -> array<vec3<f32>, 2u> {
+  var arr : array<vec3<f32>, 2u>;
   for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
     arr[i_1] = tint_symbol_8(buffer, (offset + (i_1 * 16u)));
   }
@@ -594,7 +594,7 @@ fn tint_symbol_20([[internal(disable_validation__ignore_constructible_function_p
   tint_symbol_11(buffer, (offset + 48u), value[3u]);
 }
 
-fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
+fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
   var array = value;
   for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
     tint_symbol_8(buffer, (offset + (i_1 * 16u)), array[i_1]);
@@ -770,8 +770,8 @@ fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_p
   return mat4x4<f32>(tint_symbol_12(buffer, (offset + 0u)), tint_symbol_12(buffer, (offset + 16u)), tint_symbol_12(buffer, (offset + 32u)), tint_symbol_12(buffer, (offset + 48u)));
 }
 
-fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2> {
-  var arr : array<vec3<f32>, 2>;
+fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32) -> array<vec3<f32>, 2u> {
+  var arr : array<vec3<f32>, 2u>;
   for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
     arr[i_1] = tint_symbol_9(buffer, (offset + (i_1 * 16u)));
   }
@@ -948,7 +948,7 @@ fn tint_symbol_21([[internal(disable_validation__ignore_constructible_function_p
   tint_symbol_12(buffer, (offset + 48u), value[3u]);
 }
 
-fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2>) {
+fn tint_symbol_22([[internal(disable_validation__ignore_constructible_function_parameter)]] buffer : SB, offset : u32, value : array<vec3<f32>, 2u>) {
   var array = value;
   for(var i_1 = 0u; (i_1 < 2u); i_1 = (i_1 + 1u)) {
     tint_symbol_9(buffer, (offset + (i_1 * 16u)), array[i_1]);
diff --git a/src/transform/decompose_strided_matrix_test.cc b/src/transform/decompose_strided_matrix_test.cc
index 6a205a67a8..8f5d13e88b 100644
--- a/src/transform/decompose_strided_matrix_test.cc
+++ b/src/transform/decompose_strided_matrix_test.cc
@@ -106,12 +106,12 @@ TEST_F(DecomposeStridedMatrixTest, ReadUniformMatrix) {
 struct S {
   [[size(16)]]
   padding : u32;
-  m : [[stride(32)]] array<vec2<f32>, 2>;
+  m : [[stride(32)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<uniform> s : S;
 
-fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2>) -> mat2x2<f32> {
+fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2u>) -> mat2x2<f32> {
   return mat2x2<f32>(arr[0u], arr[1u]);
 }
 
@@ -173,7 +173,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadUniformColumn) {
 struct S {
   [[size(16)]]
   padding : u32;
-  m : [[stride(32)]] array<vec2<f32>, 2>;
+  m : [[stride(32)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<uniform> s : S;
@@ -300,12 +300,12 @@ TEST_F(DecomposeStridedMatrixTest, ReadStorageMatrix) {
 struct S {
   [[size(8)]]
   padding : u32;
-  m : [[stride(32)]] array<vec2<f32>, 2>;
+  m : [[stride(32)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<storage, read_write> s : S;
 
-fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2>) -> mat2x2<f32> {
+fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2u>) -> mat2x2<f32> {
   return mat2x2<f32>(arr[0u], arr[1u]);
 }
 
@@ -367,7 +367,7 @@ TEST_F(DecomposeStridedMatrixTest, ReadStorageColumn) {
 struct S {
   [[size(16)]]
   padding : u32;
-  m : [[stride(32)]] array<vec2<f32>, 2>;
+  m : [[stride(32)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<storage, read_write> s : S;
@@ -431,13 +431,13 @@ TEST_F(DecomposeStridedMatrixTest, WriteStorageMatrix) {
 struct S {
   [[size(8)]]
   padding : u32;
-  m : [[stride(32)]] array<vec2<f32>, 2>;
+  m : [[stride(32)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<storage, read_write> s : S;
 
-fn mat2x2_stride_32_to_arr(mat : mat2x2<f32>) -> [[stride(32)]] array<vec2<f32>, 2> {
-  return [[stride(32)]] array<vec2<f32>, 2>(mat[0u], mat[1u]);
+fn mat2x2_stride_32_to_arr(mat : mat2x2<f32>) -> [[stride(32)]] array<vec2<f32>, 2u> {
+  return [[stride(32)]] array<vec2<f32>, 2u>(mat[0u], mat[1u]);
 }
 
 [[stage(compute), workgroup_size(1)]]
@@ -498,7 +498,7 @@ TEST_F(DecomposeStridedMatrixTest, WriteStorageColumn) {
 struct S {
   [[size(8)]]
   padding : u32;
-  m : [[stride(32)]] array<vec2<f32>, 2>;
+  m : [[stride(32)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<storage, read_write> s : S;
@@ -576,17 +576,17 @@ TEST_F(DecomposeStridedMatrixTest, ReadWriteViaPointerLets) {
 struct S {
   [[size(8)]]
   padding : u32;
-  m : [[stride(32)]] array<vec2<f32>, 2>;
+  m : [[stride(32)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<storage, read_write> s : S;
 
-fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2>) -> mat2x2<f32> {
+fn arr_to_mat2x2_stride_32(arr : [[stride(32)]] array<vec2<f32>, 2u>) -> mat2x2<f32> {
   return mat2x2<f32>(arr[0u], arr[1u]);
 }
 
-fn mat2x2_stride_32_to_arr(mat : mat2x2<f32>) -> [[stride(32)]] array<vec2<f32>, 2> {
-  return [[stride(32)]] array<vec2<f32>, 2>(mat[0u], mat[1u]);
+fn mat2x2_stride_32_to_arr(mat : mat2x2<f32>) -> [[stride(32)]] array<vec2<f32>, 2u> {
+  return [[stride(32)]] array<vec2<f32>, 2u>(mat[0u], mat[1u]);
 }
 
 [[stage(compute), workgroup_size(1)]]
diff --git a/src/transform/pad_array_elements.cc b/src/transform/pad_array_elements.cc
index 63bcd3af36..04358bddb7 100644
--- a/src/transform/pad_array_elements.cc
+++ b/src/transform/pad_array_elements.cc
@@ -80,7 +80,11 @@ ArrayBuilder PadArray(
 
     auto* dst = ctx.dst;
     return [=] {
-      return dst->ty.array(dst->create<ast::TypeName>(name), array->Count());
+      if (array->IsRuntimeSized()) {
+        return dst->ty.array(dst->create<ast::TypeName>(name));
+      } else {
+        return dst->ty.array(dst->create<ast::TypeName>(name), array->Count());
+      }
     };
   });
 }
diff --git a/src/transform/pad_array_elements_test.cc b/src/transform/pad_array_elements_test.cc
index 8e5fdf3559..eca088c5db 100644
--- a/src/transform/pad_array_elements_test.cc
+++ b/src/transform/pad_array_elements_test.cc
@@ -54,7 +54,31 @@ struct tint_padded_array_element {
   el : i32;
 };
 
-var<private> arr : array<tint_padded_array_element, 4>;
+var<private> arr : array<tint_padded_array_element, 4u>;
+)";
+
+  auto got = Run<PadArrayElements>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(PadArrayElementsTest, RuntimeArray) {
+  auto* src = R"(
+[[block]]
+struct S {
+  rta : [[stride(8)]] array<i32>;
+};
+)";
+  auto* expect = R"(
+struct tint_padded_array_element {
+  [[size(8)]]
+  el : i32;
+};
+
+[[block]]
+struct S {
+  rta : array<tint_padded_array_element>;
+};
 )";
 
   auto got = Run<PadArrayElements>(src);
@@ -78,9 +102,9 @@ struct tint_padded_array_element {
 };
 
 fn f() {
-  var arr : array<tint_padded_array_element, 4>;
-  arr = array<tint_padded_array_element, 4>();
-  arr = array<tint_padded_array_element, 4>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
+  var arr : array<tint_padded_array_element, 4u>;
+  arr = array<tint_padded_array_element, 4u>();
+  arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
   let x = arr[3].el;
 }
 )";
@@ -102,7 +126,7 @@ struct tint_padded_array_element {
   el : i32;
 };
 
-fn f(a : array<tint_padded_array_element, 4>) -> i32 {
+fn f(a : array<tint_padded_array_element, 4u>) -> i32 {
   return a[2].el;
 }
 )";
@@ -155,13 +179,13 @@ struct tint_padded_array_element {
   el : i32;
 };
 
-type Array = array<tint_padded_array_element, 4>;
+type Array = array<tint_padded_array_element, 4u>;
 
 fn f() {
-  var arr : array<tint_padded_array_element, 4>;
-  arr = array<tint_padded_array_element, 4>();
-  arr = array<tint_padded_array_element, 4>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
-  let vals : array<tint_padded_array_element, 4> = array<tint_padded_array_element, 4>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
+  var arr : array<tint_padded_array_element, 4u>;
+  arr = array<tint_padded_array_element, 4u>();
+  arr = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
+  let vals : array<tint_padded_array_element, 4u> = array<tint_padded_array_element, 4u>(tint_padded_array_element(1), tint_padded_array_element(2), tint_padded_array_element(3), tint_padded_array_element(4));
   arr = vals;
   let x = arr[3].el;
 }
@@ -198,10 +222,10 @@ struct tint_padded_array_element_2 {
 };
 
 struct S {
-  a : array<tint_padded_array_element, 4>;
-  b : array<tint_padded_array_element_1, 8>;
-  c : array<tint_padded_array_element, 4>;
-  d : array<tint_padded_array_element_2, 8>;
+  a : array<tint_padded_array_element, 4u>;
+  b : array<tint_padded_array_element_1, 8u>;
+  c : array<tint_padded_array_element, 4u>;
+  d : array<tint_padded_array_element_2, 8u>;
 };
 )";
 
@@ -231,7 +255,7 @@ struct tint_padded_array_element_2 {
 
 struct tint_padded_array_element_1 {
   [[size(512)]]
-  el : array<tint_padded_array_element_2, 4>;
+  el : array<tint_padded_array_element_2, 4u>;
 };
 
 struct tint_padded_array_element_5 {
@@ -241,18 +265,18 @@ struct tint_padded_array_element_5 {
 
 struct tint_padded_array_element_4 {
   [[size(64)]]
-  el : array<tint_padded_array_element_5, 4>;
+  el : array<tint_padded_array_element_5, 4u>;
 };
 
 struct tint_padded_array_element_3 {
   [[size(512)]]
-  el : array<tint_padded_array_element_4, 4>;
+  el : array<tint_padded_array_element_4, 4u>;
 };
 
 struct S {
-  a : array<tint_padded_array_element, 4>;
-  b : array<tint_padded_array_element_1, 4>;
-  c : array<tint_padded_array_element_3, 4>;
+  a : array<tint_padded_array_element, 4u>;
+  b : array<tint_padded_array_element_1, 4u>;
+  c : array<tint_padded_array_element_3, 4u>;
 };
 )";
 
@@ -286,7 +310,7 @@ struct tint_padded_array_element_2 {
 
 struct tint_padded_array_element_1 {
   [[size(512)]]
-  el : array<tint_padded_array_element_2, 4>;
+  el : array<tint_padded_array_element_2, 4u>;
 };
 
 struct tint_padded_array_element_5 {
@@ -296,18 +320,18 @@ struct tint_padded_array_element_5 {
 
 struct tint_padded_array_element_4 {
   [[size(64)]]
-  el : array<tint_padded_array_element_5, 4>;
+  el : array<tint_padded_array_element_5, 4u>;
 };
 
 struct tint_padded_array_element_3 {
   [[size(512)]]
-  el : array<tint_padded_array_element_4, 4>;
+  el : array<tint_padded_array_element_4, 4u>;
 };
 
 struct S {
-  a : array<tint_padded_array_element, 4>;
-  b : array<tint_padded_array_element_1, 4>;
-  c : array<tint_padded_array_element_3, 4>;
+  a : array<tint_padded_array_element, 4u>;
+  b : array<tint_padded_array_element_1, 4u>;
+  c : array<tint_padded_array_element_3, 4u>;
 };
 
 fn f(s : S) -> i32 {
@@ -345,7 +369,7 @@ struct tint_padded_array_element {
   el : i32;
 };
 
-type T1 = array<tint_padded_array_element, 1>;
+type T1 = array<tint_padded_array_element, 1u>;
 
 type T2 = i32;
 
@@ -354,7 +378,7 @@ struct tint_padded_array_element_1 {
   el : i32;
 };
 
-fn f1(a : array<tint_padded_array_element_1, 2>) {
+fn f1(a : array<tint_padded_array_element_1, 2u>) {
 }
 
 type T3 = i32;
@@ -365,7 +389,7 @@ struct tint_padded_array_element_2 {
 };
 
 fn f2() {
-  var v : array<tint_padded_array_element_2, 3>;
+  var v : array<tint_padded_array_element_2, 3u>;
 }
 )";
 
diff --git a/src/transform/spirv_test.cc b/src/transform/spirv_test.cc
index 25678ebd4b..6f68d493c3 100644
--- a/src/transform/spirv_test.cc
+++ b/src/transform/spirv_test.cc
@@ -35,9 +35,9 @@ fn main([[builtin(sample_index)]] sample_index : u32,
   auto* expect = R"(
 [[builtin(sample_index), internal(disable_validation__ignore_storage_class)]] var<in> sample_index_1 : u32;
 
-[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<in> mask_in_1 : array<u32, 1>;
+[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<in> mask_in_1 : array<u32, 1u>;
 
-[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<out> value : array<u32, 1>;
+[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<out> value : array<u32, 1u>;
 
 fn main_inner(sample_index : u32, mask_in : u32) -> u32 {
   return mask_in;
@@ -73,9 +73,9 @@ fn main([[builtin(sample_mask)]] mask_in : u32
 )";
 
   auto* expect = R"(
-[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<in> mask_in_1 : array<u32, 1>;
+[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<in> mask_in_1 : array<u32, 1u>;
 
-[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<out> value : array<u32, 1>;
+[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<out> value : array<u32, 1u>;
 
 fn filter(mask : u32) -> u32 {
   return (mask & 3u);
@@ -138,9 +138,9 @@ fn frag_main([[builtin(sample_index)]] sample_index : u32,
 
 [[builtin(sample_index), internal(disable_validation__ignore_storage_class)]] var<in> sample_index_1 : u32;
 
-[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<in> mask_in_1 : array<u32, 1>;
+[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<in> mask_in_1 : array<u32, 1u>;
 
-[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<out> value_1 : array<u32, 1>;
+[[builtin(sample_mask), internal(disable_validation__ignore_storage_class)]] var<out> value_1 : array<u32, 1u>;
 
 fn vert_main_inner() -> vec4<f32> {
   return vec4<f32>();
diff --git a/src/transform/transform.cc b/src/transform/transform.cc
index db6e47eb11..8b199a5574 100644
--- a/src/transform/transform.cc
+++ b/src/transform/transform.cc
@@ -120,7 +120,11 @@ ast::Type* Transform::CreateASTTypeFor(CloneContext& ctx, const sem::Type* ty) {
     if (!a->IsStrideImplicit()) {
       decos.emplace_back(ctx.dst->create<ast::StrideDecoration>(a->Stride()));
     }
-    return ctx.dst->create<ast::Array>(el, a->Count(), std::move(decos));
+    if (a->IsRuntimeSized()) {
+      return ctx.dst->ty.array(el, nullptr, std::move(decos));
+    } else {
+      return ctx.dst->ty.array(el, a->Count(), std::move(decos));
+    }
   }
   if (auto* s = ty->As<sem::Struct>()) {
     return ctx.dst->create<ast::TypeName>(ctx.Clone(s->Declaration()->name()));
diff --git a/src/transform/transform_test.cc b/src/transform/transform_test.cc
index 87b6ab261e..7cd0cdc513 100644
--- a/src/transform/transform_test.cc
+++ b/src/transform/transform_test.cc
@@ -82,8 +82,14 @@ TEST_F(CreateASTTypeForTest, ArrayImplicitStride) {
   });
   ASSERT_TRUE(arr->Is<ast::Array>());
   ASSERT_TRUE(arr->As<ast::Array>()->type()->Is<ast::F32>());
-  ASSERT_EQ(arr->As<ast::Array>()->size(), 2u);
   ASSERT_EQ(arr->As<ast::Array>()->decorations().size(), 0u);
+
+  auto* size_expr =
+      arr->As<ast::Array>()->Size()->As<ast::ScalarConstructorExpression>();
+  ASSERT_NE(size_expr, nullptr);
+  auto* size = size_expr->literal()->As<ast::IntLiteral>();
+  ASSERT_NE(size, nullptr);
+  EXPECT_EQ(size->value_as_i32(), 2);
 }
 
 TEST_F(CreateASTTypeForTest, ArrayNonImplicitStride) {
@@ -92,7 +98,6 @@ TEST_F(CreateASTTypeForTest, ArrayNonImplicitStride) {
   });
   ASSERT_TRUE(arr->Is<ast::Array>());
   ASSERT_TRUE(arr->As<ast::Array>()->type()->Is<ast::F32>());
-  ASSERT_EQ(arr->As<ast::Array>()->size(), 2u);
   ASSERT_EQ(arr->As<ast::Array>()->decorations().size(), 1u);
   ASSERT_TRUE(
       arr->As<ast::Array>()->decorations()[0]->Is<ast::StrideDecoration>());
@@ -101,6 +106,13 @@ TEST_F(CreateASTTypeForTest, ArrayNonImplicitStride) {
                 ->As<ast::StrideDecoration>()
                 ->stride(),
             64u);
+
+  auto* size_expr =
+      arr->As<ast::Array>()->Size()->As<ast::ScalarConstructorExpression>();
+  ASSERT_NE(size_expr, nullptr);
+  auto* size = size_expr->literal()->As<ast::IntLiteral>();
+  ASSERT_NE(size, nullptr);
+  EXPECT_EQ(size->value_as_i32(), 2);
 }
 
 TEST_F(CreateASTTypeForTest, Struct) {
diff --git a/src/transform/vertex_pulling.cc b/src/transform/vertex_pulling.cc
index d4a8d5e839..a44ba63d50 100644
--- a/src/transform/vertex_pulling.cc
+++ b/src/transform/vertex_pulling.cc
@@ -259,7 +259,7 @@ struct State {
         ctx.dst->Symbols().New(kStructName),
         {
             ctx.dst->Member(GetStructBufferName(),
-                            ctx.dst->ty.array<ProgramBuilder::u32, 0>(4)),
+                            ctx.dst->ty.array<ProgramBuilder::u32>(4)),
         },
         {
             ctx.dst->create<ast::StructBlockDecoration>(),
diff --git a/src/transform/wrap_arrays_in_structs.cc b/src/transform/wrap_arrays_in_structs.cc
index c16503c301..b3c315559d 100644
--- a/src/transform/wrap_arrays_in_structs.cc
+++ b/src/transform/wrap_arrays_in_structs.cc
@@ -132,8 +132,7 @@ WrapArraysInStructs::WrappedArrayInfo WrapArraysInStructs::WrapArray(
         decos.emplace_back(
             c.dst->create<ast::StrideDecoration>(array->Stride()));
       }
-      return c.dst->create<ast::Array>(el_type(c), array->Count(),
-                                       std::move(decos));
+      return c.dst->ty.array(el_type(c), array->Count(), std::move(decos));
     };
 
     // Structure() will create and append the ast::Struct to the
diff --git a/src/transform/wrap_arrays_in_structs_test.cc b/src/transform/wrap_arrays_in_structs_test.cc
index 4012e95973..a93f995726 100644
--- a/src/transform/wrap_arrays_in_structs_test.cc
+++ b/src/transform/wrap_arrays_in_structs_test.cc
@@ -40,7 +40,7 @@ var<private> arr : array<i32, 4>;
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 4>;
+  arr : array<i32, 4u>;
 };
 
 var<private> arr : tint_array_wrapper;
@@ -60,7 +60,7 @@ fn f() {
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 4>;
+  arr : array<i32, 4u>;
 };
 
 fn f() {
@@ -82,7 +82,7 @@ fn f(a : array<i32, 4>) -> i32 {
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 4>;
+  arr : array<i32, 4u>;
 };
 
 fn f(a : tint_array_wrapper) -> i32 {
@@ -103,11 +103,11 @@ fn f() -> array<i32, 4> {
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 4>;
+  arr : array<i32, 4u>;
 };
 
 fn f() -> tint_array_wrapper {
-  return tint_array_wrapper(array<i32, 4>(1, 2, 3, 4));
+  return tint_array_wrapper(array<i32, 4u>(1, 2, 3, 4));
 }
 )";
 
@@ -132,22 +132,22 @@ fn f() {
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 2>;
+  arr : array<i32, 2u>;
 };
 
 type Inner = tint_array_wrapper;
 
 struct tint_array_wrapper_1 {
-  arr : array<tint_array_wrapper, 2>;
+  arr : array<tint_array_wrapper, 2u>;
 };
 
 type Array = tint_array_wrapper_1;
 
 fn f() {
   var arr : tint_array_wrapper_1;
-  arr = tint_array_wrapper_1(array<tint_array_wrapper, 2>());
-  arr = tint_array_wrapper_1(array<tint_array_wrapper, 2>(tint_array_wrapper(array<i32, 2>(1, 2)), tint_array_wrapper(array<i32, 2>(3, 4))));
-  let vals : tint_array_wrapper_1 = tint_array_wrapper_1(array<tint_array_wrapper, 2>(tint_array_wrapper(array<i32, 2>(1, 2)), tint_array_wrapper(array<i32, 2>(3, 4))));
+  arr = tint_array_wrapper_1(array<tint_array_wrapper, 2u>());
+  arr = tint_array_wrapper_1(array<tint_array_wrapper, 2u>(tint_array_wrapper(array<i32, 2u>(1, 2)), tint_array_wrapper(array<i32, 2u>(3, 4))));
+  let vals : tint_array_wrapper_1 = tint_array_wrapper_1(array<tint_array_wrapper, 2u>(tint_array_wrapper(array<i32, 2u>(1, 2)), tint_array_wrapper(array<i32, 2u>(3, 4))));
   arr = vals;
   let x = arr.arr[3];
 }
@@ -168,11 +168,11 @@ struct S {
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 4>;
+  arr : array<i32, 4u>;
 };
 
 struct tint_array_wrapper_1 {
-  arr : array<i32, 8>;
+  arr : array<i32, 8u>;
 };
 
 struct S {
@@ -197,15 +197,15 @@ struct S {
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 4>;
+  arr : array<i32, 4u>;
 };
 
 struct tint_array_wrapper_1 {
-  arr : array<tint_array_wrapper, 4>;
+  arr : array<tint_array_wrapper, 4u>;
 };
 
 struct tint_array_wrapper_2 {
-  arr : array<tint_array_wrapper_1, 4>;
+  arr : array<tint_array_wrapper_1, 4u>;
 };
 
 struct S {
@@ -234,15 +234,15 @@ fn f(s : S) -> i32 {
 )";
   auto* expect = R"(
 struct tint_array_wrapper {
-  arr : array<i32, 4>;
+  arr : array<i32, 4u>;
 };
 
 struct tint_array_wrapper_1 {
-  arr : array<tint_array_wrapper, 4>;
+  arr : array<tint_array_wrapper, 4u>;
 };
 
 struct tint_array_wrapper_2 {
-  arr : array<tint_array_wrapper_1, 4>;
+  arr : array<tint_array_wrapper_1, 4u>;
 };
 
 struct S {
@@ -282,7 +282,7 @@ fn f2() {
 type T0 = i32;
 
 struct tint_array_wrapper {
-  arr : array<i32, 1>;
+  arr : array<i32, 1u>;
 };
 
 type T1 = tint_array_wrapper;
@@ -290,7 +290,7 @@ type T1 = tint_array_wrapper;
 type T2 = i32;
 
 struct tint_array_wrapper_1 {
-  arr : array<i32, 2>;
+  arr : array<i32, 2u>;
 };
 
 fn f1(a : tint_array_wrapper_1) {
@@ -299,7 +299,7 @@ fn f1(a : tint_array_wrapper_1) {
 type T3 = i32;
 
 struct tint_array_wrapper_2 {
-  arr : array<i32, 3>;
+  arr : array<i32, 3u>;
 };
 
 fn f2() {
diff --git a/src/writer/hlsl/generator_impl_type_test.cc b/src/writer/hlsl/generator_impl_type_test.cc
index 90df7cd86e..a5846e2a69 100644
--- a/src/writer/hlsl/generator_impl_type_test.cc
+++ b/src/writer/hlsl/generator_impl_type_test.cc
@@ -61,7 +61,7 @@ TEST_F(HlslGeneratorImplTest_Type, EmitType_ArrayOfArray) {
 // TODO(dsinclair): Is this possible? What order should it output in?
 TEST_F(HlslGeneratorImplTest_Type,
        DISABLED_EmitType_ArrayOfArrayOfRuntimeArray) {
-  auto* arr = ty.array(ty.array(ty.array<bool, 4>(), 5), 0);
+  auto* arr = ty.array(ty.array(ty.array<bool, 4>(), 5));
   Global("G", arr, ast::StorageClass::kPrivate);
 
   GeneratorImpl& gen = Build();
diff --git a/src/writer/spirv/builder_type_test.cc b/src/writer/spirv/builder_type_test.cc
index e178020322..88fb703f08 100644
--- a/src/writer/spirv/builder_type_test.cc
+++ b/src/writer/spirv/builder_type_test.cc
@@ -27,7 +27,7 @@ namespace {
 using BuilderTest_Type = TestHelper;
 
 TEST_F(BuilderTest_Type, GenerateRuntimeArray) {
-  auto* ary = ty.array(ty.i32(), 0);
+  auto* ary = ty.array(ty.i32());
   auto* str = Structure("S", {Member("x", ary)},
                         {create<ast::StructBlockDecoration>()});
   Global("a", ty.Of(str), ast::StorageClass::kStorage, ast::Access::kRead,
@@ -48,7 +48,7 @@ TEST_F(BuilderTest_Type, GenerateRuntimeArray) {
 }
 
 TEST_F(BuilderTest_Type, ReturnsGeneratedRuntimeArray) {
-  auto* ary = ty.array(ty.i32(), 0);
+  auto* ary = ty.array(ty.i32());
   auto* str = Structure("S", {Member("x", ary)},
                         {create<ast::StructBlockDecoration>()});
   Global("a", ty.Of(str), ast::StorageClass::kStorage, ast::Access::kRead,
@@ -413,7 +413,7 @@ TEST_F(BuilderTest_Type, GenerateStruct_DecoratedMembers_LayoutArraysOfMatrix) {
   // in levels of arrays.
   auto* arr_mat2x2 = ty.array(ty.mat2x2<f32>(), 1);      // Singly nested array
   auto* arr_arr_mat2x3 = ty.array(ty.mat2x3<f32>(), 1);  // Doubly nested array
-  auto* rtarr_mat4x4 = ty.array(ty.mat4x4<f32>(), 0);    // Runtime array
+  auto* rtarr_mat4x4 = ty.array(ty.mat4x4<f32>());       // Runtime array
 
   auto* s =
       Structure("S",
diff --git a/src/writer/wgsl/generator_impl.cc b/src/writer/wgsl/generator_impl.cc
index d2afa1fc6c..4db149d609 100644
--- a/src/writer/wgsl/generator_impl.cc
+++ b/src/writer/wgsl/generator_impl.cc
@@ -399,8 +399,12 @@ bool GeneratorImpl::EmitType(std::ostream& out, const ast::Type* ty) {
       return false;
     }
 
-    if (!ary->IsRuntimeArray())
-      out << ", " << ary->size();
+    if (!ary->IsRuntimeArray()) {
+      out << ", ";
+      if (!EmitExpression(out, ary->Size())) {
+        return false;
+      }
+    }
 
     out << ">";
   } else if (ty->Is<ast::Bool>()) {
diff --git a/src/writer/wgsl/generator_impl_type_test.cc b/src/writer/wgsl/generator_impl_type_test.cc
index ba8e599e85..0b422499e5 100644
--- a/src/writer/wgsl/generator_impl_type_test.cc
+++ b/src/writer/wgsl/generator_impl_type_test.cc
@@ -60,7 +60,7 @@ TEST_F(WgslGeneratorImplTest, EmitType_Array_Decoration) {
 }
 
 TEST_F(WgslGeneratorImplTest, EmitType_RuntimeArray) {
-  auto* a = ty.array(ty.bool_(), 0);
+  auto* a = ty.array(ty.bool_());
   Alias("make_type_reachable", a);
 
   GeneratorImpl& gen = Build();
diff --git a/test/array/size.wgsl b/test/array/size.wgsl
new file mode 100644
index 0000000000..c5ebee4531
--- /dev/null
+++ b/test/array/size.wgsl
@@ -0,0 +1,14 @@
+let slen = 4;
+let ulen = 4u;
+
+[[stage(fragment)]]
+fn main() {
+  var signed_literal : array<f32, 4>;
+  var unsigned_literal : array<f32, 4u>;
+  var signed_constant : array<f32, slen>;
+  var unsigned_constant : array<f32, ulen>;
+
+  // Ensure that the types are compatible.
+  signed_literal = unsigned_constant;
+  signed_constant = unsigned_literal;
+}
diff --git a/test/array/size.wgsl.expected.hlsl b/test/array/size.wgsl.expected.hlsl
new file mode 100644
index 0000000000..e5fdbd3c7e
--- /dev/null
+++ b/test/array/size.wgsl.expected.hlsl
@@ -0,0 +1,12 @@
+static const int slen = 4;
+static const uint ulen = 4u;
+
+void main() {
+  float signed_literal[4] = (float[4])0;
+  float unsigned_literal[4] = (float[4])0;
+  float signed_constant[4] = (float[4])0;
+  float unsigned_constant[4] = (float[4])0;
+  signed_literal = unsigned_constant;
+  signed_constant = unsigned_literal;
+  return;
+}
diff --git a/test/array/size.wgsl.expected.msl b/test/array/size.wgsl.expected.msl
new file mode 100644
index 0000000000..512a9b03d7
--- /dev/null
+++ b/test/array/size.wgsl.expected.msl
@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  float arr[4];
+};
+
+constant int slen = 4;
+constant uint ulen = 4u;
+fragment void tint_symbol() {
+  tint_array_wrapper signed_literal = {};
+  tint_array_wrapper unsigned_literal = {};
+  tint_array_wrapper signed_constant = {};
+  tint_array_wrapper unsigned_constant = {};
+  signed_literal = unsigned_constant;
+  signed_constant = unsigned_literal;
+  return;
+}
+
diff --git a/test/array/size.wgsl.expected.spvasm b/test/array/size.wgsl.expected.spvasm
new file mode 100644
index 0000000000..ada562b6c9
--- /dev/null
+++ b/test/array/size.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 19
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint Fragment %main "main"
+               OpExecutionMode %main OriginUpperLeft
+               OpName %slen "slen"
+               OpName %ulen "ulen"
+               OpName %main "main"
+               OpName %signed_literal "signed_literal"
+               OpName %unsigned_literal "unsigned_literal"
+               OpName %signed_constant "signed_constant"
+               OpName %unsigned_constant "unsigned_constant"
+               OpDecorate %_arr_float_ulen ArrayStride 4
+        %int = OpTypeInt 32 1
+       %slen = OpConstant %int 4
+       %uint = OpTypeInt 32 0
+       %ulen = OpConstant %uint 4
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+%_arr_float_ulen = OpTypeArray %float %ulen
+%_ptr_Function__arr_float_ulen = OpTypePointer Function %_arr_float_ulen
+         %13 = OpConstantNull %_arr_float_ulen
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+%signed_literal = OpVariable %_ptr_Function__arr_float_ulen Function %13
+%unsigned_literal = OpVariable %_ptr_Function__arr_float_ulen Function %13
+%signed_constant = OpVariable %_ptr_Function__arr_float_ulen Function %13
+%unsigned_constant = OpVariable %_ptr_Function__arr_float_ulen Function %13
+         %17 = OpLoad %_arr_float_ulen %unsigned_constant
+               OpStore %signed_literal %17
+         %18 = OpLoad %_arr_float_ulen %unsigned_literal
+               OpStore %signed_constant %18
+               OpReturn
+               OpFunctionEnd
diff --git a/test/array/size.wgsl.expected.wgsl b/test/array/size.wgsl.expected.wgsl
new file mode 100644
index 0000000000..f0df1f00b6
--- /dev/null
+++ b/test/array/size.wgsl.expected.wgsl
@@ -0,0 +1,13 @@
+let slen = 4;
+
+let ulen = 4u;
+
+[[stage(fragment)]]
+fn main() {
+  var signed_literal : array<f32, 4>;
+  var unsigned_literal : array<f32, 4u>;
+  var signed_constant : array<f32, slen>;
+  var unsigned_constant : array<f32, ulen>;
+  signed_literal = unsigned_constant;
+  signed_constant = unsigned_literal;
+}
diff --git a/test/bug/tint/1088.spvasm.expected.wgsl b/test/bug/tint/1088.spvasm.expected.wgsl
index 347b3c632c..ae16bd321e 100644
--- a/test/bug/tint/1088.spvasm.expected.wgsl
+++ b/test/bug/tint/1088.spvasm.expected.wgsl
@@ -1,6 +1,6 @@
-type Arr = [[stride(64)]] array<mat4x4<f32>, 2>;
+type Arr = [[stride(64)]] array<mat4x4<f32>, 2u>;
 
-type Arr_1 = [[stride(16)]] array<f32, 4>;
+type Arr_1 = [[stride(16)]] array<f32, 4u>;
 
 [[block]]
 struct LeftOver {
diff --git a/test/bug/tint/749.spvasm.expected.wgsl b/test/bug/tint/749.spvasm.expected.wgsl
index ef02a8319c..0334e152b2 100644
--- a/test/bug/tint/749.spvasm.expected.wgsl
+++ b/test/bug/tint/749.spvasm.expected.wgsl
@@ -1,5 +1,5 @@
 struct QuicksortObject {
-  numbers : array<i32, 10>;
+  numbers : array<i32, 10u>;
 };
 
 [[block]]
@@ -63,14 +63,14 @@ fn swap_i1_i1_(i : ptr<function, i32>, j : ptr<function, i32>) {
   let x_34 : ptr<private, i32> = &(obj.numbers[x_33]);
   let x_35 : i32 = *(x_34);
   let x_943 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_943;
   let x_527 : vec2<f32> = vec2<f32>(x_526.x, x_526.x);
   let x_36 : ptr<private, i32> = &(obj.numbers[x_32]);
   let x_528 : vec3<f32> = vec3<f32>(x_524.x, x_524.z, x_524.x);
   *(x_36) = x_35;
   let x_944 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_944;
   let x_529 : vec3<f32> = vec3<f32>(x_526.y, x_526.z, x_526.x);
   let x_945 : i32 = *(i);
@@ -93,7 +93,7 @@ fn swap_i1_i1_(i : ptr<function, i32>, j : ptr<function, i32>) {
   *(x_36) = 0;
   *(x_36) = x_949;
   let x_950 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_950;
   let x_532 : vec3<f32> = vec3<f32>(x_528.x, x_528.y, x_528.x);
   let x_951 : i32 = *(x_34);
@@ -149,7 +149,7 @@ fn performPartition_i1_i1_(l : ptr<function, i32>, h : ptr<function, i32>) -> i3
   let x_536 : vec3<f32> = vec3<f32>(x_534.x, x_534.z, x_535.x);
   j_1 = 10;
   let x_960 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_960;
   loop {
     let x_961 : i32 = pivot;
@@ -164,7 +164,7 @@ fn performPartition_i1_i1_(l : ptr<function, i32>, h : ptr<function, i32>) -> i3
     pivot = x_963;
     x_537 = vec2<f32>(vec3<f32>(1.0, 2.0, 3.0).y, vec3<f32>(1.0, 2.0, 3.0).z);
     let x_964 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_964;
     let x_56 : i32 = *(h);
     let x_965 : i32 = *(h);
@@ -198,7 +198,7 @@ fn performPartition_i1_i1_(l : ptr<function, i32>, h : ptr<function, i32>) -> i3
     param_1 = x_971;
     let x_62 : i32 = *(x_61);
     let x_972 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_972;
     let x_63 : i32 = pivot;
     let x_540 : vec2<f32> = vec2<f32>(vec3<f32>(1.0, 2.0, 3.0).y, x_534.z);
@@ -257,7 +257,7 @@ fn performPartition_i1_i1_(l : ptr<function, i32>, h : ptr<function, i32>) -> i3
       param_1 = x_985;
     }
     let x_986 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_986;
 
     continuing {
@@ -291,7 +291,7 @@ fn performPartition_i1_i1_(l : ptr<function, i32>, h : ptr<function, i32>) -> i3
   *(x_42) = x_993;
   let x_549 : vec2<f32> = vec2<f32>(x_534.x, x_534.y);
   let x_994 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_994;
   let x_995 : i32 = *(h);
   *(h) = 0;
@@ -351,15 +351,15 @@ fn quicksort_() {
   var p : i32;
   var l_1 : i32;
   var top : i32;
-  var stack : array<i32, 10>;
+  var stack : array<i32, 10u>;
   var param_5 : i32;
   l_1 = 0;
   let x_1007 : i32 = param_5;
   param_5 = 0;
   param_5 = x_1007;
   h_1 = 9;
-  let x_1008 : array<i32, 10> = stack;
-  stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+  let x_1008 : array<i32, 10u> = stack;
+  stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
   stack = x_1008;
   let x_556 : vec2<f32> = vec2<f32>(vec3<f32>(1.0, 2.0, 3.0).y, vec3<f32>(1.0, 2.0, 3.0).y);
   let x_1009 : i32 = param_5;
@@ -392,12 +392,12 @@ fn quicksort_() {
   param_4 = x_1015;
   let x_95 : i32 = l_1;
   let x_1016 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_1016;
   let x_560 : vec3<f32> = vec3<f32>(x_559.y, x_559.x, x_557.x);
   let x_96 : ptr<function, i32> = &(stack[x_94]);
-  let x_1017 : array<i32, 10> = stack;
-  stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+  let x_1017 : array<i32, 10u> = stack;
+  stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
   stack = x_1017;
   let x_561 : vec3<f32> = vec3<f32>(x_556.y, x_556.y, x_556.y);
   let x_1018 : i32 = l_1;
@@ -446,12 +446,12 @@ fn quicksort_() {
     let x_1028 : i32 = h_1;
     h_1 = 0;
     h_1 = x_1028;
-    let x_1029 : array<i32, 10> = stack;
-    stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+    let x_1029 : array<i32, 10u> = stack;
+    stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
     stack = x_1029;
     let x_106 : i32 = top;
-    let x_1030 : array<i32, 10> = stack;
-    stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+    let x_1030 : array<i32, 10u> = stack;
+    stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
     stack = x_1030;
     let x_567 : vec2<f32> = vec2<f32>(x_558.x, x_564.z);
     let x_1031 : i32 = param_4;
@@ -462,7 +462,7 @@ fn quicksort_() {
       break;
     }
     let x_1032 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_1032;
     let x_568 : vec3<f32> = vec3<f32>(x_559.y, x_559.x, x_563.y);
     let x_1033 : i32 = param_4;
@@ -486,16 +486,16 @@ fn quicksort_() {
     *(x_96) = 0;
     *(x_96) = x_1037;
     let x_111 : i32 = *(x_110);
-    let x_1038 : array<i32, 10> = stack;
-    stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+    let x_1038 : array<i32, 10u> = stack;
+    stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
     stack = x_1038;
     let x_571 : vec3<f32> = vec3<f32>(x_559.y, x_559.x, x_564.y);
     let x_1039 : i32 = l_1;
     l_1 = 0;
     l_1 = x_1039;
     h_1 = x_111;
-    let x_1040 : array<i32, 10> = stack;
-    stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+    let x_1040 : array<i32, 10u> = stack;
+    stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
     stack = x_1040;
     let x_572 : vec2<f32> = vec2<f32>(x_562.y, x_561.y);
     let x_1041 : i32 = p;
@@ -586,8 +586,8 @@ fn quicksort_() {
       let x_1061 : i32 = *(x_100);
       *(x_100) = 0;
       *(x_100) = x_1061;
-      let x_1062 : array<i32, 10> = stack;
-      stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+      let x_1062 : array<i32, 10u> = stack;
+      stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
       stack = x_1062;
       let x_584 : vec2<f32> = vec2<f32>(x_569.z, x_569.y);
       let x_585 : vec3<f32> = vec3<f32>(x_580.y, x_577.x, x_577.x);
@@ -625,8 +625,8 @@ fn quicksort_() {
       h_1 = 0;
       h_1 = x_1070;
       top = x_133;
-      let x_1071 : array<i32, 10> = stack;
-      stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+      let x_1071 : array<i32, 10u> = stack;
+      stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
       stack = x_1071;
       let x_134 : i32 = p;
       let x_590 : vec2<f32> = vec2<f32>(x_576.x, x_573.y);
@@ -651,7 +651,7 @@ fn quicksort_() {
     *(x_96) = x_1076;
     let x_592 : vec2<f32> = vec2<f32>(vec3<f32>(1.0, 2.0, 3.0).x, vec3<f32>(1.0, 2.0, 3.0).y);
     let x_1077 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_1077;
     let x_137 : i32 = p;
     let x_1078 : i32 = *(x_114);
@@ -715,16 +715,16 @@ fn quicksort_() {
       *(x_110) = x_1092;
       let x_601 : vec2<f32> = vec2<f32>(x_563.x, x_563.y);
       *(x_147) = bitcast<i32>((1u + bitcast<u32>(x_145)));
-      let x_1093 : array<i32, 10> = stack;
-      stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+      let x_1093 : array<i32, 10u> = stack;
+      stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
       stack = x_1093;
       let x_148 : i32 = top;
       let x_1094 : i32 = *(x_114);
       *(x_114) = 0;
       *(x_114) = x_1094;
       let x_602 : vec2<f32> = vec2<f32>(x_565.y, x_599.y);
-      let x_1095 : array<i32, 10> = stack;
-      stack = array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
+      let x_1095 : array<i32, 10u> = stack;
+      stack = array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0);
       stack = x_1095;
       let x_149 : i32 = (x_148 + bitcast<i32>(1u));
       let x_1096 : i32 = *(x_147);
@@ -760,7 +760,7 @@ fn quicksort_() {
       l_1 = x_1103;
       let x_604 : vec2<f32> = vec2<f32>(x_563.z, x_564.x);
       let x_1104 : QuicksortObject = obj;
-      obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+      obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
       obj = x_1104;
     }
   }
@@ -779,11 +779,11 @@ fn main_1() {
   uv = x_717;
   i_2 = 0;
   let x_721 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_721;
   if (true) {
     let x_722 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_722;
     let x_431 : vec2<f32> = vec2<f32>(vec3<f32>(1.0, 2.0, 3.0).x, vec3<f32>(1.0, 2.0, 3.0).x);
     let x_158 : i32 = i_2;
@@ -795,11 +795,11 @@ fn main_1() {
     color = x_725;
     let x_432 : vec2<f32> = vec2<f32>(x_431.y, x_431.y);
     let x_726 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_726;
   }
   let x_756 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_756;
   let x_446 : vec2<f32> = vec2<f32>(vec2<f32>(0.0, 0.0).x, vec2<f32>(0.0, 0.0).x);
   let x_757 : i32 = i_2;
@@ -807,7 +807,7 @@ fn main_1() {
   i_2 = x_757;
   quicksort_();
   let x_758 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_758;
   let x_184 : vec4<f32> = gl_FragCoord;
   let x_759 : vec2<f32> = uv;
@@ -820,14 +820,14 @@ fn main_1() {
   let x_185 : vec2<f32> = vec2<f32>(x_184.x, x_184.y);
   let x_448 : vec3<f32> = vec3<f32>(x_185.y, x_446.y, x_446.y);
   let x_761 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_761;
   let x_762 : vec2<f32> = uv;
   uv = vec2<f32>(0.0, 0.0);
   uv = x_762;
   let x_191 : vec2<f32> = x_188.resolution;
   let x_763 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_763;
   let x_449 : vec3<f32> = vec3<f32>(x_184.y, vec3<f32>(1.0, 2.0, 3.0).z, x_184.w);
   let x_764 : vec3<f32> = color;
@@ -835,7 +835,7 @@ fn main_1() {
   color = x_764;
   let x_192 : vec2<f32> = (x_185 / x_191);
   let x_765 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_765;
   let x_450 : vec2<f32> = vec2<f32>(x_447.x, x_185.y);
   let x_766 : vec3<f32> = color;
@@ -851,7 +851,7 @@ fn main_1() {
   color = x_768;
   let x_451 : vec3<f32> = vec3<f32>(x_185.x, x_185.y, x_446.y);
   let x_769 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_769;
   let x_200 : ptr<private, i32> = &(obj.numbers[0u]);
   let x_770 : i32 = *(x_200);
@@ -859,7 +859,7 @@ fn main_1() {
   *(x_200) = x_770;
   let x_201 : i32 = *(x_200);
   let x_771 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_771;
   let x_205 : ptr<function, f32> = &(color.x);
   let x_772 : i32 = *(x_200);
@@ -874,7 +874,7 @@ fn main_1() {
   i_2 = 0;
   i_2 = x_774;
   let x_775 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_775;
   let x_208 : ptr<function, f32> = &(color.x);
   let x_453 : vec3<f32> = vec3<f32>(x_451.x, x_450.x, x_450.y);
@@ -893,7 +893,7 @@ fn main_1() {
   *(x_209) = 0.0;
   *(x_209) = x_778;
   let x_779 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_779;
   if ((x_210 > 0.25)) {
     let x_780 : i32 = i_2;
@@ -908,7 +908,7 @@ fn main_1() {
     *(x_209) = x_782;
     let x_216 : i32 = obj.numbers[1];
     let x_783 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_783;
     let x_457 : vec2<f32> = vec2<f32>(x_454.x, x_454.x);
     let x_784 : vec2<f32> = uv;
@@ -916,7 +916,7 @@ fn main_1() {
     uv = x_784;
     let x_218 : ptr<function, f32> = &(color[0]);
     let x_785 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_785;
     let x_458 : vec2<f32> = vec2<f32>(vec3<f32>(1.0, 2.0, 3.0).z, vec2<f32>(0.0, 0.0).y);
     let x_786 : i32 = i_2;
@@ -1038,7 +1038,7 @@ fn main_1() {
     *(x_208) = 0.0;
     *(x_208) = x_816;
     let x_817 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_817;
     let x_468 : vec3<f32> = vec3<f32>(x_467.x, x_467.x, x_467.x);
     let x_818 : f32 = *(x_237);
@@ -1152,7 +1152,7 @@ fn main_1() {
     *(x_237) = x_844;
     let x_482 : vec3<f32> = vec3<f32>(x_455.x, x_475.y, x_455.y);
     let x_845 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_845;
     let x_846 : f32 = *(x_260);
     *(x_260) = 0.0;
@@ -1225,7 +1225,7 @@ fn main_1() {
     *(x_279) = x_863;
     let x_490 : vec2<f32> = vec2<f32>(x_480.z, x_480.z);
     let x_864 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_864;
     color.y = (f32(x_280) + x_283);
     let x_865 : f32 = *(x_208);
@@ -1243,7 +1243,7 @@ fn main_1() {
   let x_286 : ptr<function, f32> = &(uv.y);
   let x_287 : f32 = *(x_286);
   let x_868 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_868;
   let x_493 : vec2<f32> = vec2<f32>(x_475.x, x_475.y);
   let x_869 : f32 = *(x_237);
@@ -1408,7 +1408,7 @@ fn main_1() {
     *(x_209) = 0.0;
     *(x_209) = x_910;
     let x_911 : QuicksortObject = obj;
-    obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+    obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
     obj = x_911;
     let x_513 : vec3<f32> = vec3<f32>(x_505.z, x_505.x, x_448.x);
     let x_912 : i32 = *(x_300);
@@ -1461,10 +1461,10 @@ fn main_1() {
   *(x_209) = 0.0;
   *(x_209) = x_923;
   let x_924 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_924;
   let x_925 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_925;
   let x_926 : f32 = *(x_259);
   *(x_259) = 0.0;
@@ -1483,7 +1483,7 @@ fn main_1() {
   *(x_222) = x_929;
   x_GLF_color = x_330;
   let x_930 : QuicksortObject = obj;
-  obj = QuicksortObject(array<i32, 10>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
+  obj = QuicksortObject(array<i32, 10u>(0, 0, 0, 0, 0, 0, 0, 0, 0, 0));
   obj = x_930;
   let x_522 : vec3<f32> = vec3<f32>(x_330.w, x_330.y, x_493.x);
   let x_931 : f32 = *(x_208);
diff --git a/test/bug/tint/870.spvasm.expected.wgsl b/test/bug/tint/870.spvasm.expected.wgsl
index 7729cc11d6..ea206337b0 100644
--- a/test/bug/tint/870.spvasm.expected.wgsl
+++ b/test/bug/tint/870.spvasm.expected.wgsl
@@ -1,4 +1,4 @@
-type Arr = [[stride(4)]] array<i32, 6>;
+type Arr = [[stride(4)]] array<i32, 6u>;
 
 struct sspp962805860buildInformationS {
   footprint : vec4<f32>;
@@ -15,7 +15,7 @@ struct x_B4_BuildInformation {
 [[group(0), binding(2)]] var<storage, read> sspp962805860buildInformation : x_B4_BuildInformation;
 
 fn main_1() {
-  var orientation : array<i32, 6>;
+  var orientation : array<i32, 6u>;
   let x_23 : Arr = sspp962805860buildInformation.passthru.orientation;
   orientation[0] = x_23[0u];
   orientation[1] = x_23[1u];
diff --git a/test/bug/tint/943.spvasm.expected.wgsl b/test/bug/tint/943.spvasm.expected.wgsl
index 15bdc14adc..a30c20f62c 100644
--- a/test/bug/tint/943.spvasm.expected.wgsl
+++ b/test/bug/tint/943.spvasm.expected.wgsl
@@ -50,9 +50,9 @@ var<private> gl_LocalInvocationID : vec3<u32>;
 
 var<private> gl_GlobalInvocationID : vec3<u32>;
 
-var<workgroup> mm_Asub : array<array<f32, 64>, 64>;
+var<workgroup> mm_Asub : array<array<f32, 64u>, 64u>;
 
-var<workgroup> mm_Bsub : array<array<f32, 1>, 64>;
+var<workgroup> mm_Bsub : array<array<f32, 1u>, 64u>;
 
 [[group(0), binding(1)]] var<storage, read> x_165 : ssbA;
 
@@ -199,7 +199,7 @@ fn mm_matMul_i1_i1_i1_(dimAOuter : ptr<function, i32>, dimInner : ptr<function,
   var numTiles : i32;
   var innerRow : i32;
   var innerCol : i32;
-  var acc : array<array<f32, 1>, 1>;
+  var acc : array<array<f32, 1u>, 1u>;
   var tileColA : i32;
   var tileRowB : i32;
   var t : i32;
@@ -217,7 +217,7 @@ fn mm_matMul_i1_i1_i1_(dimAOuter : ptr<function, i32>, dimInner : ptr<function,
   var param_6 : i32;
   var k : i32;
   var inner : i32;
-  var BCached : array<f32, 1>;
+  var BCached : array<f32, 1u>;
   var innerRow_3 : i32;
   var ACached : f32;
   var innerCol_3 : i32;
diff --git a/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl b/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl
index 0ac837a1d7..ee03dbc823 100644
--- a/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl
+++ b/test/layout/storage/mat2x2/stride/16.spvasm.expected.wgsl
@@ -1,16 +1,16 @@
 [[block]]
 struct SSBO {
-  m : [[stride(16)]] array<vec2<f32>, 2>;
+  m : [[stride(16)]] array<vec2<f32>, 2u>;
 };
 
 [[group(0), binding(0)]] var<storage, read_write> ssbo : SSBO;
 
-fn arr_to_mat2x2_stride_16(arr : [[stride(16)]] array<vec2<f32>, 2>) -> mat2x2<f32> {
+fn arr_to_mat2x2_stride_16(arr : [[stride(16)]] array<vec2<f32>, 2u>) -> mat2x2<f32> {
   return mat2x2<f32>(arr[0u], arr[1u]);
 }
 
-fn mat2x2_stride_16_to_arr(mat : mat2x2<f32>) -> [[stride(16)]] array<vec2<f32>, 2> {
-  return [[stride(16)]] array<vec2<f32>, 2>(mat[0u], mat[1u]);
+fn mat2x2_stride_16_to_arr(mat : mat2x2<f32>) -> [[stride(16)]] array<vec2<f32>, 2u> {
+  return [[stride(16)]] array<vec2<f32>, 2u>(mat[0u], mat[1u]);
 }
 
 fn f_1() {