diff --git a/src/resolver/resolver_validation.cc b/src/resolver/resolver_validation.cc
index 1df63ba371..3218e4af49 100644
--- a/src/resolver/resolver_validation.cc
+++ b/src/resolver/resolver_validation.cc
@@ -459,42 +459,6 @@ bool Resolver::ValidateGlobalVariable(const sem::Variable* var) {
     return false;
   }
 
-  switch (var->StorageClass()) {
-    case ast::StorageClass::kStorage: {
-      // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
-      // A variable in the storage storage class is a storage buffer variable.
-      // Its store type must be a host-shareable structure type with block
-      // attribute, satisfying the storage class constraints.
-
-      auto* str = var->Type()->UnwrapRef()->As<sem::Struct>();
-      if (!str) {
-        AddError(
-            "variables declared in the <storage> storage class must be of a "
-            "structure type",
-            decl->source);
-        return false;
-      }
-      break;
-    }
-    case ast::StorageClass::kUniform: {
-      // https://gpuweb.github.io/gpuweb/wgsl/#module-scope-variables
-      // A variable in the uniform storage class is a uniform buffer variable.
-      // Its store type must be a host-shareable structure type with block
-      // attribute, satisfying the storage class constraints.
-      auto* str = var->Type()->UnwrapRef()->As<sem::Struct>();
-      if (!str) {
-        AddError(
-            "variables declared in the <uniform> storage class must be of a "
-            "structure type",
-            decl->source);
-        return false;
-      }
-      break;
-    }
-    default:
-      break;
-  }
-
   if (!decl->is_const) {
     if (!ValidateAtomicVariable(var)) {
       return false;
@@ -580,14 +544,6 @@ bool Resolver::ValidateVariable(const sem::Variable* var) {
     return false;
   }
 
-  if (auto* r = storage_ty->As<sem::Array>()) {
-    if (r->IsRuntimeSized()) {
-      AddError("runtime arrays may only appear as the last member of a struct",
-               decl->source);
-      return false;
-    }
-  }
-
   if (auto* r = storage_ty->As<sem::MultisampledTexture>()) {
     if (r->dim() != ast::TextureDimension::k2d) {
       AddError("only 2d multisampled textures are supported", decl->source);
diff --git a/src/resolver/storage_class_validation_test.cc b/src/resolver/storage_class_validation_test.cc
index f8360882aa..5d922c2f50 100644
--- a/src/resolver/storage_class_validation_test.cc
+++ b/src/resolver/storage_class_validation_test.cc
@@ -92,6 +92,40 @@ note: while analysing structure member S.m
 }
 
 TEST_F(ResolverStorageClassValidationTest, StorageBufferBool) {
+  // var<storage> g : bool;
+  Global(Source{{56, 78}}, "g", ty.bool_(), ast::StorageClass::kStorage,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(0),
+             create<ast::GroupDecoration>(0),
+         });
+
+  ASSERT_FALSE(r()->Resolve());
+
+  EXPECT_EQ(
+      r()->error(),
+      R"(56:78 error: Type 'bool' cannot be used in storage class 'storage' as it is non-host-shareable
+56:78 note: while instantiating variable g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) {
+  // var<storage> g : ptr<private, f32>;
+  Global(Source{{56, 78}}, "g",
+         ty.pointer(ty.f32(), ast::StorageClass::kPrivate),
+         ast::StorageClass::kStorage,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(0),
+             create<ast::GroupDecoration>(0),
+         });
+
+  ASSERT_FALSE(r()->Resolve());
+
+  EXPECT_EQ(
+      r()->error(),
+      R"(56:78 error: Type 'ptr<private, f32, read_write>' cannot be used in storage class 'storage' as it is non-host-shareable
+56:78 note: while instantiating variable g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, StorageBufferIntScalar) {
   // var<storage> g : i32;
   Global(Source{{56, 78}}, "g", ty.i32(), ast::StorageClass::kStorage,
          ast::DecorationList{
@@ -99,14 +133,10 @@ TEST_F(ResolverStorageClassValidationTest, StorageBufferBool) {
              create<ast::GroupDecoration>(0),
          });
 
-  ASSERT_FALSE(r()->Resolve());
-
-  EXPECT_EQ(
-      r()->error(),
-      R"(56:78 error: variables declared in the <storage> storage class must be of a structure type)");
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
-TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) {
+TEST_F(ResolverStorageClassValidationTest, StorageBufferVector) {
   // var<storage> g : vec4<f32>;
   Global(Source{{56, 78}}, "g", ty.vec4<f32>(), ast::StorageClass::kStorage,
          ast::DecorationList{
@@ -114,11 +144,7 @@ TEST_F(ResolverStorageClassValidationTest, StorageBufferPointer) {
              create<ast::GroupDecoration>(0),
          });
 
-  ASSERT_FALSE(r()->Resolve());
-
-  EXPECT_EQ(
-      r()->error(),
-      R"(56:78 error: variables declared in the <storage> storage class must be of a structure type)");
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverStorageClassValidationTest, StorageBufferArray) {
@@ -132,11 +158,7 @@ TEST_F(ResolverStorageClassValidationTest, StorageBufferArray) {
              create<ast::GroupDecoration>(0),
          });
 
-  ASSERT_FALSE(r()->Resolve());
-
-  EXPECT_EQ(
-      r()->error(),
-      R"(56:78 error: variables declared in the <storage> storage class must be of a structure type)");
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverStorageClassValidationTest, StorageBufferBoolAlias) {
@@ -240,8 +262,10 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferBool) {
 }
 
 TEST_F(ResolverStorageClassValidationTest, UniformBufferPointer) {
-  // var<uniform> g : vec4<f32>;
-  Global(Source{{56, 78}}, "g", ty.vec4<f32>(), ast::StorageClass::kUniform,
+  // var<uniform> g : ptr<private, f32>;
+  Global(Source{{56, 78}}, "g",
+         ty.pointer(ty.f32(), ast::StorageClass::kPrivate),
+         ast::StorageClass::kUniform,
          ast::DecorationList{
              create<ast::BindingDecoration>(0),
              create<ast::GroupDecoration>(0),
@@ -251,7 +275,30 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferPointer) {
 
   EXPECT_EQ(
       r()->error(),
-      R"(56:78 error: variables declared in the <uniform> storage class must be of a structure type)");
+      R"(56:78 error: Type 'ptr<private, f32, read_write>' cannot be used in storage class 'uniform' as it is non-host-shareable
+56:78 note: while instantiating variable g)");
+}
+
+TEST_F(ResolverStorageClassValidationTest, UniformBufferIntScalar) {
+  // var<uniform> g : i32;
+  Global(Source{{56, 78}}, "g", ty.i32(), ast::StorageClass::kUniform,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(0),
+             create<ast::GroupDecoration>(0),
+         });
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
+}
+
+TEST_F(ResolverStorageClassValidationTest, UniformBufferVector) {
+  // var<uniform> g : vec4<f32>;
+  Global(Source{{56, 78}}, "g", ty.vec4<f32>(), ast::StorageClass::kUniform,
+         ast::DecorationList{
+             create<ast::BindingDecoration>(0),
+             create<ast::GroupDecoration>(0),
+         });
+
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverStorageClassValidationTest, UniformBufferArray) {
@@ -264,11 +311,7 @@ TEST_F(ResolverStorageClassValidationTest, UniformBufferArray) {
              create<ast::GroupDecoration>(0),
          });
 
-  ASSERT_FALSE(r()->Resolve());
-
-  EXPECT_EQ(
-      r()->error(),
-      R"(56:78 error: variables declared in the <uniform> storage class must be of a structure type)");
+  ASSERT_TRUE(r()->Resolve()) << r()->error();
 }
 
 TEST_F(ResolverStorageClassValidationTest, UniformBufferBoolAlias) {
diff --git a/src/transform/add_spirv_block_decoration.cc b/src/transform/add_spirv_block_decoration.cc
index e4829c924b..3995c839e1 100644
--- a/src/transform/add_spirv_block_decoration.cc
+++ b/src/transform/add_spirv_block_decoration.cc
@@ -52,9 +52,9 @@ void AddSpirvBlockDecoration::Run(CloneContext& ctx, const DataMap&, DataMap&) {
     }
   }
 
-  // A map from a struct in the source program to a block-decorated wrapper that
+  // A map from a type in the source program to a block-decorated wrapper that
   // contains it in the destination program.
-  std::unordered_map<const sem::Struct*, const ast::Struct*> wrapper_structs;
+  std::unordered_map<const sem::Type*, const ast::Struct*> wrapper_structs;
 
   // Process global variables that are buffers.
   for (auto* var : ctx.src->AST().GlobalVariables()) {
@@ -64,40 +64,33 @@ void AddSpirvBlockDecoration::Run(CloneContext& ctx, const DataMap&, DataMap&) {
       continue;
     }
 
-    auto* str = sem.Get<sem::Struct>(var->type);
-    if (!str) {
-      // TODO(jrprice): We'll need to wrap these too, when WGSL supports this.
-      TINT_ICE(Transform, ctx.dst->Diagnostics())
-          << "non-struct buffer types are not yet supported";
-      continue;
-    }
+    auto* ty = sem.Get(var->type);
+    auto* str = ty->As<sem::Struct>();
+    if (!str || nested_structs.count(str)) {
+      const char* kMemberName = "inner";
 
-    if (nested_structs.count(str)) {
-      const char* kInnerStructMemberName = "inner";
-
-      // This struct is nested somewhere else, so we need to wrap it first.
-      auto* wrapper = utils::GetOrCreate(wrapper_structs, str, [&]() {
+      // This is a non-struct or a struct that is nested somewhere else, so we
+      // need to wrap it first.
+      auto* wrapper = utils::GetOrCreate(wrapper_structs, ty, [&]() {
         auto* block =
             ctx.dst->ASTNodes().Create<SpirvBlockDecoration>(ctx.dst->ID());
-        auto wrapper_name =
-            ctx.src->Symbols().NameFor(str->Declaration()->name) + "_block";
+        auto wrapper_name = ctx.src->Symbols().NameFor(var->symbol) + "_block";
         auto* ret = ctx.dst->create<ast::Struct>(
             ctx.dst->Symbols().New(wrapper_name),
-            ast::StructMemberList{ctx.dst->Member(kInnerStructMemberName,
-                                                  CreateASTTypeFor(ctx, str))},
+            ast::StructMemberList{
+                ctx.dst->Member(kMemberName, CreateASTTypeFor(ctx, ty))},
             ast::DecorationList{block});
-        ctx.InsertAfter(ctx.src->AST().GlobalDeclarations(), str->Declaration(),
-                        ret);
+        ctx.InsertBefore(ctx.src->AST().GlobalDeclarations(), var, ret);
         return ret;
       });
       ctx.Replace(var->type, ctx.dst->ty.Of(wrapper));
 
-      // Insert a member accessor to get the original struct from the wrapper at
+      // Insert a member accessor to get the original type from the wrapper at
       // any usage of the original variable.
       for (auto* user : sem_var->Users()) {
-        ctx.Replace(user->Declaration(),
-                    ctx.dst->MemberAccessor(ctx.Clone(var->symbol),
-                                            kInnerStructMemberName));
+        ctx.Replace(
+            user->Declaration(),
+            ctx.dst->MemberAccessor(ctx.Clone(var->symbol), kMemberName));
       }
     } else {
       // Add a block decoration to this struct directly.
diff --git a/src/transform/add_spirv_block_decoration_test.cc b/src/transform/add_spirv_block_decoration_test.cc
index 3780c209b4..74e4b33386 100644
--- a/src/transform/add_spirv_block_decoration_test.cc
+++ b/src/transform/add_spirv_block_decoration_test.cc
@@ -73,7 +73,98 @@ fn main() -> S {
   EXPECT_EQ(expect, str(got));
 }
 
-TEST_F(AddSpirvBlockDecorationTest, Basic) {
+TEST_F(AddSpirvBlockDecorationTest, BasicScalar) {
+  auto* src = R"(
+[[group(0), binding(0)]]
+var<uniform> u : f32;
+
+[[stage(fragment)]]
+fn main() {
+  let f = u;
+}
+)";
+  auto* expect = R"(
+[[internal(spirv_block)]]
+struct u_block {
+  inner : f32;
+};
+
+[[group(0), binding(0)]] var<uniform> u : u_block;
+
+[[stage(fragment)]]
+fn main() {
+  let f = u.inner;
+}
+)";
+
+  auto got = Run<AddSpirvBlockDecoration>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddSpirvBlockDecorationTest, BasicArray) {
+  auto* src = R"(
+[[group(0), binding(0)]]
+var<uniform> u : array<vec4<f32>, 4u>;
+
+[[stage(fragment)]]
+fn main() {
+  let a = u;
+}
+)";
+  auto* expect = R"(
+[[internal(spirv_block)]]
+struct u_block {
+  inner : array<vec4<f32>, 4u>;
+};
+
+[[group(0), binding(0)]] var<uniform> u : u_block;
+
+[[stage(fragment)]]
+fn main() {
+  let a = u.inner;
+}
+)";
+
+  auto got = Run<AddSpirvBlockDecoration>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddSpirvBlockDecorationTest, BasicArray_Alias) {
+  auto* src = R"(
+type Numbers = array<vec4<f32>, 4u>;
+
+[[group(0), binding(0)]]
+var<uniform> u : Numbers;
+
+[[stage(fragment)]]
+fn main() {
+  let a = u;
+}
+)";
+  auto* expect = R"(
+type Numbers = array<vec4<f32>, 4u>;
+
+[[internal(spirv_block)]]
+struct u_block {
+  inner : array<vec4<f32>, 4u>;
+};
+
+[[group(0), binding(0)]] var<uniform> u : u_block;
+
+[[stage(fragment)]]
+fn main() {
+  let a = u.inner;
+}
+)";
+
+  auto got = Run<AddSpirvBlockDecoration>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(AddSpirvBlockDecorationTest, BasicStruct) {
   auto* src = R"(
 struct S {
   f : f32;
@@ -174,11 +265,6 @@ struct Inner {
   f : f32;
 };
 
-[[internal(spirv_block)]]
-struct Inner_block {
-  inner : Inner;
-};
-
 [[internal(spirv_block)]]
 struct Outer {
   i : Inner;
@@ -186,7 +272,12 @@ struct Outer {
 
 [[group(0), binding(0)]] var<uniform> u0 : Outer;
 
-[[group(0), binding(1)]] var<uniform> u1 : Inner_block;
+[[internal(spirv_block)]]
+struct u1_block {
+  inner : Inner;
+};
+
+[[group(0), binding(1)]] var<uniform> u1 : u1_block;
 
 [[stage(fragment)]]
 fn main() {
@@ -226,18 +317,18 @@ struct Inner {
   f : f32;
 };
 
-[[internal(spirv_block)]]
-struct Inner_block {
-  inner : Inner;
-};
-
 struct Outer {
   i : Inner;
 };
 
 var<private> p : Outer;
 
-[[group(0), binding(1)]] var<uniform> u : Inner_block;
+[[internal(spirv_block)]]
+struct u_block {
+  inner : Inner;
+};
+
+[[group(0), binding(1)]] var<uniform> u : u_block;
 
 [[stage(fragment)]]
 fn main() {
@@ -282,11 +373,6 @@ struct Inner {
   f : f32;
 };
 
-[[internal(spirv_block)]]
-struct Inner_block {
-  inner : Inner;
-};
-
 [[internal(spirv_block)]]
 struct S {
   i : Inner;
@@ -294,9 +380,14 @@ struct S {
 
 [[group(0), binding(0)]] var<uniform> u0 : S;
 
-[[group(0), binding(1)]] var<uniform> u1 : Inner_block;
+[[internal(spirv_block)]]
+struct u1_block {
+  inner : Inner;
+};
 
-[[group(0), binding(2)]] var<uniform> u2 : Inner_block;
+[[group(0), binding(1)]] var<uniform> u1 : u1_block;
+
+[[group(0), binding(2)]] var<uniform> u2 : u1_block;
 
 [[stage(fragment)]]
 fn main() {
@@ -332,11 +423,11 @@ struct S {
 };
 
 [[internal(spirv_block)]]
-struct S_block {
+struct u_block {
   inner : S;
 };
 
-[[group(0), binding(0)]] var<uniform> u : S_block;
+[[group(0), binding(0)]] var<uniform> u : u_block;
 
 [[stage(fragment)]]
 fn main() {
@@ -375,13 +466,13 @@ struct S {
 };
 
 [[internal(spirv_block)]]
-struct S_block {
+struct u0_block {
   inner : S;
 };
 
-[[group(0), binding(0)]] var<uniform> u0 : S_block;
+[[group(0), binding(0)]] var<uniform> u0 : u0_block;
 
-[[group(0), binding(1)]] var<uniform> u1 : S_block;
+[[group(0), binding(1)]] var<uniform> u1 : u0_block;
 
 [[stage(fragment)]]
 fn main() {
@@ -427,11 +518,6 @@ struct Inner {
   f : f32;
 };
 
-[[internal(spirv_block)]]
-struct Inner_block {
-  inner : Inner;
-};
-
 type MyInner = Inner;
 
 [[internal(spirv_block)]]
@@ -443,7 +529,12 @@ type MyOuter = Outer;
 
 [[group(0), binding(0)]] var<uniform> u0 : MyOuter;
 
-[[group(0), binding(1)]] var<uniform> u1 : Inner_block;
+[[internal(spirv_block)]]
+struct u1_block {
+  inner : Inner;
+};
+
+[[group(0), binding(1)]] var<uniform> u1 : u1_block;
 
 [[stage(fragment)]]
 fn main() {
diff --git a/src/transform/module_scope_var_to_entry_point_param.cc b/src/transform/module_scope_var_to_entry_point_param.cc
index ba61160940..960fc84601 100644
--- a/src/transform/module_scope_var_to_entry_point_param.cc
+++ b/src/transform/module_scope_var_to_entry_point_param.cc
@@ -157,6 +157,7 @@ struct ModuleScopeVarToEntryPointParam::State {
 
       for (auto* var : func_sem->TransitivelyReferencedGlobals()) {
         auto sc = var->StorageClass();
+        auto* ty = var->Type()->UnwrapRef();
         if (sc == ast::StorageClass::kNone) {
           continue;
         }
@@ -174,13 +175,15 @@ struct ModuleScopeVarToEntryPointParam::State {
         auto new_var_symbol = ctx.dst->Sym();
 
         // Helper to create an AST node for the store type of the variable.
-        auto store_type = [&]() {
-          return CreateASTTypeFor(ctx, var->Type()->UnwrapRef());
-        };
+        auto store_type = [&]() { return CreateASTTypeFor(ctx, ty); };
 
         // Track whether the new variable is a pointer or not.
         bool is_pointer = false;
 
+        // Track whether the new variable was wrapped in a struct or not.
+        bool is_wrapped = false;
+        const char* kWrappedArrayMemberName = "arr";
+
         if (is_entry_point) {
           if (var->Type()->UnwrapRef()->is_handle()) {
             // For a texture or sampler variable, redeclare it as an entry point
@@ -200,8 +203,23 @@ struct ModuleScopeVarToEntryPointParam::State {
                 ast::DisabledValidation::kEntryPointParameter));
             attributes.push_back(
                 ctx.dst->Disable(ast::DisabledValidation::kIgnoreStorageClass));
-            auto* param_type = ctx.dst->ty.pointer(
-                store_type(), sc, var->Declaration()->declared_access);
+
+            auto* param_type = store_type();
+            if (auto* arr = ty->As<sem::Array>();
+                arr && arr->IsRuntimeSized()) {
+              // Wrap runtime-sized arrays in structures, so that we can declare
+              // pointers to them. Ideally we'd just emit the array itself as a
+              // pointer, but this is not representable in Tint's AST.
+              CloneStructTypes(ty);
+              auto* wrapper = ctx.dst->Structure(
+                  ctx.dst->Sym(),
+                  {ctx.dst->Member(kWrappedArrayMemberName, param_type)});
+              param_type = ctx.dst->ty.Of(wrapper);
+              is_wrapped = true;
+            }
+
+            param_type = ctx.dst->ty.pointer(
+                param_type, sc, var->Declaration()->declared_access);
             auto* param =
                 ctx.dst->Param(new_var_symbol, param_type, attributes);
             ctx.InsertFront(func_ast->params, param);
@@ -283,6 +301,10 @@ struct ModuleScopeVarToEntryPointParam::State {
 
               expr = ctx.dst->Deref(expr);
             }
+            if (is_wrapped) {
+              // Get the member from the wrapper structure.
+              expr = ctx.dst->MemberAccessor(expr, kWrappedArrayMemberName);
+            }
             ctx.Replace(user->Declaration(), expr);
           }
         }
diff --git a/src/transform/module_scope_var_to_entry_point_param_test.cc b/src/transform/module_scope_var_to_entry_point_param_test.cc
index a1620e787b..0f7d518a37 100644
--- a/src/transform/module_scope_var_to_entry_point_param_test.cc
+++ b/src/transform/module_scope_var_to_entry_point_param_test.cc
@@ -232,6 +232,99 @@ fn main([[group(0), binding(0), internal(disable_validation__entry_point_paramet
   EXPECT_EQ(expect, str(got));
 }
 
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray) {
+  auto* src = R"(
+[[group(0), binding(0)]]
+var<storage> buffer : array<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  _ = buffer[0];
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol_1 {
+  arr : array<f32>;
+};
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter), internal(disable_validation__ignore_storage_class)]] tint_symbol : ptr<storage, tint_symbol_1>) {
+  _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_RuntimeArray_Alias) {
+  auto* src = R"(
+type myarray = array<f32>;
+
+[[group(0), binding(0)]]
+var<storage> buffer : myarray;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  _ = buffer[0];
+}
+)";
+
+  auto* expect = R"(
+struct tint_symbol_1 {
+  arr : array<f32>;
+};
+
+type myarray = array<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter), internal(disable_validation__ignore_storage_class)]] tint_symbol : ptr<storage, tint_symbol_1>) {
+  _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
+TEST_F(ModuleScopeVarToEntryPointParamTest, Buffer_ArrayOfStruct) {
+  auto* src = R"(
+struct S {
+  f : f32;
+};
+
+[[group(0), binding(0)]]
+var<storage> buffer : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  _ = buffer[0];
+}
+)";
+
+  auto* expect = R"(
+struct S {
+  f : f32;
+};
+
+struct tint_symbol_1 {
+  arr : array<S>;
+};
+
+[[stage(compute), workgroup_size(1)]]
+fn main([[group(0), binding(0), internal(disable_validation__entry_point_parameter), internal(disable_validation__ignore_storage_class)]] tint_symbol : ptr<storage, tint_symbol_1>) {
+  _ = (*(tint_symbol)).arr[0];
+}
+)";
+
+  auto got = Run<ModuleScopeVarToEntryPointParam>(src);
+
+  EXPECT_EQ(expect, str(got));
+}
+
 TEST_F(ModuleScopeVarToEntryPointParamTest, Buffers_FunctionCalls) {
   auto* src = R"(
 struct S {
diff --git a/src/writer/hlsl/generator_impl.cc b/src/writer/hlsl/generator_impl.cc
index 76dc529d61..bae3044f39 100644
--- a/src/writer/hlsl/generator_impl.cc
+++ b/src/writer/hlsl/generator_impl.cc
@@ -2869,14 +2869,6 @@ bool GeneratorImpl::EmitUniformVariable(const sem::Variable* var) {
   auto* decl = var->Declaration();
   auto binding_point = decl->BindingPoint();
   auto* type = var->Type()->UnwrapRef();
-
-  auto* str = type->As<sem::Struct>();
-  if (!str) {
-    // https://www.w3.org/TR/WGSL/#module-scope-variables
-    TINT_ICE(Writer, diagnostics_)
-        << "variables with uniform storage must be structure";
-  }
-
   auto name = builder_.Symbols().NameFor(decl->symbol);
   line() << "cbuffer cbuffer_" << name << RegisterAndSpace('b', binding_point)
          << " {";
@@ -3513,13 +3505,7 @@ bool GeneratorImpl::EmitType(std::ostream& out,
       out << "ByteAddressBuffer";
       return true;
     case ast::StorageClass::kUniform: {
-      auto* str = type->As<sem::Struct>();
-      if (!str) {
-        // https://www.w3.org/TR/WGSL/#module-scope-variables
-        TINT_ICE(Writer, diagnostics_)
-            << "variables with uniform storage must be structure";
-      }
-      auto array_length = (str->Size() + 15) / 16;
+      auto array_length = (type->Size() + 15) / 16;
       out << "uint4 " << name << "[" << array_length << "]";
       if (name_printed) {
         *name_printed = true;
diff --git a/test/buffer/storage/types/array.wgsl b/test/buffer/storage/types/array.wgsl
new file mode 100644
index 0000000000..6105a26727
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : array<f32, 4>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : array<f32, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/array.wgsl.expected.hlsl b/test/buffer/storage/types/array.wgsl.expected.hlsl
new file mode 100644
index 0000000000..4b02ccdd25
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.hlsl
@@ -0,0 +1,28 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float value[4]) {
+  float array[4] = value;
+  {
+    [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      buffer.Store((offset + (i * 4u)), asuint(array[i]));
+    }
+  }
+}
+
+typedef float tint_symbol_4_ret[4];
+tint_symbol_4_ret tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  float arr[4] = (float[4])0;
+  {
+    [loop] for(uint i_1 = 0u; (i_1 < 4u); i_1 = (i_1 + 1u)) {
+      arr[i_1] = asfloat(buffer.Load((offset + (i_1 * 4u))));
+    }
+  }
+  return arr;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}
diff --git a/test/buffer/storage/types/array.wgsl.expected.msl b/test/buffer/storage/types/array.wgsl.expected.msl
new file mode 100644
index 0000000000..a95963fecf
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  /* 0x0000 */ float arr[4];
+};
+
+kernel void tint_symbol(device tint_array_wrapper* tint_symbol_1 [[buffer(1)]], const device tint_array_wrapper* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/array.wgsl.expected.spvasm b/test/buffer/storage/types/array.wgsl.expected.spvasm
new file mode 100644
index 0000000000..65f933d554
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.spvasm
@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %_arr_float_uint_4 ArrayStride 4
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_float_uint_4 = OpTypeArray %float %uint_4
+   %in_block = OpTypeStruct %_arr_float_uint_4
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer__arr_float_uint_4 = OpTypePointer StorageBuffer %_arr_float_uint_4
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_4 %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer__arr_float_uint_4 %in %uint_0
+         %17 = OpLoad %_arr_float_uint_4 %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/array.wgsl.expected.wgsl b/test/buffer/storage/types/array.wgsl.expected.wgsl
new file mode 100644
index 0000000000..0aaacaedd2
--- /dev/null
+++ b/test/buffer/storage/types/array.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : array<f32, 4>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : array<f32, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/f32.wgsl b/test/buffer/storage/types/f32.wgsl
new file mode 100644
index 0000000000..d66cdb0cea
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : f32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/f32.wgsl.expected.hlsl b/test/buffer/storage/types/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..e6dca9ed16
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store(0u, asuint(asfloat(tint_symbol.Load(0u))));
+  return;
+}
diff --git a/test/buffer/storage/types/f32.wgsl.expected.msl b/test/buffer/storage/types/f32.wgsl.expected.msl
new file mode 100644
index 0000000000..033b5e2feb
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float* tint_symbol_1 [[buffer(1)]], const device float* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/f32.wgsl.expected.spvasm b/test/buffer/storage/types/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..dea11b4b0e
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+   %in_block = OpTypeStruct %float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_float = OpTypePointer StorageBuffer %float
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %13 = OpAccessChain %_ptr_StorageBuffer_float %out %uint_0
+         %14 = OpAccessChain %_ptr_StorageBuffer_float %in %uint_0
+         %15 = OpLoad %float %14
+               OpStore %13 %15
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/f32.wgsl.expected.wgsl b/test/buffer/storage/types/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..a385a75d5e
--- /dev/null
+++ b/test/buffer/storage/types/f32.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : f32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/i32.wgsl b/test/buffer/storage/types/i32.wgsl
new file mode 100644
index 0000000000..2ffd35fe0f
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : i32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/i32.wgsl.expected.hlsl b/test/buffer/storage/types/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..3c53558620
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store(0u, asuint(asint(tint_symbol.Load(0u))));
+  return;
+}
diff --git a/test/buffer/storage/types/i32.wgsl.expected.msl b/test/buffer/storage/types/i32.wgsl.expected.msl
new file mode 100644
index 0000000000..1de9e037ca
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device int* tint_symbol_1 [[buffer(1)]], const device int* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/i32.wgsl.expected.spvasm b/test/buffer/storage/types/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..4abab36eac
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+        %int = OpTypeInt 32 1
+   %in_block = OpTypeStruct %int
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %13 = OpAccessChain %_ptr_StorageBuffer_int %out %uint_0
+         %14 = OpAccessChain %_ptr_StorageBuffer_int %in %uint_0
+         %15 = OpLoad %int %14
+               OpStore %13 %15
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/i32.wgsl.expected.wgsl b/test/buffer/storage/types/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..3424f2e6f1
--- /dev/null
+++ b/test/buffer/storage/types/i32.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : i32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat2x2.wgsl b/test/buffer/storage/types/mat2x2.wgsl
new file mode 100644
index 0000000000..60ff548adb
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat2x2<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.hlsl b/test/buffer/storage/types/mat2x2.wgsl.expected.hlsl
new file mode 100644
index 0000000000..a1fc1b643f
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float2x2 value) {
+  buffer.Store2((offset + 0u), asuint(value[0u]));
+  buffer.Store2((offset + 8u), asuint(value[1u]));
+}
+
+float2x2 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float2x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.msl b/test/buffer/storage/types/mat2x2.wgsl.expected.msl
new file mode 100644
index 0000000000..c9ffff792a
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float2x2* tint_symbol_1 [[buffer(1)]], const device float2x2* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.spvasm b/test/buffer/storage/types/mat2x2.wgsl.expected.spvasm
new file mode 100644
index 0000000000..a372c43cfa
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 8
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+   %in_block = OpTypeStruct %mat2v2float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v2float = OpTypePointer StorageBuffer %mat2v2float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat2v2float %in %uint_0
+         %17 = OpLoad %mat2v2float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/mat2x2.wgsl.expected.wgsl b/test/buffer/storage/types/mat2x2.wgsl.expected.wgsl
new file mode 100644
index 0000000000..8d2cbae323
--- /dev/null
+++ b/test/buffer/storage/types/mat2x2.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat2x2<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat2x3.wgsl b/test/buffer/storage/types/mat2x3.wgsl
new file mode 100644
index 0000000000..151db83808
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat2x3<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.hlsl b/test/buffer/storage/types/mat2x3.wgsl.expected.hlsl
new file mode 100644
index 0000000000..baff4c5904
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float2x3 value) {
+  buffer.Store3((offset + 0u), asuint(value[0u]));
+  buffer.Store3((offset + 16u), asuint(value[1u]));
+}
+
+float2x3 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float2x3(asfloat(buffer.Load3((offset + 0u))), asfloat(buffer.Load3((offset + 16u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.msl b/test/buffer/storage/types/mat2x3.wgsl.expected.msl
new file mode 100644
index 0000000000..1b704d4a15
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float2x3* tint_symbol_1 [[buffer(1)]], const device float2x3* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.spvasm b/test/buffer/storage/types/mat2x3.wgsl.expected.spvasm
new file mode 100644
index 0000000000..31ddb997df
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 16
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+   %in_block = OpTypeStruct %mat2v3float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat2v3float = OpTypePointer StorageBuffer %mat2v3float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat2v3float %in %uint_0
+         %17 = OpLoad %mat2v3float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/mat2x3.wgsl.expected.wgsl b/test/buffer/storage/types/mat2x3.wgsl.expected.wgsl
new file mode 100644
index 0000000000..ab3a5fd966
--- /dev/null
+++ b/test/buffer/storage/types/mat2x3.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat2x3<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat3x2.wgsl b/test/buffer/storage/types/mat3x2.wgsl
new file mode 100644
index 0000000000..d77ac81d22
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat3x2<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.hlsl b/test/buffer/storage/types/mat3x2.wgsl.expected.hlsl
new file mode 100644
index 0000000000..31d606bf03
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.hlsl
@@ -0,0 +1,18 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float3x2 value) {
+  buffer.Store2((offset + 0u), asuint(value[0u]));
+  buffer.Store2((offset + 8u), asuint(value[1u]));
+  buffer.Store2((offset + 16u), asuint(value[2u]));
+}
+
+float3x2 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float3x2(asfloat(buffer.Load2((offset + 0u))), asfloat(buffer.Load2((offset + 8u))), asfloat(buffer.Load2((offset + 16u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.msl b/test/buffer/storage/types/mat3x2.wgsl.expected.msl
new file mode 100644
index 0000000000..b8765f08ab
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float3x2* tint_symbol_1 [[buffer(1)]], const device float3x2* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.spvasm b/test/buffer/storage/types/mat3x2.wgsl.expected.spvasm
new file mode 100644
index 0000000000..c352e184c9
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 8
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+   %in_block = OpTypeStruct %mat3v2float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat3v2float = OpTypePointer StorageBuffer %mat3v2float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat3v2float %in %uint_0
+         %17 = OpLoad %mat3v2float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/mat3x2.wgsl.expected.wgsl b/test/buffer/storage/types/mat3x2.wgsl.expected.wgsl
new file mode 100644
index 0000000000..4c87f4ac1e
--- /dev/null
+++ b/test/buffer/storage/types/mat3x2.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat3x2<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat4x4.wgsl b/test/buffer/storage/types/mat4x4.wgsl
new file mode 100644
index 0000000000..bf18a256f9
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : mat4x4<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.hlsl b/test/buffer/storage/types/mat4x4.wgsl.expected.hlsl
new file mode 100644
index 0000000000..e818c0922d
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.hlsl
@@ -0,0 +1,19 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, float4x4 value) {
+  buffer.Store4((offset + 0u), asuint(value[0u]));
+  buffer.Store4((offset + 16u), asuint(value[1u]));
+  buffer.Store4((offset + 32u), asuint(value[2u]));
+  buffer.Store4((offset + 48u), asuint(value[3u]));
+}
+
+float4x4 tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  return float4x4(asfloat(buffer.Load4((offset + 0u))), asfloat(buffer.Load4((offset + 16u))), asfloat(buffer.Load4((offset + 32u))), asfloat(buffer.Load4((offset + 48u))));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_4(tint_symbol, 0u));
+  return;
+}
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.msl b/test/buffer/storage/types/mat4x4.wgsl.expected.msl
new file mode 100644
index 0000000000..6b33874fa0
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float4x4* tint_symbol_1 [[buffer(1)]], const device float4x4* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.spvasm b/test/buffer/storage/types/mat4x4.wgsl.expected.spvasm
new file mode 100644
index 0000000000..512341416c
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %in_block 0 ColMajor
+               OpMemberDecorate %in_block 0 MatrixStride 16
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat4v4float = OpTypeMatrix %v4float 4
+   %in_block = OpTypeStruct %mat4v4float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_mat4v4float = OpTypePointer StorageBuffer %mat4v4float
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %15 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %out %uint_0
+         %16 = OpAccessChain %_ptr_StorageBuffer_mat4v4float %in %uint_0
+         %17 = OpLoad %mat4v4float %16
+               OpStore %15 %17
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/mat4x4.wgsl.expected.wgsl b/test/buffer/storage/types/mat4x4.wgsl.expected.wgsl
new file mode 100644
index 0000000000..7c1ebe9e7a
--- /dev/null
+++ b/test/buffer/storage/types/mat4x4.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : mat4x4<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/runtime_array.wgsl b/test/buffer/storage/types/runtime_array.wgsl
new file mode 100644
index 0000000000..f403fd377d
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl
@@ -0,0 +1,14 @@
+struct S {
+  f : f32;
+};
+
+[[group(0), binding(0)]]
+var<storage, read> in : array<S>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out[0] = in[0];
+}
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.hlsl b/test/buffer/storage/types/runtime_array.wgsl.expected.hlsl
new file mode 100644
index 0000000000..035878ad15
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.hlsl
@@ -0,0 +1,21 @@
+struct S {
+  float f;
+};
+
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
+  buffer.Store((offset + 0u), asuint(value.f));
+}
+
+S tint_symbol_4(ByteAddressBuffer buffer, uint offset) {
+  const S tint_symbol_6 = {asfloat(buffer.Load((offset + 0u)))};
+  return tint_symbol_6;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, (4u * uint(0)), tint_symbol_4(tint_symbol, (4u * uint(0))));
+  return;
+}
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.msl b/test/buffer/storage/types/runtime_array.wgsl.expected.msl
new file mode 100644
index 0000000000..ff6a283788
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  /* 0x0000 */ float f;
+};
+struct tint_symbol_2 {
+  /* 0x0000 */ S arr[1];
+};
+struct tint_symbol_4 {
+  /* 0x0000 */ S arr[1];
+};
+
+kernel void tint_symbol(device tint_symbol_2* tint_symbol_1 [[buffer(1)]], const device tint_symbol_4* tint_symbol_3 [[buffer(0)]]) {
+  (*(tint_symbol_1)).arr[0] = (*(tint_symbol_3)).arr[0];
+  return;
+}
+
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.spvasm b/test/buffer/storage/types/runtime_array.wgsl.expected.spvasm
new file mode 100644
index 0000000000..5a63e67be4
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.spvasm
@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 20
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %S "S"
+               OpMemberName %S 0 "f"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpMemberDecorate %S 0 Offset 0
+               OpDecorate %_runtimearr_S ArrayStride 4
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+          %S = OpTypeStruct %float
+%_runtimearr_S = OpTypeRuntimeArray %S
+   %in_block = OpTypeStruct %_runtimearr_S
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+        %int = OpTypeInt 32 1
+      %int_0 = OpConstant %int 0
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %17 = OpAccessChain %_ptr_StorageBuffer_S %out %uint_0 %int_0
+         %18 = OpAccessChain %_ptr_StorageBuffer_S %in %uint_0 %int_0
+         %19 = OpLoad %S %18
+               OpStore %17 %19
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/runtime_array.wgsl.expected.wgsl b/test/buffer/storage/types/runtime_array.wgsl.expected.wgsl
new file mode 100644
index 0000000000..418c67d64f
--- /dev/null
+++ b/test/buffer/storage/types/runtime_array.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+struct S {
+  f : f32;
+};
+
+[[group(0), binding(0)]] var<storage, read> in : array<S>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : array<S>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out[0] = in[0];
+}
diff --git a/test/buffer/storage/types/struct.wgsl b/test/buffer/storage/types/struct.wgsl
new file mode 100644
index 0000000000..42acfbce5e
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl
@@ -0,0 +1,17 @@
+struct Inner {
+  f : f32;
+};
+struct S {
+  inner : Inner;
+};
+
+[[group(0), binding(0)]]
+var<storage, read> in : S;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/struct.wgsl.expected.hlsl b/test/buffer/storage/types/struct.wgsl.expected.hlsl
new file mode 100644
index 0000000000..05be6006f9
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.hlsl
@@ -0,0 +1,33 @@
+struct Inner {
+  float f;
+};
+struct S {
+  Inner inner;
+};
+
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+void tint_symbol_3(RWByteAddressBuffer buffer, uint offset, Inner value) {
+  buffer.Store((offset + 0u), asuint(value.f));
+}
+
+void tint_symbol_2(RWByteAddressBuffer buffer, uint offset, S value) {
+  tint_symbol_3(buffer, (offset + 0u), value.inner);
+}
+
+Inner tint_symbol_6(ByteAddressBuffer buffer, uint offset) {
+  const Inner tint_symbol_8 = {asfloat(buffer.Load((offset + 0u)))};
+  return tint_symbol_8;
+}
+
+S tint_symbol_5(ByteAddressBuffer buffer, uint offset) {
+  const S tint_symbol_9 = {tint_symbol_6(buffer, (offset + 0u))};
+  return tint_symbol_9;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_2(tint_symbol_1, 0u, tint_symbol_5(tint_symbol, 0u));
+  return;
+}
diff --git a/test/buffer/storage/types/struct.wgsl.expected.msl b/test/buffer/storage/types/struct.wgsl.expected.msl
new file mode 100644
index 0000000000..d5c7427870
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.msl
@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ float f;
+};
+struct S {
+  /* 0x0000 */ Inner inner;
+};
+
+kernel void tint_symbol(device S* tint_symbol_1 [[buffer(0)]], const device S* tint_symbol_2 [[buffer(1)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/struct.wgsl.expected.spvasm b/test/buffer/storage/types/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000000..058a3e5427
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.spvasm
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 12
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "inner"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "f"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+      %Inner = OpTypeStruct %float
+          %S = OpTypeStruct %Inner
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+         %in = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %11 = OpLoad %S %in
+               OpStore %out %11
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/struct.wgsl.expected.wgsl b/test/buffer/storage/types/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000000..f86c26b429
--- /dev/null
+++ b/test/buffer/storage/types/struct.wgsl.expected.wgsl
@@ -0,0 +1,16 @@
+struct Inner {
+  f : f32;
+};
+
+struct S {
+  inner : Inner;
+};
+
+[[group(0), binding(0)]] var<storage, read> in : S;
+
+[[group(0), binding(1)]] var<storage, read_write> out : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/u32.wgsl b/test/buffer/storage/types/u32.wgsl
new file mode 100644
index 0000000000..4664d29d6c
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : u32;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/u32.wgsl.expected.hlsl b/test/buffer/storage/types/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..8574dafdc8
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store(0u, asuint(tint_symbol.Load(0u)));
+  return;
+}
diff --git a/test/buffer/storage/types/u32.wgsl.expected.msl b/test/buffer/storage/types/u32.wgsl.expected.msl
new file mode 100644
index 0000000000..bb65310eec
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device uint* tint_symbol_1 [[buffer(1)]], const device uint* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/u32.wgsl.expected.spvasm b/test/buffer/storage/types/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..201cc294b3
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.spvasm
@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+       %uint = OpTypeInt 32 0
+   %in_block = OpTypeStruct %uint
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_uint = OpTypePointer StorageBuffer %uint
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %12 = OpAccessChain %_ptr_StorageBuffer_uint %out %uint_0
+         %13 = OpAccessChain %_ptr_StorageBuffer_uint %in %uint_0
+         %14 = OpLoad %uint %13
+               OpStore %12 %14
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/u32.wgsl.expected.wgsl b/test/buffer/storage/types/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..46dc9d383e
--- /dev/null
+++ b/test/buffer/storage/types/u32.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : u32;
+
+[[group(0), binding(1)]] var<storage, read_write> out : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/vec2.wgsl b/test/buffer/storage/types/vec2.wgsl
new file mode 100644
index 0000000000..9841fd4181
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec2<i32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.hlsl b/test/buffer/storage/types/vec2.wgsl.expected.hlsl
new file mode 100644
index 0000000000..3c5c7d7ebe
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store2(0u, asuint(asint(tint_symbol.Load2(0u))));
+  return;
+}
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.msl b/test/buffer/storage/types/vec2.wgsl.expected.msl
new file mode 100644
index 0000000000..c730e05c48
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device int2* tint_symbol_1 [[buffer(1)]], const device int2* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.spvasm b/test/buffer/storage/types/vec2.wgsl.expected.spvasm
new file mode 100644
index 0000000000..85b79d3146
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.spvasm
@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+   %in_block = OpTypeStruct %v2int
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v2int = OpTypePointer StorageBuffer %v2int
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_StorageBuffer_v2int %out %uint_0
+         %15 = OpAccessChain %_ptr_StorageBuffer_v2int %in %uint_0
+         %16 = OpLoad %v2int %15
+               OpStore %14 %16
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/vec2.wgsl.expected.wgsl b/test/buffer/storage/types/vec2.wgsl.expected.wgsl
new file mode 100644
index 0000000000..a945b3228b
--- /dev/null
+++ b/test/buffer/storage/types/vec2.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec2<i32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/vec3.wgsl b/test/buffer/storage/types/vec3.wgsl
new file mode 100644
index 0000000000..5335059754
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec3<u32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.hlsl b/test/buffer/storage/types/vec3.wgsl.expected.hlsl
new file mode 100644
index 0000000000..9af25f4d63
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store3(0u, asuint(tint_symbol.Load3(0u)));
+  return;
+}
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.msl b/test/buffer/storage/types/vec3.wgsl.expected.msl
new file mode 100644
index 0000000000..6d55796494
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device uint3* tint_symbol_1 [[buffer(1)]], const device uint3* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.spvasm b/test/buffer/storage/types/vec3.wgsl.expected.spvasm
new file mode 100644
index 0000000000..70a2dafdb0
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.spvasm
@@ -0,0 +1,39 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+   %in_block = OpTypeStruct %v3uint
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v3uint = OpTypePointer StorageBuffer %v3uint
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %13 = OpAccessChain %_ptr_StorageBuffer_v3uint %out %uint_0
+         %14 = OpAccessChain %_ptr_StorageBuffer_v3uint %in %uint_0
+         %15 = OpLoad %v3uint %14
+               OpStore %13 %15
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/vec3.wgsl.expected.wgsl b/test/buffer/storage/types/vec3.wgsl.expected.wgsl
new file mode 100644
index 0000000000..71377b0d4d
--- /dev/null
+++ b/test/buffer/storage/types/vec3.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec3<u32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/vec4.wgsl b/test/buffer/storage/types/vec4.wgsl
new file mode 100644
index 0000000000..494a89abdc
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl
@@ -0,0 +1,10 @@
+[[group(0), binding(0)]]
+var<storage, read> in : vec4<f32>;
+
+[[group(0), binding(1)]]
+var<storage, read_write> out : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.hlsl b/test/buffer/storage/types/vec4.wgsl.expected.hlsl
new file mode 100644
index 0000000000..dd068ba25a
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.hlsl
@@ -0,0 +1,8 @@
+ByteAddressBuffer tint_symbol : register(t0, space0);
+RWByteAddressBuffer tint_symbol_1 : register(u1, space0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_symbol_1.Store4(0u, asuint(asfloat(tint_symbol.Load4(0u))));
+  return;
+}
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.msl b/test/buffer/storage/types/vec4.wgsl.expected.msl
new file mode 100644
index 0000000000..4af349b7cf
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(device float4* tint_symbol_1 [[buffer(1)]], const device float4* tint_symbol_2 [[buffer(0)]]) {
+  *(tint_symbol_1) = *(tint_symbol_2);
+  return;
+}
+
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.spvasm b/test/buffer/storage/types/vec4.wgsl.expected.spvasm
new file mode 100644
index 0000000000..d74f2981da
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.spvasm
@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %in_block "in_block"
+               OpMemberName %in_block 0 "inner"
+               OpName %in "in"
+               OpName %out "out"
+               OpName %main "main"
+               OpDecorate %in_block Block
+               OpMemberDecorate %in_block 0 Offset 0
+               OpDecorate %in NonWritable
+               OpDecorate %in DescriptorSet 0
+               OpDecorate %in Binding 0
+               OpDecorate %out DescriptorSet 0
+               OpDecorate %out Binding 1
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+   %in_block = OpTypeStruct %v4float
+%_ptr_StorageBuffer_in_block = OpTypePointer StorageBuffer %in_block
+         %in = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+        %out = OpVariable %_ptr_StorageBuffer_in_block StorageBuffer
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_v4float = OpTypePointer StorageBuffer %v4float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_StorageBuffer_v4float %out %uint_0
+         %15 = OpAccessChain %_ptr_StorageBuffer_v4float %in %uint_0
+         %16 = OpLoad %v4float %15
+               OpStore %14 %16
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/storage/types/vec4.wgsl.expected.wgsl b/test/buffer/storage/types/vec4.wgsl.expected.wgsl
new file mode 100644
index 0000000000..7b278bd0fc
--- /dev/null
+++ b/test/buffer/storage/types/vec4.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+[[group(0), binding(0)]] var<storage, read> in : vec4<f32>;
+
+[[group(0), binding(1)]] var<storage, read_write> out : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  out = in;
+}
diff --git a/test/buffer/uniform/types/array.wgsl b/test/buffer/uniform/types/array.wgsl
new file mode 100644
index 0000000000..7151b53956
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : array<vec4<f32>, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/array.wgsl.expected.hlsl b/test/buffer/uniform/types/array.wgsl.expected.hlsl
new file mode 100644
index 0000000000..34e3ef1c6a
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.hlsl
@@ -0,0 +1,21 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[4];
+};
+
+typedef float4 tint_symbol_ret[4];
+tint_symbol_ret tint_symbol(uint4 buffer[4], uint offset) {
+  float4 arr[4] = (float4[4])0;
+  {
+    [loop] for(uint i = 0u; (i < 4u); i = (i + 1u)) {
+      const uint scalar_offset = ((offset + (i * 16u))) / 4;
+      arr[i] = asfloat(buffer[scalar_offset / 4]);
+    }
+  }
+  return arr;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float4 x[4] = tint_symbol(u, 0u);
+  return;
+}
diff --git a/test/buffer/uniform/types/array.wgsl.expected.msl b/test/buffer/uniform/types/array.wgsl.expected.msl
new file mode 100644
index 0000000000..71c0b99b0f
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.msl
@@ -0,0 +1,12 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  /* 0x0000 */ float4 arr[4];
+};
+
+kernel void tint_symbol(const constant tint_array_wrapper* tint_symbol_1 [[buffer(0)]]) {
+  tint_array_wrapper const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/array.wgsl.expected.spvasm b/test/buffer/uniform/types/array.wgsl.expected.spvasm
new file mode 100644
index 0000000000..ef26f1dffb
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 17
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %_arr_v4float_uint_4 ArrayStride 16
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+       %uint = OpTypeInt 32 0
+     %uint_4 = OpConstant %uint 4
+%_arr_v4float_uint_4 = OpTypeArray %v4float %uint_4
+    %u_block = OpTypeStruct %_arr_v4float_uint_4
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform__arr_v4float_uint_4 = OpTypePointer Uniform %_arr_v4float_uint_4
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+         %15 = OpAccessChain %_ptr_Uniform__arr_v4float_uint_4 %u %uint_0
+         %16 = OpLoad %_arr_v4float_uint_4 %15
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/array.wgsl.expected.wgsl b/test/buffer/uniform/types/array.wgsl.expected.wgsl
new file mode 100644
index 0000000000..3cfa21e32d
--- /dev/null
+++ b/test/buffer/uniform/types/array.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : array<vec4<f32>, 4>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/f32.wgsl b/test/buffer/uniform/types/f32.wgsl
new file mode 100644
index 0000000000..1f3fbfa74c
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.hlsl b/test/buffer/uniform/types/f32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..496e5c9fd4
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float x = asfloat(u[0].x);
+  return;
+}
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.msl b/test/buffer/uniform/types/f32.wgsl.expected.msl
new file mode 100644
index 0000000000..81695e10d4
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float* tint_symbol_1 [[buffer(0)]]) {
+  float const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.spvasm b/test/buffer/uniform/types/f32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..bb09d33562
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %u_block = OpTypeStruct %float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_float = OpTypePointer Uniform %float
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %12 = OpAccessChain %_ptr_Uniform_float %u %uint_0
+         %13 = OpLoad %float %12
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/f32.wgsl.expected.wgsl b/test/buffer/uniform/types/f32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..6bd2751aaf
--- /dev/null
+++ b/test/buffer/uniform/types/f32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : f32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/i32.wgsl b/test/buffer/uniform/types/i32.wgsl
new file mode 100644
index 0000000000..ca165b4601
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.hlsl b/test/buffer/uniform/types/i32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..e6ed76ea43
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int x = asint(u[0].x);
+  return;
+}
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.msl b/test/buffer/uniform/types/i32.wgsl.expected.msl
new file mode 100644
index 0000000000..3d4cdde979
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant int* tint_symbol_1 [[buffer(0)]]) {
+  int const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.spvasm b/test/buffer/uniform/types/i32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..1f7b1f7701
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+        %int = OpTypeInt 32 1
+    %u_block = OpTypeStruct %int
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %12 = OpAccessChain %_ptr_Uniform_int %u %uint_0
+         %13 = OpLoad %int %12
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/i32.wgsl.expected.wgsl b/test/buffer/uniform/types/i32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..0cf7b04a61
--- /dev/null
+++ b/test/buffer/uniform/types/i32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x2.wgsl b/test/buffer/uniform/types/mat2x2.wgsl
new file mode 100644
index 0000000000..d1cc3d3d93
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl b/test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl
new file mode 100644
index 0000000000..c34b4c53df
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+float2x2 tint_symbol(uint4 buffer[1], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  uint4 ubo_load = buffer[scalar_offset / 4];
+  const uint scalar_offset_1 = ((offset + 8u)) / 4;
+  uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
+  return float2x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float2x2 x = tint_symbol(u, 0u);
+  return;
+}
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.msl b/test/buffer/uniform/types/mat2x2.wgsl.expected.msl
new file mode 100644
index 0000000000..226e54180d
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float2x2* tint_symbol_1 [[buffer(0)]]) {
+  float2x2 const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm b/test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm
new file mode 100644
index 0000000000..b5fb3179b8
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 8
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat2v2float = OpTypeMatrix %v2float 2
+    %u_block = OpTypeStruct %mat2v2float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat2v2float = OpTypePointer Uniform %mat2v2float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat2v2float %u %uint_0
+         %15 = OpLoad %mat2v2float %14
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl b/test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl
new file mode 100644
index 0000000000..8c42684b4f
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x2.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat2x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x3.wgsl b/test/buffer/uniform/types/mat2x3.wgsl
new file mode 100644
index 0000000000..0fb978d61f
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl b/test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl
new file mode 100644
index 0000000000..e92e916135
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.hlsl
@@ -0,0 +1,15 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[2];
+};
+
+float2x3 tint_symbol(uint4 buffer[2], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  return float2x3(asfloat(buffer[scalar_offset / 4].xyz), asfloat(buffer[scalar_offset_1 / 4].xyz));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float2x3 x = tint_symbol(u, 0u);
+  return;
+}
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.msl b/test/buffer/uniform/types/mat2x3.wgsl.expected.msl
new file mode 100644
index 0000000000..1c226a8283
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float2x3* tint_symbol_1 [[buffer(0)]]) {
+  float2x3 const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm b/test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm
new file mode 100644
index 0000000000..3a86048f06
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 16
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+    %u_block = OpTypeStruct %mat2v3float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat2v3float = OpTypePointer Uniform %mat2v3float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat2v3float %u %uint_0
+         %15 = OpLoad %mat2v3float %14
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl b/test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl
new file mode 100644
index 0000000000..5a74e9478a
--- /dev/null
+++ b/test/buffer/uniform/types/mat2x3.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat3x2.wgsl b/test/buffer/uniform/types/mat3x2.wgsl
new file mode 100644
index 0000000000..59b1b8059c
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl b/test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl
new file mode 100644
index 0000000000..44f2aa4431
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.hlsl
@@ -0,0 +1,19 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[2];
+};
+
+float3x2 tint_symbol(uint4 buffer[2], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  uint4 ubo_load = buffer[scalar_offset / 4];
+  const uint scalar_offset_1 = ((offset + 8u)) / 4;
+  uint4 ubo_load_1 = buffer[scalar_offset_1 / 4];
+  const uint scalar_offset_2 = ((offset + 16u)) / 4;
+  uint4 ubo_load_2 = buffer[scalar_offset_2 / 4];
+  return float3x2(asfloat(((scalar_offset & 2) ? ubo_load.zw : ubo_load.xy)), asfloat(((scalar_offset_1 & 2) ? ubo_load_1.zw : ubo_load_1.xy)), asfloat(((scalar_offset_2 & 2) ? ubo_load_2.zw : ubo_load_2.xy)));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float3x2 x = tint_symbol(u, 0u);
+  return;
+}
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.msl b/test/buffer/uniform/types/mat3x2.wgsl.expected.msl
new file mode 100644
index 0000000000..b39621d687
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float3x2* tint_symbol_1 [[buffer(0)]]) {
+  float3x2 const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm b/test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm
new file mode 100644
index 0000000000..550ed15496
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 8
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v2float = OpTypeVector %float 2
+%mat3v2float = OpTypeMatrix %v2float 3
+    %u_block = OpTypeStruct %mat3v2float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat3v2float = OpTypePointer Uniform %mat3v2float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat3v2float %u %uint_0
+         %15 = OpLoad %mat3v2float %14
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl b/test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl
new file mode 100644
index 0000000000..5f16f764b7
--- /dev/null
+++ b/test/buffer/uniform/types/mat3x2.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat3x2<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat4x4.wgsl b/test/buffer/uniform/types/mat4x4.wgsl
new file mode 100644
index 0000000000..7d2a57b0b1
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl b/test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl
new file mode 100644
index 0000000000..0b14de0372
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.hlsl
@@ -0,0 +1,17 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[4];
+};
+
+float4x4 tint_symbol(uint4 buffer[4], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const uint scalar_offset_1 = ((offset + 16u)) / 4;
+  const uint scalar_offset_2 = ((offset + 32u)) / 4;
+  const uint scalar_offset_3 = ((offset + 48u)) / 4;
+  return float4x4(asfloat(buffer[scalar_offset / 4]), asfloat(buffer[scalar_offset_1 / 4]), asfloat(buffer[scalar_offset_2 / 4]), asfloat(buffer[scalar_offset_3 / 4]));
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float4x4 x = tint_symbol(u, 0u);
+  return;
+}
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.msl b/test/buffer/uniform/types/mat4x4.wgsl.expected.msl
new file mode 100644
index 0000000000..5a16521c42
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float4x4* tint_symbol_1 [[buffer(0)]]) {
+  float4x4 const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm b/test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm
new file mode 100644
index 0000000000..eb84aa3485
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.spvasm
@@ -0,0 +1,37 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 16
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpMemberDecorate %u_block 0 ColMajor
+               OpMemberDecorate %u_block 0 MatrixStride 16
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+%mat4v4float = OpTypeMatrix %v4float 4
+    %u_block = OpTypeStruct %mat4v4float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_mat4v4float = OpTypePointer Uniform %mat4v4float
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %14 = OpAccessChain %_ptr_Uniform_mat4v4float %u %uint_0
+         %15 = OpLoad %mat4v4float %14
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl b/test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl
new file mode 100644
index 0000000000..a1188d409b
--- /dev/null
+++ b/test/buffer/uniform/types/mat4x4.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : mat4x4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/struct.wgsl b/test/buffer/uniform/types/struct.wgsl
new file mode 100644
index 0000000000..135bcab6fd
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl
@@ -0,0 +1,14 @@
+struct Inner {
+  f : f32;
+};
+struct S {
+  inner : Inner;
+};
+
+[[group(0), binding(0)]]
+var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.hlsl b/test/buffer/uniform/types/struct.wgsl.expected.hlsl
new file mode 100644
index 0000000000..d15eb31450
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.hlsl
@@ -0,0 +1,27 @@
+struct Inner {
+  float f;
+};
+struct S {
+  Inner inner;
+};
+
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+Inner tint_symbol_1(uint4 buffer[1], uint offset) {
+  const uint scalar_offset = ((offset + 0u)) / 4;
+  const Inner tint_symbol_3 = {asfloat(buffer[scalar_offset / 4][scalar_offset % 4])};
+  return tint_symbol_3;
+}
+
+S tint_symbol(uint4 buffer[1], uint offset) {
+  const S tint_symbol_4 = {tint_symbol_1(buffer, (offset + 0u))};
+  return tint_symbol_4;
+}
+
+[numthreads(1, 1, 1)]
+void main() {
+  const S x = tint_symbol(u, 0u);
+  return;
+}
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.msl b/test/buffer/uniform/types/struct.wgsl.expected.msl
new file mode 100644
index 0000000000..4e2f5c9797
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.msl
@@ -0,0 +1,15 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct Inner {
+  /* 0x0000 */ float f;
+};
+struct S {
+  /* 0x0000 */ Inner inner;
+};
+
+kernel void tint_symbol(const constant S* tint_symbol_1 [[buffer(0)]]) {
+  S const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.spvasm b/test/buffer/uniform/types/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000000..b8d2ad8fce
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 11
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "inner"
+               OpName %Inner "Inner"
+               OpMemberName %Inner 0 "f"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %S Block
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %Inner 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+      %Inner = OpTypeStruct %float
+          %S = OpTypeStruct %Inner
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+          %u = OpVariable %_ptr_Uniform_S Uniform
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %10 = OpLoad %S %u
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/struct.wgsl.expected.wgsl b/test/buffer/uniform/types/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000000..fed4de3d78
--- /dev/null
+++ b/test/buffer/uniform/types/struct.wgsl.expected.wgsl
@@ -0,0 +1,14 @@
+struct Inner {
+  f : f32;
+};
+
+struct S {
+  inner : Inner;
+};
+
+[[group(0), binding(0)]] var<uniform> u : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/u32.wgsl b/test/buffer/uniform/types/u32.wgsl
new file mode 100644
index 0000000000..a1c80a397b
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.hlsl b/test/buffer/uniform/types/u32.wgsl.expected.hlsl
new file mode 100644
index 0000000000..f033091a5a
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const uint x = u[0].x;
+  return;
+}
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.msl b/test/buffer/uniform/types/u32.wgsl.expected.msl
new file mode 100644
index 0000000000..4d088dd484
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant uint* tint_symbol_1 [[buffer(0)]]) {
+  uint const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.spvasm b/test/buffer/uniform/types/u32.wgsl.expected.spvasm
new file mode 100644
index 0000000000..b73dd01a1f
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.spvasm
@@ -0,0 +1,32 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 13
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+       %uint = OpTypeInt 32 0
+    %u_block = OpTypeStruct %uint
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_uint = OpTypePointer Uniform %uint
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %11 = OpAccessChain %_ptr_Uniform_uint %u %uint_0
+         %12 = OpLoad %uint %11
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/u32.wgsl.expected.wgsl b/test/buffer/uniform/types/u32.wgsl.expected.wgsl
new file mode 100644
index 0000000000..6751a05e52
--- /dev/null
+++ b/test/buffer/uniform/types/u32.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : u32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/vec2.wgsl b/test/buffer/uniform/types/vec2.wgsl
new file mode 100644
index 0000000000..8f2debec27
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.hlsl b/test/buffer/uniform/types/vec2.wgsl.expected.hlsl
new file mode 100644
index 0000000000..a36e25c3a9
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const int2 x = asint(u[0].xy);
+  return;
+}
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.msl b/test/buffer/uniform/types/vec2.wgsl.expected.msl
new file mode 100644
index 0000000000..3f1079516e
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant int2* tint_symbol_1 [[buffer(0)]]) {
+  int2 const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.spvasm b/test/buffer/uniform/types/vec2.wgsl.expected.spvasm
new file mode 100644
index 0000000000..913911a522
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+        %int = OpTypeInt 32 1
+      %v2int = OpTypeVector %int 2
+    %u_block = OpTypeStruct %v2int
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v2int = OpTypePointer Uniform %v2int
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %13 = OpAccessChain %_ptr_Uniform_v2int %u %uint_0
+         %14 = OpLoad %v2int %13
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/vec2.wgsl.expected.wgsl b/test/buffer/uniform/types/vec2.wgsl.expected.wgsl
new file mode 100644
index 0000000000..61695c1daf
--- /dev/null
+++ b/test/buffer/uniform/types/vec2.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : vec2<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/vec3.wgsl b/test/buffer/uniform/types/vec3.wgsl
new file mode 100644
index 0000000000..012d269fa4
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.hlsl b/test/buffer/uniform/types/vec3.wgsl.expected.hlsl
new file mode 100644
index 0000000000..8e081c8a73
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const uint3 x = u[0].xyz;
+  return;
+}
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.msl b/test/buffer/uniform/types/vec3.wgsl.expected.msl
new file mode 100644
index 0000000000..407c0e7545
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant uint3* tint_symbol_1 [[buffer(0)]]) {
+  uint3 const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.spvasm b/test/buffer/uniform/types/vec3.wgsl.expected.spvasm
new file mode 100644
index 0000000000..52fd93a432
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.spvasm
@@ -0,0 +1,33 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 14
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+       %uint = OpTypeInt 32 0
+     %v3uint = OpTypeVector %uint 3
+    %u_block = OpTypeStruct %v3uint
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v3uint = OpTypePointer Uniform %v3uint
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %12 = OpAccessChain %_ptr_Uniform_v3uint %u %uint_0
+         %13 = OpLoad %v3uint %12
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/vec3.wgsl.expected.wgsl b/test/buffer/uniform/types/vec3.wgsl.expected.wgsl
new file mode 100644
index 0000000000..d8d99e8ea4
--- /dev/null
+++ b/test/buffer/uniform/types/vec3.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : vec3<u32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/vec4.wgsl b/test/buffer/uniform/types/vec4.wgsl
new file mode 100644
index 0000000000..24d63f09fc
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl
@@ -0,0 +1,7 @@
+[[group(0), binding(0)]]
+var<uniform> u : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.hlsl b/test/buffer/uniform/types/vec4.wgsl.expected.hlsl
new file mode 100644
index 0000000000..2ae2ec1970
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.hlsl
@@ -0,0 +1,9 @@
+cbuffer cbuffer_u : register(b0, space0) {
+  uint4 u[1];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  const float4 x = asfloat(u[0]);
+  return;
+}
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.msl b/test/buffer/uniform/types/vec4.wgsl.expected.msl
new file mode 100644
index 0000000000..28d31ecb39
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.msl
@@ -0,0 +1,8 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(const constant float4* tint_symbol_1 [[buffer(0)]]) {
+  float4 const x = *(tint_symbol_1);
+  return;
+}
+
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.spvasm b/test/buffer/uniform/types/vec4.wgsl.expected.spvasm
new file mode 100644
index 0000000000..f95df57838
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main"
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %u_block "u_block"
+               OpMemberName %u_block 0 "inner"
+               OpName %u "u"
+               OpName %main "main"
+               OpDecorate %u_block Block
+               OpMemberDecorate %u_block 0 Offset 0
+               OpDecorate %u NonWritable
+               OpDecorate %u DescriptorSet 0
+               OpDecorate %u Binding 0
+      %float = OpTypeFloat 32
+    %v4float = OpTypeVector %float 4
+    %u_block = OpTypeStruct %v4float
+%_ptr_Uniform_u_block = OpTypePointer Uniform %u_block
+          %u = OpVariable %_ptr_Uniform_u_block Uniform
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %uint = OpTypeInt 32 0
+     %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_v4float = OpTypePointer Uniform %v4float
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %13 = OpAccessChain %_ptr_Uniform_v4float %u %uint_0
+         %14 = OpLoad %v4float %13
+               OpReturn
+               OpFunctionEnd
diff --git a/test/buffer/uniform/types/vec4.wgsl.expected.wgsl b/test/buffer/uniform/types/vec4.wgsl.expected.wgsl
new file mode 100644
index 0000000000..fe83de0469
--- /dev/null
+++ b/test/buffer/uniform/types/vec4.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[group(0), binding(0)]] var<uniform> u : vec4<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  let x = u;
+}