From fb4e751258c65a9d8481d9d0026e1bf98e582d5b Mon Sep 17 00:00:00 2001 From: Ben Clayton Date: Wed, 30 Jun 2021 15:59:40 +0000 Subject: [PATCH] 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 Reviewed-by: Corentin Wallez Commit-Queue: Ben Clayton Auto-Submit: Ben Clayton Kokoro: Kokoro --- test/var/initialization/function/array.wgsl | 5 + .../function/array.wgsl.expected.hlsl | 10 + .../function/array.wgsl.expected.msl | 13 ++ .../function/array.wgsl.expected.spvasm | 26 +++ .../function/array.wgsl.expected.wgsl | 5 + test/var/initialization/function/matrix.wgsl | 5 + .../function/matrix.wgsl.expected.hlsl | 6 + .../function/matrix.wgsl.expected.msl | 9 + .../function/matrix.wgsl.expected.spvasm | 24 +++ .../function/matrix.wgsl.expected.wgsl | 5 + test/var/initialization/function/scalar.wgsl | 5 + .../function/scalar.wgsl.expected.hlsl | 6 + .../function/scalar.wgsl.expected.msl | 9 + .../function/scalar.wgsl.expected.spvasm | 22 ++ .../function/scalar.wgsl.expected.wgsl | 5 + test/var/initialization/function/struct.wgsl | 10 + .../function/struct.wgsl.expected.hlsl | 11 + .../function/struct.wgsl.expected.msl | 14 ++ .../function/struct.wgsl.expected.spvasm | 29 +++ .../function/struct.wgsl.expected.wgsl | 10 + test/var/initialization/function/vector.wgsl | 5 + .../function/vector.wgsl.expected.hlsl | 6 + .../function/vector.wgsl.expected.msl | 9 + .../function/vector.wgsl.expected.spvasm | 23 +++ .../function/vector.wgsl.expected.wgsl | 5 + test/var/initialization/private/array.wgsl | 6 + .../private/array.wgsl.expected.hlsl | 11 + .../private/array.wgsl.expected.msl | 13 ++ .../private/array.wgsl.expected.spvasm | 26 +++ .../private/array.wgsl.expected.wgsl | 6 + test/var/initialization/private/matrix.wgsl | 6 + .../private/matrix.wgsl.expected.hlsl | 7 + .../private/matrix.wgsl.expected.msl | 9 + .../private/matrix.wgsl.expected.spvasm | 24 +++ .../private/matrix.wgsl.expected.wgsl | 6 + test/var/initialization/private/scalar.wgsl | 6 + .../private/scalar.wgsl.expected.hlsl | 7 + .../private/scalar.wgsl.expected.msl | 9 + .../private/scalar.wgsl.expected.spvasm | 22 ++ .../private/scalar.wgsl.expected.wgsl | 6 + test/var/initialization/private/struct.wgsl | 11 + .../private/struct.wgsl.expected.hlsl | 12 ++ .../private/struct.wgsl.expected.msl | 14 ++ .../private/struct.wgsl.expected.spvasm | 29 +++ .../private/struct.wgsl.expected.wgsl | 11 + test/var/initialization/private/vector.wgsl | 6 + .../private/vector.wgsl.expected.hlsl | 7 + .../private/vector.wgsl.expected.msl | 9 + .../private/vector.wgsl.expected.spvasm | 23 +++ .../private/vector.wgsl.expected.wgsl | 6 + test/var/initialization/workgroup/array.wgsl | 6 + .../workgroup/array.wgsl.expected.hlsl | 21 ++ .../workgroup/array.wgsl.expected.msl | 18 ++ .../workgroup/array.wgsl.expected.spvasm | 43 ++++ .../workgroup/array.wgsl.expected.wgsl | 6 + test/var/initialization/workgroup/matrix.wgsl | 6 + .../workgroup/matrix.wgsl.expected.hlsl | 16 ++ .../workgroup/matrix.wgsl.expected.msl | 192 ++++++++++++++++++ .../workgroup/matrix.wgsl.expected.spvasm | 42 ++++ .../workgroup/matrix.wgsl.expected.wgsl | 6 + test/var/initialization/workgroup/scalar.wgsl | 6 + .../workgroup/scalar.wgsl.expected.hlsl | 16 ++ .../workgroup/scalar.wgsl.expected.msl | 13 ++ .../workgroup/scalar.wgsl.expected.spvasm | 40 ++++ .../workgroup/scalar.wgsl.expected.wgsl | 6 + test/var/initialization/workgroup/struct.wgsl | 11 + .../workgroup/struct.wgsl.expected.hlsl | 22 ++ .../workgroup/struct.wgsl.expected.msl | 19 ++ .../workgroup/struct.wgsl.expected.spvasm | 47 +++++ .../workgroup/struct.wgsl.expected.wgsl | 11 + test/var/initialization/workgroup/vector.wgsl | 6 + .../workgroup/vector.wgsl.expected.hlsl | 16 ++ .../workgroup/vector.wgsl.expected.msl | 13 ++ .../workgroup/vector.wgsl.expected.spvasm | 41 ++++ .../workgroup/vector.wgsl.expected.wgsl | 6 + 75 files changed, 1198 insertions(+) create mode 100644 test/var/initialization/function/array.wgsl create mode 100644 test/var/initialization/function/array.wgsl.expected.hlsl create mode 100644 test/var/initialization/function/array.wgsl.expected.msl create mode 100644 test/var/initialization/function/array.wgsl.expected.spvasm create mode 100644 test/var/initialization/function/array.wgsl.expected.wgsl create mode 100644 test/var/initialization/function/matrix.wgsl create mode 100644 test/var/initialization/function/matrix.wgsl.expected.hlsl create mode 100644 test/var/initialization/function/matrix.wgsl.expected.msl create mode 100644 test/var/initialization/function/matrix.wgsl.expected.spvasm create mode 100644 test/var/initialization/function/matrix.wgsl.expected.wgsl create mode 100644 test/var/initialization/function/scalar.wgsl create mode 100644 test/var/initialization/function/scalar.wgsl.expected.hlsl create mode 100644 test/var/initialization/function/scalar.wgsl.expected.msl create mode 100644 test/var/initialization/function/scalar.wgsl.expected.spvasm create mode 100644 test/var/initialization/function/scalar.wgsl.expected.wgsl create mode 100644 test/var/initialization/function/struct.wgsl create mode 100644 test/var/initialization/function/struct.wgsl.expected.hlsl create mode 100644 test/var/initialization/function/struct.wgsl.expected.msl create mode 100644 test/var/initialization/function/struct.wgsl.expected.spvasm create mode 100644 test/var/initialization/function/struct.wgsl.expected.wgsl create mode 100644 test/var/initialization/function/vector.wgsl create mode 100644 test/var/initialization/function/vector.wgsl.expected.hlsl create mode 100644 test/var/initialization/function/vector.wgsl.expected.msl create mode 100644 test/var/initialization/function/vector.wgsl.expected.spvasm create mode 100644 test/var/initialization/function/vector.wgsl.expected.wgsl create mode 100644 test/var/initialization/private/array.wgsl create mode 100644 test/var/initialization/private/array.wgsl.expected.hlsl create mode 100644 test/var/initialization/private/array.wgsl.expected.msl create mode 100644 test/var/initialization/private/array.wgsl.expected.spvasm create mode 100644 test/var/initialization/private/array.wgsl.expected.wgsl create mode 100644 test/var/initialization/private/matrix.wgsl create mode 100644 test/var/initialization/private/matrix.wgsl.expected.hlsl create mode 100644 test/var/initialization/private/matrix.wgsl.expected.msl create mode 100644 test/var/initialization/private/matrix.wgsl.expected.spvasm create mode 100644 test/var/initialization/private/matrix.wgsl.expected.wgsl create mode 100644 test/var/initialization/private/scalar.wgsl create mode 100644 test/var/initialization/private/scalar.wgsl.expected.hlsl create mode 100644 test/var/initialization/private/scalar.wgsl.expected.msl create mode 100644 test/var/initialization/private/scalar.wgsl.expected.spvasm create mode 100644 test/var/initialization/private/scalar.wgsl.expected.wgsl create mode 100644 test/var/initialization/private/struct.wgsl create mode 100644 test/var/initialization/private/struct.wgsl.expected.hlsl create mode 100644 test/var/initialization/private/struct.wgsl.expected.msl create mode 100644 test/var/initialization/private/struct.wgsl.expected.spvasm create mode 100644 test/var/initialization/private/struct.wgsl.expected.wgsl create mode 100644 test/var/initialization/private/vector.wgsl create mode 100644 test/var/initialization/private/vector.wgsl.expected.hlsl create mode 100644 test/var/initialization/private/vector.wgsl.expected.msl create mode 100644 test/var/initialization/private/vector.wgsl.expected.spvasm create mode 100644 test/var/initialization/private/vector.wgsl.expected.wgsl create mode 100644 test/var/initialization/workgroup/array.wgsl create mode 100644 test/var/initialization/workgroup/array.wgsl.expected.hlsl create mode 100644 test/var/initialization/workgroup/array.wgsl.expected.msl create mode 100644 test/var/initialization/workgroup/array.wgsl.expected.spvasm create mode 100644 test/var/initialization/workgroup/array.wgsl.expected.wgsl create mode 100644 test/var/initialization/workgroup/matrix.wgsl create mode 100644 test/var/initialization/workgroup/matrix.wgsl.expected.hlsl create mode 100644 test/var/initialization/workgroup/matrix.wgsl.expected.msl create mode 100644 test/var/initialization/workgroup/matrix.wgsl.expected.spvasm create mode 100644 test/var/initialization/workgroup/matrix.wgsl.expected.wgsl create mode 100644 test/var/initialization/workgroup/scalar.wgsl create mode 100644 test/var/initialization/workgroup/scalar.wgsl.expected.hlsl create mode 100644 test/var/initialization/workgroup/scalar.wgsl.expected.msl create mode 100644 test/var/initialization/workgroup/scalar.wgsl.expected.spvasm create mode 100644 test/var/initialization/workgroup/scalar.wgsl.expected.wgsl create mode 100644 test/var/initialization/workgroup/struct.wgsl create mode 100644 test/var/initialization/workgroup/struct.wgsl.expected.hlsl create mode 100644 test/var/initialization/workgroup/struct.wgsl.expected.msl create mode 100644 test/var/initialization/workgroup/struct.wgsl.expected.spvasm create mode 100644 test/var/initialization/workgroup/struct.wgsl.expected.wgsl create mode 100644 test/var/initialization/workgroup/vector.wgsl create mode 100644 test/var/initialization/workgroup/vector.wgsl.expected.hlsl create mode 100644 test/var/initialization/workgroup/vector.wgsl.expected.msl create mode 100644 test/var/initialization/workgroup/vector.wgsl.expected.spvasm create mode 100644 test/var/initialization/workgroup/vector.wgsl.expected.wgsl diff --git a/test/var/initialization/function/array.wgsl b/test/var/initialization/function/array.wgsl new file mode 100644 index 0000000000..2fc36d991c --- /dev/null +++ b/test/var/initialization/function/array.wgsl @@ -0,0 +1,5 @@ +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : array; + 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 0000000000..ecb9e9f469 --- /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 0000000000..8ca89e829d --- /dev/null +++ b/test/var/initialization/function/array.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +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 0000000000..ced29ed95f --- /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 0000000000..3db0d56a33 --- /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; + ignore(v); +} diff --git a/test/var/initialization/function/matrix.wgsl b/test/var/initialization/function/matrix.wgsl new file mode 100644 index 0000000000..eb65b1e541 --- /dev/null +++ b/test/var/initialization/function/matrix.wgsl @@ -0,0 +1,5 @@ +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : mat2x3; + 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 0000000000..ec24c8bec8 --- /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 0000000000..951b502761 --- /dev/null +++ b/test/var/initialization/function/matrix.wgsl.expected.msl @@ -0,0 +1,9 @@ +#include + +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 0000000000..5b2ef08c7a --- /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 0000000000..17e072bb61 --- /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; + ignore(v); +} diff --git a/test/var/initialization/function/scalar.wgsl b/test/var/initialization/function/scalar.wgsl new file mode 100644 index 0000000000..81ccc32bb4 --- /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 0000000000..d96e258e64 --- /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 0000000000..d0221cad77 --- /dev/null +++ b/test/var/initialization/function/scalar.wgsl.expected.msl @@ -0,0 +1,9 @@ +#include + +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 0000000000..c67d83d5e1 --- /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 0000000000..73c1726602 --- /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 0000000000..35fd43d6ba --- /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 0000000000..2493fcb1cf --- /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 0000000000..ffec229457 --- /dev/null +++ b/test/var/initialization/function/struct.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +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 0000000000..9c865a6c7d --- /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 0000000000..57ec0e64cc --- /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 0000000000..24c52b456d --- /dev/null +++ b/test/var/initialization/function/vector.wgsl @@ -0,0 +1,5 @@ +[[stage(compute), workgroup_size(1)]] +fn main() { + var v : vec3; + 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 0000000000..f595b78661 --- /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 0000000000..c1b8f4d562 --- /dev/null +++ b/test/var/initialization/function/vector.wgsl.expected.msl @@ -0,0 +1,9 @@ +#include + +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 0000000000..aa8b4794c8 --- /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 0000000000..6fd0442797 --- /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; + ignore(v); +} diff --git a/test/var/initialization/private/array.wgsl b/test/var/initialization/private/array.wgsl new file mode 100644 index 0000000000..9dae8f7d16 --- /dev/null +++ b/test/var/initialization/private/array.wgsl @@ -0,0 +1,6 @@ +var v : array; + +[[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 0000000000..28733784fa --- /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 0000000000..ab06870491 --- /dev/null +++ b/test/var/initialization/private/array.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +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 0000000000..1996520984 --- /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 0000000000..3b54ea90a4 --- /dev/null +++ b/test/var/initialization/private/array.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var v : array; + +[[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 0000000000..b167840202 --- /dev/null +++ b/test/var/initialization/private/matrix.wgsl @@ -0,0 +1,6 @@ +var v : mat2x3; + +[[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 0000000000..df9907e6aa --- /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 0000000000..76f97114df --- /dev/null +++ b/test/var/initialization/private/matrix.wgsl.expected.msl @@ -0,0 +1,9 @@ +#include + +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 0000000000..04a8f59985 --- /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 0000000000..ce0680822e --- /dev/null +++ b/test/var/initialization/private/matrix.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var v : mat2x3; + +[[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 0000000000..117968603a --- /dev/null +++ b/test/var/initialization/private/scalar.wgsl @@ -0,0 +1,6 @@ +var 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 0000000000..0daf53c88a --- /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 0000000000..c96b5991d2 --- /dev/null +++ b/test/var/initialization/private/scalar.wgsl.expected.msl @@ -0,0 +1,9 @@ +#include + +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 0000000000..3deea8cf2b --- /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 0000000000..e634b37da0 --- /dev/null +++ b/test/var/initialization/private/scalar.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var 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 0000000000..0c64c99e5e --- /dev/null +++ b/test/var/initialization/private/struct.wgsl @@ -0,0 +1,11 @@ +struct S { + a : i32; + b : f32; +}; + +var 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 0000000000..ca8d6ca640 --- /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 0000000000..948d0b709a --- /dev/null +++ b/test/var/initialization/private/struct.wgsl.expected.msl @@ -0,0 +1,14 @@ +#include + +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 0000000000..ff2048adf6 --- /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 0000000000..f6777565de --- /dev/null +++ b/test/var/initialization/private/struct.wgsl.expected.wgsl @@ -0,0 +1,11 @@ +struct S { + a : i32; + b : f32; +}; + +var 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 0000000000..db1904e7e9 --- /dev/null +++ b/test/var/initialization/private/vector.wgsl @@ -0,0 +1,6 @@ +var v : vec3; + +[[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 0000000000..1f4a920388 --- /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 0000000000..6398a63d5e --- /dev/null +++ b/test/var/initialization/private/vector.wgsl.expected.msl @@ -0,0 +1,9 @@ +#include + +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 0000000000..f5086e9459 --- /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 0000000000..e5d11f7f80 --- /dev/null +++ b/test/var/initialization/private/vector.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var v : vec3; + +[[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 0000000000..9ba52fcfa3 --- /dev/null +++ b/test/var/initialization/workgroup/array.wgsl @@ -0,0 +1,6 @@ +var v : array; + +[[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 0000000000..e711cb67c5 --- /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 0000000000..f1774531cd --- /dev/null +++ b/test/var/initialization/workgroup/array.wgsl.expected.msl @@ -0,0 +1,18 @@ +#include + +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 0000000000..c1ef12c6f9 --- /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 0000000000..5cfadc1071 --- /dev/null +++ b/test/var/initialization/workgroup/array.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var v : array; + +[[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 0000000000..ddc128c140 --- /dev/null +++ b/test/var/initialization/workgroup/matrix.wgsl @@ -0,0 +1,6 @@ +var v : mat2x3; + +[[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 0000000000..1590817a2f --- /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 0000000000..52f49c4e19 --- /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') + 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... 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... 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... 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) 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> cols, _integer_sequence) 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 elems, _integer_sequence) 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) 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> cols, _integer_sequence) 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 elems, _integer_sequence) 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) 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> cols, _integer_sequence) 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 elems, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &that, _integer_sequence) 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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 &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 0000000000..7101d8074d --- /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 0000000000..6477351913 --- /dev/null +++ b/test/var/initialization/workgroup/matrix.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var v : mat2x3; + +[[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 0000000000..58f11480ac --- /dev/null +++ b/test/var/initialization/workgroup/scalar.wgsl @@ -0,0 +1,6 @@ +var 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 0000000000..8f52058e6e --- /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 0000000000..f63e941cf9 --- /dev/null +++ b/test/var/initialization/workgroup/scalar.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +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 0000000000..aede25ba32 --- /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 0000000000..e92577b7cf --- /dev/null +++ b/test/var/initialization/workgroup/scalar.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var 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 0000000000..81a88b2fdc --- /dev/null +++ b/test/var/initialization/workgroup/struct.wgsl @@ -0,0 +1,11 @@ +struct S { + a : i32; + b : f32; +}; + +var 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 0000000000..84ee4d6951 --- /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 0000000000..88dc937fe7 --- /dev/null +++ b/test/var/initialization/workgroup/struct.wgsl.expected.msl @@ -0,0 +1,19 @@ +#include + +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 0000000000..83f2f1b704 --- /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 0000000000..77244de177 --- /dev/null +++ b/test/var/initialization/workgroup/struct.wgsl.expected.wgsl @@ -0,0 +1,11 @@ +struct S { + a : i32; + b : f32; +}; + +var 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 0000000000..efad7b11b4 --- /dev/null +++ b/test/var/initialization/workgroup/vector.wgsl @@ -0,0 +1,6 @@ +var v : vec3; + +[[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 0000000000..bec61923a7 --- /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 0000000000..55cda8ec87 --- /dev/null +++ b/test/var/initialization/workgroup/vector.wgsl.expected.msl @@ -0,0 +1,13 @@ +#include + +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 0000000000..a5a02e2443 --- /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 0000000000..78f4426818 --- /dev/null +++ b/test/var/initialization/workgroup/vector.wgsl.expected.wgsl @@ -0,0 +1,6 @@ +var v : vec3; + +[[stage(compute), workgroup_size(1)]] +fn main() { + ignore(v); +}