test: Add tests to verify variables are zero initialized
Bug: tint:938
Fixed: tint:759
Change-Id: I5794d7f40ec154c55240b163316ed45c7f14b190
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/56547
Reviewed-by: James Price <jrprice@google.com>
Reviewed-by: Corentin Wallez <cwallez@chromium.org>
Commit-Queue: Ben Clayton <bclayton@google.com>
Auto-Submit: Ben Clayton <bclayton@google.com>
Kokoro: Kokoro <noreply+kokoro@google.com>
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);
+}
