[tint] Add var tests for nested structs
Bug: tint:2142
Change-Id: I685e2345fac1edbb1bc27513fadc956f2e8830d1
Reviewed-on: https://dawn-review.googlesource.com/c/dawn/+/169980
Auto-Submit: Ben Clayton <bclayton@google.com>
Commit-Queue: dan sinclair <dsinclair@chromium.org>
Reviewed-by: dan sinclair <dsinclair@chromium.org>
diff --git a/test/tint/var/initialization/function/nested_structs.wgsl b/test/tint/var/initialization/function/nested_structs.wgsl
new file mode 100644
index 0000000..950318f
--- /dev/null
+++ b/test/tint/var/initialization/function/nested_structs.wgsl
@@ -0,0 +1,19 @@
+struct S1 {
+ i : i32,
+}
+struct S2 {
+ s1 : S1,
+}
+struct S3 {
+ s2 : S2,
+}
+
+fn f(s3 : S3) -> i32 { return s3.s2.s1.i; }
+
+@group(0) @binding(0) var<storage, read_write> out : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+ const C = 42;
+ out = f(S3(S2(S1(C))));
+}
diff --git a/test/tint/var/initialization/function/nested_structs.wgsl.expected.dxc.hlsl b/test/tint/var/initialization/function/nested_structs.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..268ee6d
--- /dev/null
+++ b/test/tint/var/initialization/function/nested_structs.wgsl.expected.dxc.hlsl
@@ -0,0 +1,22 @@
+struct S1 {
+ int i;
+};
+struct S2 {
+ S1 s1;
+};
+struct S3 {
+ S2 s2;
+};
+
+int f(S3 s3) {
+ return s3.s2.s1.i;
+}
+
+RWByteAddressBuffer tint_symbol : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ S3 tint_symbol_1 = {{{42}}};
+ tint_symbol.Store(0u, asuint(f(tint_symbol_1)));
+ return;
+}
diff --git a/test/tint/var/initialization/function/nested_structs.wgsl.expected.fxc.hlsl b/test/tint/var/initialization/function/nested_structs.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..268ee6d
--- /dev/null
+++ b/test/tint/var/initialization/function/nested_structs.wgsl.expected.fxc.hlsl
@@ -0,0 +1,22 @@
+struct S1 {
+ int i;
+};
+struct S2 {
+ S1 s1;
+};
+struct S3 {
+ S2 s2;
+};
+
+int f(S3 s3) {
+ return s3.s2.s1.i;
+}
+
+RWByteAddressBuffer tint_symbol : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ S3 tint_symbol_1 = {{{42}}};
+ tint_symbol.Store(0u, asuint(f(tint_symbol_1)));
+ return;
+}
diff --git a/test/tint/var/initialization/function/nested_structs.wgsl.expected.glsl b/test/tint/var/initialization/function/nested_structs.wgsl.expected.glsl
new file mode 100644
index 0000000..a8af172
--- /dev/null
+++ b/test/tint/var/initialization/function/nested_structs.wgsl.expected.glsl
@@ -0,0 +1,32 @@
+#version 310 es
+
+struct S1 {
+ int i;
+};
+
+struct S2 {
+ S1 s1;
+};
+
+struct S3 {
+ S2 s2;
+};
+
+int f(S3 s3) {
+ return s3.s2.s1.i;
+}
+
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ int inner;
+} tint_symbol;
+
+void tint_symbol_1() {
+ S3 tint_symbol_2 = S3(S2(S1(42)));
+ tint_symbol.inner = f(tint_symbol_2);
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_1();
+ return;
+}
diff --git a/test/tint/var/initialization/function/nested_structs.wgsl.expected.msl b/test/tint/var/initialization/function/nested_structs.wgsl.expected.msl
new file mode 100644
index 0000000..797a735
--- /dev/null
+++ b/test/tint/var/initialization/function/nested_structs.wgsl.expected.msl
@@ -0,0 +1,25 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S1 {
+ int i;
+};
+
+struct S2 {
+ S1 s1;
+};
+
+struct S3 {
+ S2 s2;
+};
+
+int f(S3 s3) {
+ return s3.s2.s1.i;
+}
+
+kernel void tint_symbol(device int* tint_symbol_2 [[buffer(0)]]) {
+ S3 const tint_symbol_1 = S3{.s2=S2{.s1=S1{.i=42}}};
+ *(tint_symbol_2) = f(tint_symbol_1);
+ return;
+}
+
diff --git a/test/tint/var/initialization/function/nested_structs.wgsl.expected.spvasm b/test/tint/var/initialization/function/nested_structs.wgsl.expected.spvasm
new file mode 100644
index 0000000..345c1b6
--- /dev/null
+++ b/test/tint/var/initialization/function/nested_structs.wgsl.expected.spvasm
@@ -0,0 +1,60 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 28
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %out_block "out_block"
+ OpMemberName %out_block 0 "inner"
+ OpName %out "out"
+ OpName %S3 "S3"
+ OpMemberName %S3 0 "s2"
+ OpName %S2 "S2"
+ OpMemberName %S2 0 "s1"
+ OpName %S1 "S1"
+ OpMemberName %S1 0 "i"
+ OpName %f "f"
+ OpName %s3 "s3"
+ OpName %main "main"
+ OpDecorate %out_block Block
+ OpMemberDecorate %out_block 0 Offset 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 0
+ OpMemberDecorate %S3 0 Offset 0
+ OpMemberDecorate %S2 0 Offset 0
+ OpMemberDecorate %S1 0 Offset 0
+ %int = OpTypeInt 32 1
+ %out_block = OpTypeStruct %int
+%_ptr_StorageBuffer_out_block = OpTypePointer StorageBuffer %out_block
+ %out = OpVariable %_ptr_StorageBuffer_out_block StorageBuffer
+ %S1 = OpTypeStruct %int
+ %S2 = OpTypeStruct %S1
+ %S3 = OpTypeStruct %S2
+ %5 = OpTypeFunction %int %S3
+ %void = OpTypeVoid
+ %15 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %int_42 = OpConstant %int 42
+ %25 = OpConstantComposite %S1 %int_42
+ %26 = OpConstantComposite %S2 %25
+ %27 = OpConstantComposite %S3 %26
+ %f = OpFunction %int None %5
+ %s3 = OpFunctionParameter %S3
+ %11 = OpLabel
+ %12 = OpCompositeExtract %S2 %s3 0
+ %13 = OpCompositeExtract %S1 %12 0
+ %14 = OpCompositeExtract %int %13 0
+ OpReturnValue %14
+ OpFunctionEnd
+ %main = OpFunction %void None %15
+ %18 = OpLabel
+ %22 = OpAccessChain %_ptr_StorageBuffer_int %out %uint_0
+ %23 = OpFunctionCall %int %f %27
+ OpStore %22 %23
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/var/initialization/function/nested_structs.wgsl.expected.wgsl b/test/tint/var/initialization/function/nested_structs.wgsl.expected.wgsl
new file mode 100644
index 0000000..1263c89
--- /dev/null
+++ b/test/tint/var/initialization/function/nested_structs.wgsl.expected.wgsl
@@ -0,0 +1,23 @@
+struct S1 {
+ i : i32,
+}
+
+struct S2 {
+ s1 : S1,
+}
+
+struct S3 {
+ s2 : S2,
+}
+
+fn f(s3 : S3) -> i32 {
+ return s3.s2.s1.i;
+}
+
+@group(0) @binding(0) var<storage, read_write> out : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+ const C = 42;
+ out = f(S3(S2(S1(C))));
+}
diff --git a/test/tint/var/initialization/private/nested_structs.wgsl b/test/tint/var/initialization/private/nested_structs.wgsl
new file mode 100644
index 0000000..f6be695
--- /dev/null
+++ b/test/tint/var/initialization/private/nested_structs.wgsl
@@ -0,0 +1,19 @@
+struct S1 {
+ i : i32,
+}
+struct S2 {
+ s1 : S1,
+}
+struct S3 {
+ s2 : S2,
+}
+
+const C = 42;
+var<private> P = S3(S2(S1(C)));
+
+@group(0) @binding(0) var<storage, read_write> out : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+ out = P.s2.s1.i;
+}
diff --git a/test/tint/var/initialization/private/nested_structs.wgsl.expected.dxc.hlsl b/test/tint/var/initialization/private/nested_structs.wgsl.expected.dxc.hlsl
new file mode 100644
index 0000000..4f76d41
--- /dev/null
+++ b/test/tint/var/initialization/private/nested_structs.wgsl.expected.dxc.hlsl
@@ -0,0 +1,21 @@
+struct S1 {
+ int i;
+};
+struct S2 {
+ S1 s1;
+};
+struct S3 {
+ S2 s2;
+};
+
+static const S1 c_2 = {42};
+static const S2 c_1 = {c_2};
+static const S3 c = {c_1};
+static S3 P = c;
+RWByteAddressBuffer tint_symbol : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol.Store(0u, asuint(P.s2.s1.i));
+ return;
+}
diff --git a/test/tint/var/initialization/private/nested_structs.wgsl.expected.fxc.hlsl b/test/tint/var/initialization/private/nested_structs.wgsl.expected.fxc.hlsl
new file mode 100644
index 0000000..4f76d41
--- /dev/null
+++ b/test/tint/var/initialization/private/nested_structs.wgsl.expected.fxc.hlsl
@@ -0,0 +1,21 @@
+struct S1 {
+ int i;
+};
+struct S2 {
+ S1 s1;
+};
+struct S3 {
+ S2 s2;
+};
+
+static const S1 c_2 = {42};
+static const S2 c_1 = {c_2};
+static const S3 c = {c_1};
+static S3 P = c;
+RWByteAddressBuffer tint_symbol : register(u0);
+
+[numthreads(1, 1, 1)]
+void main() {
+ tint_symbol.Store(0u, asuint(P.s2.s1.i));
+ return;
+}
diff --git a/test/tint/var/initialization/private/nested_structs.wgsl.expected.glsl b/test/tint/var/initialization/private/nested_structs.wgsl.expected.glsl
new file mode 100644
index 0000000..e6f59d6
--- /dev/null
+++ b/test/tint/var/initialization/private/nested_structs.wgsl.expected.glsl
@@ -0,0 +1,28 @@
+#version 310 es
+
+struct S1 {
+ int i;
+};
+
+struct S2 {
+ S1 s1;
+};
+
+struct S3 {
+ S2 s2;
+};
+
+S3 P = S3(S2(S1(42)));
+layout(binding = 0, std430) buffer tint_symbol_block_ssbo {
+ int inner;
+} tint_symbol;
+
+void tint_symbol_1() {
+ tint_symbol.inner = P.s2.s1.i;
+}
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+void main() {
+ tint_symbol_1();
+ return;
+}
diff --git a/test/tint/var/initialization/private/nested_structs.wgsl.expected.msl b/test/tint/var/initialization/private/nested_structs.wgsl.expected.msl
new file mode 100644
index 0000000..a4c8ea3
--- /dev/null
+++ b/test/tint/var/initialization/private/nested_structs.wgsl.expected.msl
@@ -0,0 +1,26 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S1 {
+ int i;
+};
+
+struct S2 {
+ S1 s1;
+};
+
+struct S3 {
+ S2 s2;
+};
+
+struct tint_private_vars_struct {
+ S3 P;
+};
+
+kernel void tint_symbol(device int* tint_symbol_1 [[buffer(0)]]) {
+ thread tint_private_vars_struct tint_private_vars = {};
+ tint_private_vars.P = S3{.s2=S2{.s1=S1{.i=42}}};
+ *(tint_symbol_1) = tint_private_vars.P.s2.s1.i;
+ return;
+}
+
diff --git a/test/tint/var/initialization/private/nested_structs.wgsl.expected.spvasm b/test/tint/var/initialization/private/nested_structs.wgsl.expected.spvasm
new file mode 100644
index 0000000..7469455
--- /dev/null
+++ b/test/tint/var/initialization/private/nested_structs.wgsl.expected.spvasm
@@ -0,0 +1,54 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 25
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S3 "S3"
+ OpMemberName %S3 0 "s2"
+ OpName %S2 "S2"
+ OpMemberName %S2 0 "s1"
+ OpName %S1 "S1"
+ OpMemberName %S1 0 "i"
+ OpName %P "P"
+ OpName %out_block "out_block"
+ OpMemberName %out_block 0 "inner"
+ OpName %out "out"
+ OpName %main "main"
+ OpMemberDecorate %S3 0 Offset 0
+ OpMemberDecorate %S2 0 Offset 0
+ OpMemberDecorate %S1 0 Offset 0
+ OpDecorate %out_block Block
+ OpMemberDecorate %out_block 0 Offset 0
+ OpDecorate %out DescriptorSet 0
+ OpDecorate %out Binding 0
+ %int = OpTypeInt 32 1
+ %S1 = OpTypeStruct %int
+ %S2 = OpTypeStruct %S1
+ %S3 = OpTypeStruct %S2
+ %int_42 = OpConstant %int 42
+ %6 = OpConstantComposite %S1 %int_42
+ %7 = OpConstantComposite %S2 %6
+ %8 = OpConstantComposite %S3 %7
+%_ptr_Private_S3 = OpTypePointer Private %S3
+ %P = OpVariable %_ptr_Private_S3 Private %8
+ %out_block = OpTypeStruct %int
+%_ptr_StorageBuffer_out_block = OpTypePointer StorageBuffer %out_block
+ %out = OpVariable %_ptr_StorageBuffer_out_block StorageBuffer
+ %void = OpTypeVoid
+ %14 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+%_ptr_Private_int = OpTypePointer Private %int
+ %main = OpFunction %void None %14
+ %17 = OpLabel
+ %21 = OpAccessChain %_ptr_StorageBuffer_int %out %uint_0
+ %23 = OpAccessChain %_ptr_Private_int %P %uint_0 %uint_0 %uint_0
+ %24 = OpLoad %int %23
+ OpStore %21 %24
+ OpReturn
+ OpFunctionEnd
diff --git a/test/tint/var/initialization/private/nested_structs.wgsl.expected.wgsl b/test/tint/var/initialization/private/nested_structs.wgsl.expected.wgsl
new file mode 100644
index 0000000..afc3d49
--- /dev/null
+++ b/test/tint/var/initialization/private/nested_structs.wgsl.expected.wgsl
@@ -0,0 +1,22 @@
+struct S1 {
+ i : i32,
+}
+
+struct S2 {
+ s1 : S1,
+}
+
+struct S3 {
+ s2 : S2,
+}
+
+const C = 42;
+
+var<private> P = S3(S2(S1(C)));
+
+@group(0) @binding(0) var<storage, read_write> out : i32;
+
+@compute @workgroup_size(1)
+fn main() {
+ out = P.s2.s1.i;
+}
