diff --git a/test/var/initialization/function/array.wgsl b/test/var/initialization/function/array.wgsl
new file mode 100644
index 0000000..2fc36d9
--- /dev/null
+++ b/test/var/initialization/function/array.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    var v : array<i32, 3>;
+    ignore(v);
+}
diff --git a/test/var/initialization/function/array.wgsl.expected.hlsl b/test/var/initialization/function/array.wgsl.expected.hlsl
new file mode 100644
index 0000000..ecb9e9f
--- /dev/null
+++ b/test/var/initialization/function/array.wgsl.expected.hlsl
@@ -0,0 +1,10 @@
+struct tint_array_wrapper {
+  int arr[3];
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  tint_array_wrapper v = {{0, 0, 0}};
+  v;
+  return;
+}
diff --git a/test/var/initialization/function/array.wgsl.expected.msl b/test/var/initialization/function/array.wgsl.expected.msl
new file mode 100644
index 0000000..8ca89e8
--- /dev/null
+++ b/test/var/initialization/function/array.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  int arr[3];
+};
+
+kernel void tint_symbol() {
+  tint_array_wrapper v = {};
+  (void) v;
+  return;
+}
+
diff --git a/test/var/initialization/function/array.wgsl.expected.spvasm b/test/var/initialization/function/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..ced29ed
--- /dev/null
+++ b/test/var/initialization/function/array.wgsl.expected.spvasm
@@ -0,0 +1,26 @@
+; 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 %main "main"
+               OpName %v "v"
+               OpDecorate %_arr_int_uint_3 ArrayStride 4
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_3 = OpConstant %uint 3
+%_arr_int_uint_3 = OpTypeArray %int %uint_3
+%_ptr_Function__arr_int_uint_3 = OpTypePointer Function %_arr_int_uint_3
+         %11 = OpConstantNull %_arr_int_uint_3
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function__arr_int_uint_3 Function %11
+         %13 = OpLoad %_arr_int_uint_3 %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/function/array.wgsl.expected.wgsl b/test/var/initialization/function/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..3db0d56
--- /dev/null
+++ b/test/var/initialization/function/array.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : array<i32, 3>;
+  ignore(v);
+}
diff --git a/test/var/initialization/function/matrix.wgsl b/test/var/initialization/function/matrix.wgsl
new file mode 100644
index 0000000..eb65b1e
--- /dev/null
+++ b/test/var/initialization/function/matrix.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    var v : mat2x3<f32>;
+    ignore(v);
+}
diff --git a/test/var/initialization/function/matrix.wgsl.expected.hlsl b/test/var/initialization/function/matrix.wgsl.expected.hlsl
new file mode 100644
index 0000000..ec24c8b
--- /dev/null
+++ b/test/var/initialization/function/matrix.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void main() {
+  float2x3 v = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  v;
+  return;
+}
diff --git a/test/var/initialization/function/matrix.wgsl.expected.msl b/test/var/initialization/function/matrix.wgsl.expected.msl
new file mode 100644
index 0000000..951b502
--- /dev/null
+++ b/test/var/initialization/function/matrix.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  float2x3 v = float2x3(0.0f);
+  (void) v;
+  return;
+}
+
diff --git a/test/var/initialization/function/matrix.wgsl.expected.spvasm b/test/var/initialization/function/matrix.wgsl.expected.spvasm
new file mode 100644
index 0000000..5b2ef08
--- /dev/null
+++ b/test/var/initialization/function/matrix.wgsl.expected.spvasm
@@ -0,0 +1,24 @@
+; 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 %main "main"
+               OpName %v "v"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Function_mat2v3float = OpTypePointer Function %mat2v3float
+         %10 = OpConstantNull %mat2v3float
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_mat2v3float Function %10
+         %12 = OpLoad %mat2v3float %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/function/matrix.wgsl.expected.wgsl b/test/var/initialization/function/matrix.wgsl.expected.wgsl
new file mode 100644
index 0000000..17e072b
--- /dev/null
+++ b/test/var/initialization/function/matrix.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : mat2x3<f32>;
+  ignore(v);
+}
diff --git a/test/var/initialization/function/scalar.wgsl b/test/var/initialization/function/scalar.wgsl
new file mode 100644
index 0000000..81ccc32
--- /dev/null
+++ b/test/var/initialization/function/scalar.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    var v : i32;
+    ignore(v);
+}
diff --git a/test/var/initialization/function/scalar.wgsl.expected.hlsl b/test/var/initialization/function/scalar.wgsl.expected.hlsl
new file mode 100644
index 0000000..d96e258
--- /dev/null
+++ b/test/var/initialization/function/scalar.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void main() {
+  int v = 0;
+  v;
+  return;
+}
diff --git a/test/var/initialization/function/scalar.wgsl.expected.msl b/test/var/initialization/function/scalar.wgsl.expected.msl
new file mode 100644
index 0000000..d0221ca
--- /dev/null
+++ b/test/var/initialization/function/scalar.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  int v = 0;
+  (void) v;
+  return;
+}
+
diff --git a/test/var/initialization/function/scalar.wgsl.expected.spvasm b/test/var/initialization/function/scalar.wgsl.expected.spvasm
new file mode 100644
index 0000000..c67d83d
--- /dev/null
+++ b/test/var/initialization/function/scalar.wgsl.expected.spvasm
@@ -0,0 +1,22 @@
+; 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 %main "main"
+               OpName %v "v"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+%_ptr_Function_int = OpTypePointer Function %int
+          %8 = OpConstantNull %int
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_int Function %8
+         %10 = OpLoad %int %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/function/scalar.wgsl.expected.wgsl b/test/var/initialization/function/scalar.wgsl.expected.wgsl
new file mode 100644
index 0000000..73c1726
--- /dev/null
+++ b/test/var/initialization/function/scalar.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : i32;
+  ignore(v);
+}
diff --git a/test/var/initialization/function/struct.wgsl b/test/var/initialization/function/struct.wgsl
new file mode 100644
index 0000000..35fd43d
--- /dev/null
+++ b/test/var/initialization/function/struct.wgsl
@@ -0,0 +1,10 @@
+struct S {
+    a : i32;
+    b : f32;
+};
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    var v : S;
+    ignore(v);
+}
diff --git a/test/var/initialization/function/struct.wgsl.expected.hlsl b/test/var/initialization/function/struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..2493fcb
--- /dev/null
+++ b/test/var/initialization/function/struct.wgsl.expected.hlsl
@@ -0,0 +1,11 @@
+struct S {
+  int a;
+  float b;
+};
+
+[numthreads(1, 1, 1)]
+void main() {
+  S v = {0, 0.0f};
+  v;
+  return;
+}
diff --git a/test/var/initialization/function/struct.wgsl.expected.msl b/test/var/initialization/function/struct.wgsl.expected.msl
new file mode 100644
index 0000000..ffec229
--- /dev/null
+++ b/test/var/initialization/function/struct.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int a;
+  float b;
+};
+
+kernel void tint_symbol() {
+  S v = {};
+  (void) v;
+  return;
+}
+
diff --git a/test/var/initialization/function/struct.wgsl.expected.spvasm b/test/var/initialization/function/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..9c865a6
--- /dev/null
+++ b/test/var/initialization/function/struct.wgsl.expected.spvasm
@@ -0,0 +1,29 @@
+; 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 %main "main"
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpName %v "v"
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %float = OpTypeFloat 32
+          %S = OpTypeStruct %int %float
+%_ptr_Function_S = OpTypePointer Function %S
+         %10 = OpConstantNull %S
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_S Function %10
+         %12 = OpLoad %S %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/function/struct.wgsl.expected.wgsl b/test/var/initialization/function/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..57ec0e6
--- /dev/null
+++ b/test/var/initialization/function/struct.wgsl.expected.wgsl
@@ -0,0 +1,10 @@
+struct S {
+  a : i32;
+  b : f32;
+};
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : S;
+  ignore(v);
+}
diff --git a/test/var/initialization/function/vector.wgsl b/test/var/initialization/function/vector.wgsl
new file mode 100644
index 0000000..24c52b4
--- /dev/null
+++ b/test/var/initialization/function/vector.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    var v : vec3<i32>;
+    ignore(v);
+}
diff --git a/test/var/initialization/function/vector.wgsl.expected.hlsl b/test/var/initialization/function/vector.wgsl.expected.hlsl
new file mode 100644
index 0000000..f595b78
--- /dev/null
+++ b/test/var/initialization/function/vector.wgsl.expected.hlsl
@@ -0,0 +1,6 @@
+[numthreads(1, 1, 1)]
+void main() {
+  int3 v = int3(0, 0, 0);
+  v;
+  return;
+}
diff --git a/test/var/initialization/function/vector.wgsl.expected.msl b/test/var/initialization/function/vector.wgsl.expected.msl
new file mode 100644
index 0000000..c1b8f4d
--- /dev/null
+++ b/test/var/initialization/function/vector.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  int3 v = 0;
+  (void) v;
+  return;
+}
+
diff --git a/test/var/initialization/function/vector.wgsl.expected.spvasm b/test/var/initialization/function/vector.wgsl.expected.spvasm
new file mode 100644
index 0000000..aa8b479
--- /dev/null
+++ b/test/var/initialization/function/vector.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+; 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 %main "main"
+               OpName %v "v"
+       %void = OpTypeVoid
+          %1 = OpTypeFunction %void
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_ptr_Function_v3int = OpTypePointer Function %v3int
+          %9 = OpConstantNull %v3int
+       %main = OpFunction %void None %1
+          %4 = OpLabel
+          %v = OpVariable %_ptr_Function_v3int Function %9
+         %11 = OpLoad %v3int %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/function/vector.wgsl.expected.wgsl b/test/var/initialization/function/vector.wgsl.expected.wgsl
new file mode 100644
index 0000000..6fd0442
--- /dev/null
+++ b/test/var/initialization/function/vector.wgsl.expected.wgsl
@@ -0,0 +1,5 @@
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  var v : vec3<i32>;
+  ignore(v);
+}
diff --git a/test/var/initialization/private/array.wgsl b/test/var/initialization/private/array.wgsl
new file mode 100644
index 0000000..9dae8f7
--- /dev/null
+++ b/test/var/initialization/private/array.wgsl
@@ -0,0 +1,6 @@
+var<private> v : array<i32, 3>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/private/array.wgsl.expected.hlsl b/test/var/initialization/private/array.wgsl.expected.hlsl
new file mode 100644
index 0000000..2873378
--- /dev/null
+++ b/test/var/initialization/private/array.wgsl.expected.hlsl
@@ -0,0 +1,11 @@
+struct tint_array_wrapper {
+  int arr[3];
+};
+
+static tint_array_wrapper v = {{0, 0, 0}};
+
+[numthreads(1, 1, 1)]
+void main() {
+  v;
+  return;
+}
diff --git a/test/var/initialization/private/array.wgsl.expected.msl b/test/var/initialization/private/array.wgsl.expected.msl
new file mode 100644
index 0000000..ab06870
--- /dev/null
+++ b/test/var/initialization/private/array.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  int arr[3];
+};
+
+kernel void tint_symbol() {
+  thread tint_array_wrapper tint_symbol_1 = {};
+  (void) tint_symbol_1;
+  return;
+}
+
diff --git a/test/var/initialization/private/array.wgsl.expected.spvasm b/test/var/initialization/private/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..1996520
--- /dev/null
+++ b/test/var/initialization/private/array.wgsl.expected.spvasm
@@ -0,0 +1,26 @@
+; 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 %v "v"
+               OpName %main "main"
+               OpDecorate %_arr_int_uint_3 ArrayStride 4
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_3 = OpConstant %uint 3
+%_arr_int_uint_3 = OpTypeArray %int %uint_3
+%_ptr_Private__arr_int_uint_3 = OpTypePointer Private %_arr_int_uint_3
+          %7 = OpConstantNull %_arr_int_uint_3
+          %v = OpVariable %_ptr_Private__arr_int_uint_3 Private %7
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %13 = OpLoad %_arr_int_uint_3 %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/private/array.wgsl.expected.wgsl b/test/var/initialization/private/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..3b54ea9
--- /dev/null
+++ b/test/var/initialization/private/array.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<private> v : array<i32, 3>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/private/matrix.wgsl b/test/var/initialization/private/matrix.wgsl
new file mode 100644
index 0000000..b167840
--- /dev/null
+++ b/test/var/initialization/private/matrix.wgsl
@@ -0,0 +1,6 @@
+var<private> v : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/private/matrix.wgsl.expected.hlsl b/test/var/initialization/private/matrix.wgsl.expected.hlsl
new file mode 100644
index 0000000..df9907e
--- /dev/null
+++ b/test/var/initialization/private/matrix.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+static float2x3 v = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+
+[numthreads(1, 1, 1)]
+void main() {
+  v;
+  return;
+}
diff --git a/test/var/initialization/private/matrix.wgsl.expected.msl b/test/var/initialization/private/matrix.wgsl.expected.msl
new file mode 100644
index 0000000..76f9711
--- /dev/null
+++ b/test/var/initialization/private/matrix.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  thread float2x3 tint_symbol_1 = float2x3(0.0f);
+  (void) tint_symbol_1;
+  return;
+}
+
diff --git a/test/var/initialization/private/matrix.wgsl.expected.spvasm b/test/var/initialization/private/matrix.wgsl.expected.spvasm
new file mode 100644
index 0000000..04a8f59
--- /dev/null
+++ b/test/var/initialization/private/matrix.wgsl.expected.spvasm
@@ -0,0 +1,24 @@
+; 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 %v "v"
+               OpName %main "main"
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Private_mat2v3float = OpTypePointer Private %mat2v3float
+          %6 = OpConstantNull %mat2v3float
+          %v = OpVariable %_ptr_Private_mat2v3float Private %6
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %12 = OpLoad %mat2v3float %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/private/matrix.wgsl.expected.wgsl b/test/var/initialization/private/matrix.wgsl.expected.wgsl
new file mode 100644
index 0000000..ce06808
--- /dev/null
+++ b/test/var/initialization/private/matrix.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<private> v : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/private/scalar.wgsl b/test/var/initialization/private/scalar.wgsl
new file mode 100644
index 0000000..1179686
--- /dev/null
+++ b/test/var/initialization/private/scalar.wgsl
@@ -0,0 +1,6 @@
+var<private> v : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/private/scalar.wgsl.expected.hlsl b/test/var/initialization/private/scalar.wgsl.expected.hlsl
new file mode 100644
index 0000000..0daf53c
--- /dev/null
+++ b/test/var/initialization/private/scalar.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+static int v = 0;
+
+[numthreads(1, 1, 1)]
+void main() {
+  v;
+  return;
+}
diff --git a/test/var/initialization/private/scalar.wgsl.expected.msl b/test/var/initialization/private/scalar.wgsl.expected.msl
new file mode 100644
index 0000000..c96b599
--- /dev/null
+++ b/test/var/initialization/private/scalar.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  thread int tint_symbol_1 = 0;
+  (void) tint_symbol_1;
+  return;
+}
+
diff --git a/test/var/initialization/private/scalar.wgsl.expected.spvasm b/test/var/initialization/private/scalar.wgsl.expected.spvasm
new file mode 100644
index 0000000..3deea8c
--- /dev/null
+++ b/test/var/initialization/private/scalar.wgsl.expected.spvasm
@@ -0,0 +1,22 @@
+; 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 %v "v"
+               OpName %main "main"
+        %int = OpTypeInt 32 1
+%_ptr_Private_int = OpTypePointer Private %int
+          %4 = OpConstantNull %int
+          %v = OpVariable %_ptr_Private_int Private %4
+       %void = OpTypeVoid
+          %5 = OpTypeFunction %void
+       %main = OpFunction %void None %5
+          %8 = OpLabel
+         %10 = OpLoad %int %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/private/scalar.wgsl.expected.wgsl b/test/var/initialization/private/scalar.wgsl.expected.wgsl
new file mode 100644
index 0000000..e634b37
--- /dev/null
+++ b/test/var/initialization/private/scalar.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<private> v : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/private/struct.wgsl b/test/var/initialization/private/struct.wgsl
new file mode 100644
index 0000000..0c64c99
--- /dev/null
+++ b/test/var/initialization/private/struct.wgsl
@@ -0,0 +1,11 @@
+struct S {
+    a : i32;
+    b : f32;
+};
+
+var<private> v : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/private/struct.wgsl.expected.hlsl b/test/var/initialization/private/struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..ca8d6ca
--- /dev/null
+++ b/test/var/initialization/private/struct.wgsl.expected.hlsl
@@ -0,0 +1,12 @@
+struct S {
+  int a;
+  float b;
+};
+
+static S v = {0, 0.0f};
+
+[numthreads(1, 1, 1)]
+void main() {
+  v;
+  return;
+}
diff --git a/test/var/initialization/private/struct.wgsl.expected.msl b/test/var/initialization/private/struct.wgsl.expected.msl
new file mode 100644
index 0000000..948d0b7
--- /dev/null
+++ b/test/var/initialization/private/struct.wgsl.expected.msl
@@ -0,0 +1,14 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int a;
+  float b;
+};
+
+kernel void tint_symbol() {
+  thread S tint_symbol_1 = {};
+  (void) tint_symbol_1;
+  return;
+}
+
diff --git a/test/var/initialization/private/struct.wgsl.expected.spvasm b/test/var/initialization/private/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..ff2048a
--- /dev/null
+++ b/test/var/initialization/private/struct.wgsl.expected.spvasm
@@ -0,0 +1,29 @@
+; 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 %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpName %v "v"
+               OpName %main "main"
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+        %int = OpTypeInt 32 1
+      %float = OpTypeFloat 32
+          %S = OpTypeStruct %int %float
+%_ptr_Private_S = OpTypePointer Private %S
+          %6 = OpConstantNull %S
+          %v = OpVariable %_ptr_Private_S Private %6
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %12 = OpLoad %S %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/private/struct.wgsl.expected.wgsl b/test/var/initialization/private/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..f677756
--- /dev/null
+++ b/test/var/initialization/private/struct.wgsl.expected.wgsl
@@ -0,0 +1,11 @@
+struct S {
+  a : i32;
+  b : f32;
+};
+
+var<private> v : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/private/vector.wgsl b/test/var/initialization/private/vector.wgsl
new file mode 100644
index 0000000..db1904e
--- /dev/null
+++ b/test/var/initialization/private/vector.wgsl
@@ -0,0 +1,6 @@
+var<private> v : vec3<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/private/vector.wgsl.expected.hlsl b/test/var/initialization/private/vector.wgsl.expected.hlsl
new file mode 100644
index 0000000..1f4a920
--- /dev/null
+++ b/test/var/initialization/private/vector.wgsl.expected.hlsl
@@ -0,0 +1,7 @@
+static int3 v = int3(0, 0, 0);
+
+[numthreads(1, 1, 1)]
+void main() {
+  v;
+  return;
+}
diff --git a/test/var/initialization/private/vector.wgsl.expected.msl b/test/var/initialization/private/vector.wgsl.expected.msl
new file mode 100644
index 0000000..6398a63
--- /dev/null
+++ b/test/var/initialization/private/vector.wgsl.expected.msl
@@ -0,0 +1,9 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+  thread int3 tint_symbol_1 = 0;
+  (void) tint_symbol_1;
+  return;
+}
+
diff --git a/test/var/initialization/private/vector.wgsl.expected.spvasm b/test/var/initialization/private/vector.wgsl.expected.spvasm
new file mode 100644
index 0000000..f5086e9
--- /dev/null
+++ b/test/var/initialization/private/vector.wgsl.expected.spvasm
@@ -0,0 +1,23 @@
+; 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 %v "v"
+               OpName %main "main"
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_ptr_Private_v3int = OpTypePointer Private %v3int
+          %5 = OpConstantNull %v3int
+          %v = OpVariable %_ptr_Private_v3int Private %5
+       %void = OpTypeVoid
+          %6 = OpTypeFunction %void
+       %main = OpFunction %void None %6
+          %9 = OpLabel
+         %11 = OpLoad %v3int %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/private/vector.wgsl.expected.wgsl b/test/var/initialization/private/vector.wgsl.expected.wgsl
new file mode 100644
index 0000000..e5d11f7
--- /dev/null
+++ b/test/var/initialization/private/vector.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<private> v : vec3<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/workgroup/array.wgsl b/test/var/initialization/workgroup/array.wgsl
new file mode 100644
index 0000000..9ba52fc
--- /dev/null
+++ b/test/var/initialization/workgroup/array.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : array<i32, 3>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/workgroup/array.wgsl.expected.hlsl b/test/var/initialization/workgroup/array.wgsl.expected.hlsl
new file mode 100644
index 0000000..e711cb6
--- /dev/null
+++ b/test/var/initialization/workgroup/array.wgsl.expected.hlsl
@@ -0,0 +1,21 @@
+struct tint_array_wrapper {
+  int arr[3];
+};
+
+groupshared tint_array_wrapper v;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint local_invocation_index = tint_symbol.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    const tint_array_wrapper tint_symbol_2 = {{0, 0, 0}};
+    v = tint_symbol_2;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  v;
+  return;
+}
diff --git a/test/var/initialization/workgroup/array.wgsl.expected.msl b/test/var/initialization/workgroup/array.wgsl.expected.msl
new file mode 100644
index 0000000..f177453
--- /dev/null
+++ b/test/var/initialization/workgroup/array.wgsl.expected.msl
@@ -0,0 +1,18 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct tint_array_wrapper {
+  int arr[3];
+};
+
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup tint_array_wrapper tint_symbol_3;
+  if ((local_invocation_index == 0u)) {
+    tint_array_wrapper const tint_symbol_2 = {.arr={}};
+    tint_symbol_3 = tint_symbol_2;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (void) tint_symbol_3;
+  return;
+}
+
diff --git a/test/var/initialization/workgroup/array.wgsl.expected.spvasm b/test/var/initialization/workgroup/array.wgsl.expected.spvasm
new file mode 100644
index 0000000..c1ef12c
--- /dev/null
+++ b/test/var/initialization/workgroup/array.wgsl.expected.spvasm
@@ -0,0 +1,43 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %v "v"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpDecorate %_arr_int_uint_3 ArrayStride 4
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+       %uint = OpTypeInt 32 0
+     %uint_3 = OpConstant %uint 3
+%_arr_int_uint_3 = OpTypeArray %int %uint_3
+%_ptr_Workgroup__arr_int_uint_3 = OpTypePointer Workgroup %_arr_int_uint_3
+          %v = OpVariable %_ptr_Workgroup__arr_int_uint_3 Workgroup
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %19 = OpConstantNull %_arr_int_uint_3
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+         %13 = OpLoad %uint %tint_symbol
+         %15 = OpIEqual %bool %13 %uint_0
+               OpSelectionMerge %17 None
+               OpBranchConditional %15 %18 %17
+         %18 = OpLabel
+               OpStore %v %19
+               OpBranch %17
+         %17 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %24 = OpLoad %_arr_int_uint_3 %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/workgroup/array.wgsl.expected.wgsl b/test/var/initialization/workgroup/array.wgsl.expected.wgsl
new file mode 100644
index 0000000..5cfadc1
--- /dev/null
+++ b/test/var/initialization/workgroup/array.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : array<i32, 3>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/workgroup/matrix.wgsl b/test/var/initialization/workgroup/matrix.wgsl
new file mode 100644
index 0000000..ddc128c
--- /dev/null
+++ b/test/var/initialization/workgroup/matrix.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/workgroup/matrix.wgsl.expected.hlsl b/test/var/initialization/workgroup/matrix.wgsl.expected.hlsl
new file mode 100644
index 0000000..1590817
--- /dev/null
+++ b/test/var/initialization/workgroup/matrix.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+groupshared float2x3 v;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint local_invocation_index = tint_symbol.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    v = float2x3(0.0f, 0.0f, 0.0f, 0.0f, 0.0f, 0.0f);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  v;
+  return;
+}
diff --git a/test/var/initialization/workgroup/matrix.wgsl.expected.msl b/test/var/initialization/workgroup/matrix.wgsl.expected.msl
new file mode 100644
index 0000000..52f49c4
--- /dev/null
+++ b/test/var/initialization/workgroup/matrix.wgsl.expected.msl
@@ -0,0 +1,192 @@
+SKIP: crbug.com/tint/938
+
+
+
+Validation Failure:
+
+Compilation failed:
+
+program_source:5:24: error: no matching constructor for initialization of 'threadgroup metal::float2x3' (aka 'threadgroup matrix<float, 2, 3>')
+  threadgroup float2x3 tint_symbol_2;
+                       ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:269:23: note: candidate template ignored: requirement 'sizeof...(U) == 2' was not satisfied [with U = <>]
+  METAL_FUNC explicit matrix(initializer_list<U>... cols) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:281:23: note: candidate template ignored: requirement '(sizeof...(U) == 2) || (sizeof...(U) == 2 * 3)' was not satisfied [with U = <>]
+  METAL_FUNC explicit matrix(U... vals) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:299:23: note: candidate template ignored: requirement 'sizeof...(U) == 2' was not satisfied [with U = <>]
+  METAL_FUNC explicit matrix(initializer_list<U>... cols) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:311:23: note: candidate template ignored: requirement '(sizeof...(U) == 2) || (sizeof...(U) == 2 * 3)' was not satisfied [with U = <>]
+  METAL_FUNC explicit matrix(U... vals) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:330:23: note: candidate template ignored: requirement 'sizeof...(U) == 2' was not satisfied [with U = <>]
+  METAL_FUNC explicit matrix(initializer_list<U>... cols) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:342:23: note: candidate template ignored: requirement '(sizeof...(U) == 2) || (sizeof...(U) == 2 * 3)' was not satisfied [with U = <>]
+  METAL_FUNC explicit matrix(U... vals) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:56:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(T val, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:62:23: note: candidate constructor template not viable: requires 3 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(cols_init_tag, initializer_list<vec<T, Rows>> cols, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:68:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
+  METAL_FUNC explicit matrix(cols_all_tag, U... cols) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:75:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
+  METAL_FUNC explicit matrix(elems_all_tag, U... elems) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:80:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(initializer_list<T> elems, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:86:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(T val, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:92:23: note: candidate constructor template not viable: requires 3 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(cols_init_tag, initializer_list<vec<T, Rows>> cols, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:98:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
+  METAL_FUNC explicit matrix(cols_all_tag, U... cols) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:105:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
+  METAL_FUNC explicit matrix(elems_all_tag, U... elems) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:110:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(initializer_list<T> elems, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:117:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(T val, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:123:23: note: candidate constructor template not viable: requires 3 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(cols_init_tag, initializer_list<vec<T, Rows>> cols, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:129:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
+  METAL_FUNC explicit matrix(cols_all_tag, U... cols) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:136:23: note: candidate constructor template not viable: requires at least 1 argument, but 0 were provided
+  METAL_FUNC explicit matrix(elems_all_tag, U... elems) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:141:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(initializer_list<T> elems, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:149:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:154:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:159:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:164:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:170:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:177:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:183:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:188:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:193:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:198:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:204:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:211:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:218:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:225:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:232:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:239:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:247:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:255:23: note: candidate constructor template not viable: requires 2 arguments, but 0 were provided
+  METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that, _integer_sequence<int, C...>) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:358:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:365:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:372:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:379:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:387:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:396:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that) thread
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:404:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:411:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:418:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:425:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:433:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:442:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that) constant
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:451:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const thread matrix<U, Cols, Rows> &that) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:460:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const device matrix<U, Cols, Rows> &that) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:469:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const constant matrix<U, Cols, Rows> &that) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:478:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const threadgroup matrix<U, Cols, Rows> &that) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:488:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const threadgroup_imageblock matrix<U, Cols, Rows> &that) ray_data
+                      ^
+/System/Library/PrivateFrameworks/GPUCompiler.framework/Versions/31001/Libraries/lib/clang/31001.189/include/metal/metal_matrix:498:23: note: candidate constructor template not viable: requires single argument 'that', but no arguments were provided
+  METAL_FUNC explicit matrix(const ray_data matrix<U, Cols, Rows> &that) ray_data
+                      ^
+program_source:6:31: warning: equality comparison with extraneous parentheses
+  if ((local_invocation_index == 0u)) {
+       ~~~~~~~~~~~~~~~~~~~~~~~^~~~~
+program_source:6:31: note: remove extraneous parentheses around the comparison to silence this warning
+  if ((local_invocation_index == 0u)) {
+      ~                       ^    ~
+program_source:6:31: note: use '=' to turn this equality comparison into an assignment
+  if ((local_invocation_index == 0u)) {
+                              ^~
+                              =
diff --git a/test/var/initialization/workgroup/matrix.wgsl.expected.spvasm b/test/var/initialization/workgroup/matrix.wgsl.expected.spvasm
new file mode 100644
index 0000000..7101d80
--- /dev/null
+++ b/test/var/initialization/workgroup/matrix.wgsl.expected.spvasm
@@ -0,0 +1,42 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %v "v"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+      %float = OpTypeFloat 32
+    %v3float = OpTypeVector %float 3
+%mat2v3float = OpTypeMatrix %v3float 2
+%_ptr_Workgroup_mat2v3float = OpTypePointer Workgroup %mat2v3float
+          %v = OpVariable %_ptr_Workgroup_mat2v3float Workgroup
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %19 = OpConstantNull %mat2v3float
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+         %13 = OpLoad %uint %tint_symbol
+         %15 = OpIEqual %bool %13 %uint_0
+               OpSelectionMerge %17 None
+               OpBranchConditional %15 %18 %17
+         %18 = OpLabel
+               OpStore %v %19
+               OpBranch %17
+         %17 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %24 = OpLoad %mat2v3float %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/workgroup/matrix.wgsl.expected.wgsl b/test/var/initialization/workgroup/matrix.wgsl.expected.wgsl
new file mode 100644
index 0000000..6477351
--- /dev/null
+++ b/test/var/initialization/workgroup/matrix.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : mat2x3<f32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/workgroup/scalar.wgsl b/test/var/initialization/workgroup/scalar.wgsl
new file mode 100644
index 0000000..58f1148
--- /dev/null
+++ b/test/var/initialization/workgroup/scalar.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/workgroup/scalar.wgsl.expected.hlsl b/test/var/initialization/workgroup/scalar.wgsl.expected.hlsl
new file mode 100644
index 0000000..8f52058
--- /dev/null
+++ b/test/var/initialization/workgroup/scalar.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+groupshared int v;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint local_invocation_index = tint_symbol.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    v = 0;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  v;
+  return;
+}
diff --git a/test/var/initialization/workgroup/scalar.wgsl.expected.msl b/test/var/initialization/workgroup/scalar.wgsl.expected.msl
new file mode 100644
index 0000000..f63e941
--- /dev/null
+++ b/test/var/initialization/workgroup/scalar.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup int tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    tint_symbol_2 = int();
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (void) tint_symbol_2;
+  return;
+}
+
diff --git a/test/var/initialization/workgroup/scalar.wgsl.expected.spvasm b/test/var/initialization/workgroup/scalar.wgsl.expected.spvasm
new file mode 100644
index 0000000..aede25b
--- /dev/null
+++ b/test/var/initialization/workgroup/scalar.wgsl.expected.spvasm
@@ -0,0 +1,40 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 23
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %v "v"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+          %v = OpVariable %_ptr_Workgroup_int Workgroup
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+          %7 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %17 = OpConstantNull %int
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %7
+         %10 = OpLabel
+         %11 = OpLoad %uint %tint_symbol
+         %13 = OpIEqual %bool %11 %uint_0
+               OpSelectionMerge %15 None
+               OpBranchConditional %13 %16 %15
+         %16 = OpLabel
+               OpStore %v %17
+               OpBranch %15
+         %15 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %22 = OpLoad %int %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/workgroup/scalar.wgsl.expected.wgsl b/test/var/initialization/workgroup/scalar.wgsl.expected.wgsl
new file mode 100644
index 0000000..e92577b
--- /dev/null
+++ b/test/var/initialization/workgroup/scalar.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : i32;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/workgroup/struct.wgsl b/test/var/initialization/workgroup/struct.wgsl
new file mode 100644
index 0000000..81a88b2
--- /dev/null
+++ b/test/var/initialization/workgroup/struct.wgsl
@@ -0,0 +1,11 @@
+struct S {
+    a : i32;
+    b : f32;
+};
+
+var<workgroup> v : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/workgroup/struct.wgsl.expected.hlsl b/test/var/initialization/workgroup/struct.wgsl.expected.hlsl
new file mode 100644
index 0000000..84ee4d6
--- /dev/null
+++ b/test/var/initialization/workgroup/struct.wgsl.expected.hlsl
@@ -0,0 +1,22 @@
+struct S {
+  int a;
+  float b;
+};
+
+groupshared S v;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint local_invocation_index = tint_symbol.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    const S tint_symbol_2 = {0, 0.0f};
+    v = tint_symbol_2;
+  }
+  GroupMemoryBarrierWithGroupSync();
+  v;
+  return;
+}
diff --git a/test/var/initialization/workgroup/struct.wgsl.expected.msl b/test/var/initialization/workgroup/struct.wgsl.expected.msl
new file mode 100644
index 0000000..88dc937
--- /dev/null
+++ b/test/var/initialization/workgroup/struct.wgsl.expected.msl
@@ -0,0 +1,19 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+  int a;
+  float b;
+};
+
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup S tint_symbol_3;
+  if ((local_invocation_index == 0u)) {
+    S const tint_symbol_2 = {};
+    tint_symbol_3 = tint_symbol_2;
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (void) tint_symbol_3;
+  return;
+}
+
diff --git a/test/var/initialization/workgroup/struct.wgsl.expected.spvasm b/test/var/initialization/workgroup/struct.wgsl.expected.spvasm
new file mode 100644
index 0000000..83f2f1b
--- /dev/null
+++ b/test/var/initialization/workgroup/struct.wgsl.expected.spvasm
@@ -0,0 +1,47 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %S "S"
+               OpMemberName %S 0 "a"
+               OpMemberName %S 1 "b"
+               OpName %v "v"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpMemberDecorate %S 0 Offset 0
+               OpMemberDecorate %S 1 Offset 4
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %float = OpTypeFloat 32
+          %S = OpTypeStruct %int %float
+%_ptr_Workgroup_S = OpTypePointer Workgroup %S
+          %v = OpVariable %_ptr_Workgroup_S Workgroup
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+          %9 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %19 = OpConstantNull %S
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %9
+         %12 = OpLabel
+         %13 = OpLoad %uint %tint_symbol
+         %15 = OpIEqual %bool %13 %uint_0
+               OpSelectionMerge %17 None
+               OpBranchConditional %15 %18 %17
+         %18 = OpLabel
+               OpStore %v %19
+               OpBranch %17
+         %17 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %24 = OpLoad %S %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/workgroup/struct.wgsl.expected.wgsl b/test/var/initialization/workgroup/struct.wgsl.expected.wgsl
new file mode 100644
index 0000000..77244de
--- /dev/null
+++ b/test/var/initialization/workgroup/struct.wgsl.expected.wgsl
@@ -0,0 +1,11 @@
+struct S {
+  a : i32;
+  b : f32;
+};
+
+var<workgroup> v : S;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}
diff --git a/test/var/initialization/workgroup/vector.wgsl b/test/var/initialization/workgroup/vector.wgsl
new file mode 100644
index 0000000..efad7b1
--- /dev/null
+++ b/test/var/initialization/workgroup/vector.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : vec3<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+    ignore(v);
+}
diff --git a/test/var/initialization/workgroup/vector.wgsl.expected.hlsl b/test/var/initialization/workgroup/vector.wgsl.expected.hlsl
new file mode 100644
index 0000000..bec6192
--- /dev/null
+++ b/test/var/initialization/workgroup/vector.wgsl.expected.hlsl
@@ -0,0 +1,16 @@
+groupshared int3 v;
+
+struct tint_symbol_1 {
+  uint local_invocation_index : SV_GroupIndex;
+};
+
+[numthreads(1, 1, 1)]
+void main(tint_symbol_1 tint_symbol) {
+  const uint local_invocation_index = tint_symbol.local_invocation_index;
+  if ((local_invocation_index == 0u)) {
+    v = int3(0, 0, 0);
+  }
+  GroupMemoryBarrierWithGroupSync();
+  v;
+  return;
+}
diff --git a/test/var/initialization/workgroup/vector.wgsl.expected.msl b/test/var/initialization/workgroup/vector.wgsl.expected.msl
new file mode 100644
index 0000000..55cda8e
--- /dev/null
+++ b/test/var/initialization/workgroup/vector.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol(uint local_invocation_index [[thread_index_in_threadgroup]]) {
+  threadgroup int3 tint_symbol_2;
+  if ((local_invocation_index == 0u)) {
+    tint_symbol_2 = int3();
+  }
+  threadgroup_barrier(mem_flags::mem_threadgroup);
+  (void) tint_symbol_2;
+  return;
+}
+
diff --git a/test/var/initialization/workgroup/vector.wgsl.expected.spvasm b/test/var/initialization/workgroup/vector.wgsl.expected.spvasm
new file mode 100644
index 0000000..a5a02e2
--- /dev/null
+++ b/test/var/initialization/workgroup/vector.wgsl.expected.spvasm
@@ -0,0 +1,41 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 24
+; Schema: 0
+               OpCapability Shader
+               OpMemoryModel Logical GLSL450
+               OpEntryPoint GLCompute %main "main" %tint_symbol
+               OpExecutionMode %main LocalSize 1 1 1
+               OpName %v "v"
+               OpName %tint_symbol "tint_symbol"
+               OpName %main "main"
+               OpDecorate %tint_symbol BuiltIn LocalInvocationIndex
+        %int = OpTypeInt 32 1
+      %v3int = OpTypeVector %int 3
+%_ptr_Workgroup_v3int = OpTypePointer Workgroup %v3int
+          %v = OpVariable %_ptr_Workgroup_v3int Workgroup
+       %uint = OpTypeInt 32 0
+%_ptr_Input_uint = OpTypePointer Input %uint
+%tint_symbol = OpVariable %_ptr_Input_uint Input
+       %void = OpTypeVoid
+          %8 = OpTypeFunction %void
+     %uint_0 = OpConstant %uint 0
+       %bool = OpTypeBool
+         %18 = OpConstantNull %v3int
+     %uint_2 = OpConstant %uint 2
+   %uint_264 = OpConstant %uint 264
+       %main = OpFunction %void None %8
+         %11 = OpLabel
+         %12 = OpLoad %uint %tint_symbol
+         %14 = OpIEqual %bool %12 %uint_0
+               OpSelectionMerge %16 None
+               OpBranchConditional %14 %17 %16
+         %17 = OpLabel
+               OpStore %v %18
+               OpBranch %16
+         %16 = OpLabel
+               OpControlBarrier %uint_2 %uint_2 %uint_264
+         %23 = OpLoad %v3int %v
+               OpReturn
+               OpFunctionEnd
diff --git a/test/var/initialization/workgroup/vector.wgsl.expected.wgsl b/test/var/initialization/workgroup/vector.wgsl.expected.wgsl
new file mode 100644
index 0000000..78f4426
--- /dev/null
+++ b/test/var/initialization/workgroup/vector.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+var<workgroup> v : vec3<i32>;
+
+[[stage(compute), workgroup_size(1)]]
+fn main() {
+  ignore(v);
+}