writer/msl: Generate address spaces for pointers
Add more E2E tests to cover pointers with different storage classes.
Fixed: tint:815
Change-Id: I224a794cdf60648ce71dc9a0922d489542995be1
Reviewed-on: https://dawn-review.googlesource.com/c/tint/+/51404
Auto-Submit: James Price <jrprice@google.com>
Commit-Queue: Ben Clayton <bclayton@chromium.org>
Kokoro: Kokoro <noreply+kokoro@google.com>
Reviewed-by: Ben Clayton <bclayton@chromium.org>
diff --git a/test/ptr_ref/load/local/ptr_function.wgsl b/test/ptr_ref/load/local/ptr_function.wgsl
new file mode 100644
index 0000000..4c7f82c
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_function.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute)]]
+fn main() {
+ var i : i32 = 123;
+ let p : ptr<function, i32> = &i;
+ let use : i32 = *p + 1;
+}
diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl
new file mode 100644
index 0000000..b3db42f
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl
@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL
diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl
new file mode 100644
index 0000000..ce60e0c
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.msl
@@ -0,0 +1,10 @@
+#include <metal_stdlib>
+
+using namespace metal;
+kernel void tint_symbol() {
+ int i = 123;
+ thread int* const p = &(i);
+ int const use = (*(p) + 1);
+ return;
+}
+
diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_function.wgsl.expected.spvasm
new file mode 100644
index 0000000..3cb667b
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.spvasm
@@ -0,0 +1,26 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 15
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %main "main"
+ OpName %i "i"
+ %void = OpTypeVoid
+ %1 = OpTypeFunction %void
+ %int = OpTypeInt 32 1
+ %int_123 = OpConstant %int 123
+%_ptr_Function_int = OpTypePointer Function %int
+ %9 = OpConstantNull %int
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %1
+ %4 = OpLabel
+ %i = OpVariable %_ptr_Function_int Function %9
+ OpStore %i %int_123
+ %12 = OpLoad %int %i
+ %14 = OpIAdd %int %12 %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/ptr_ref/load/local/ptr_function.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_function.wgsl.expected.wgsl
new file mode 100644
index 0000000..be1837a
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_function.wgsl.expected.wgsl
@@ -0,0 +1,6 @@
+[[stage(compute)]]
+fn main() {
+ var i : i32 = 123;
+ let p : ptr<function, i32> = &(i);
+ let use : i32 = (*(p) + 1);
+}
diff --git a/test/ptr_ref/load/local/ptr_private.wgsl b/test/ptr_ref/load/local/ptr_private.wgsl
new file mode 100644
index 0000000..bd93db3
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_private.wgsl
@@ -0,0 +1,7 @@
+var<private> i : i32 = 123;
+
+[[stage(compute)]]
+fn main() {
+ let p : ptr<private, i32> = &i;
+ let use : i32 = *p + 1;
+}
diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl
new file mode 100644
index 0000000..b3db42f
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl
@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL
diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl
new file mode 100644
index 0000000..1b2d3db
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.msl
@@ -0,0 +1 @@
+SKIP: TINT_UNIMPLEMENTED crbug.com/tint/726: module-scope private and workgroup variables not yet implemented
diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_private.wgsl.expected.spvasm
new file mode 100644
index 0000000..7bf6bf9
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.spvasm
@@ -0,0 +1,24 @@
+; 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 %i "i"
+ OpName %main "main"
+ %int = OpTypeInt 32 1
+ %int_123 = OpConstant %int 123
+%_ptr_Private_int = OpTypePointer Private %int
+ %i = OpVariable %_ptr_Private_int Private %int_123
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %11 = OpLoad %int %i
+ %13 = OpIAdd %int %11 %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/ptr_ref/load/local/ptr_private.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_private.wgsl.expected.wgsl
new file mode 100644
index 0000000..a865378
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_private.wgsl.expected.wgsl
@@ -0,0 +1,7 @@
+var<private> i : i32 = 123;
+
+[[stage(compute)]]
+fn main() {
+ let p : ptr<private, i32> = &(i);
+ let use : i32 = (*(p) + 1);
+}
diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl b/test/ptr_ref/load/local/ptr_storage.wgsl
new file mode 100644
index 0000000..d334f34
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_storage.wgsl
@@ -0,0 +1,13 @@
+[[block]]
+struct S {
+ a : i32;
+};
+
+[[group(0), binding(0)]]
+var<storage> v : [[access(read_write)]] S;
+
+[[stage(compute)]]
+fn main() {
+ let p : ptr<storage, i32> = &v.a;
+ let use : i32 = *p + 1;
+}
diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl
new file mode 100644
index 0000000..c77b5fd
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl
@@ -0,0 +1 @@
+SKIP: error: cannot take the address of expression
diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl
new file mode 100644
index 0000000..17d0c1b
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ int a;
+};
+
+kernel void tint_symbol(device S& v [[buffer(0)]]) {
+ device int* const p = &(v.a);
+ int const use = (*(p) + 1);
+ return;
+}
+
diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.spvasm
new file mode 100644
index 0000000..e0e29bc
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpName %v "v"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %v DescriptorSet 0
+ OpDecorate %v Binding 0
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int
+%_ptr_StorageBuffer_S = OpTypePointer StorageBuffer %S
+ %v = OpVariable %_ptr_StorageBuffer_S StorageBuffer
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_StorageBuffer_int = OpTypePointer StorageBuffer %int
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %13 = OpAccessChain %_ptr_StorageBuffer_int %v %uint_0
+ %15 = OpLoad %int %13
+ %17 = OpIAdd %int %15 %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl
new file mode 100644
index 0000000..0c88535
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_storage.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+ a : i32;
+};
+
+[[group(0), binding(0)]] var<storage> v : [[access(read_write)]] S;
+
+[[stage(compute)]]
+fn main() {
+ let p : ptr<storage, i32> = &(v.a);
+ let use : i32 = (*(p) + 1);
+}
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl b/test/ptr_ref/load/local/ptr_uniform.wgsl
new file mode 100644
index 0000000..3dd38d4
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl
@@ -0,0 +1,13 @@
+[[block]]
+struct S {
+ a : i32;
+};
+
+[[group(0), binding(0)]]
+var<uniform> v : S;
+
+[[stage(compute)]]
+fn main() {
+ let p : ptr<uniform, i32> = &v.a;
+ let use : i32 = *p + 1;
+}
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
new file mode 100644
index 0000000..b3db42f
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl
@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl
new file mode 100644
index 0000000..9008b2f
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.msl
@@ -0,0 +1,13 @@
+#include <metal_stdlib>
+
+using namespace metal;
+struct S {
+ /* 0x0000 */ int a;
+};
+
+kernel void tint_symbol(constant S& v [[buffer(0)]]) {
+ constant int* const p = &(v.a);
+ int const use = (*(p) + 1);
+ return;
+}
+
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.spvasm
new file mode 100644
index 0000000..6810fb2
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.spvasm
@@ -0,0 +1,34 @@
+; SPIR-V
+; Version: 1.3
+; Generator: Google Tint Compiler; 0
+; Bound: 18
+; Schema: 0
+ OpCapability Shader
+ OpMemoryModel Logical GLSL450
+ OpEntryPoint GLCompute %main "main"
+ OpExecutionMode %main LocalSize 1 1 1
+ OpName %S "S"
+ OpMemberName %S 0 "a"
+ OpName %v "v"
+ OpName %main "main"
+ OpDecorate %S Block
+ OpMemberDecorate %S 0 Offset 0
+ OpDecorate %v DescriptorSet 0
+ OpDecorate %v Binding 0
+ %int = OpTypeInt 32 1
+ %S = OpTypeStruct %int
+%_ptr_Uniform_S = OpTypePointer Uniform %S
+ %v = OpVariable %_ptr_Uniform_S Uniform
+ %void = OpTypeVoid
+ %5 = OpTypeFunction %void
+ %uint = OpTypeInt 32 0
+ %uint_0 = OpConstant %uint 0
+%_ptr_Uniform_int = OpTypePointer Uniform %int
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %5
+ %8 = OpLabel
+ %13 = OpAccessChain %_ptr_Uniform_int %v %uint_0
+ %15 = OpLoad %int %13
+ %17 = OpIAdd %int %15 %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.wgsl
new file mode 100644
index 0000000..6cc9983
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_uniform.wgsl.expected.wgsl
@@ -0,0 +1,12 @@
+[[block]]
+struct S {
+ a : i32;
+};
+
+[[group(0), binding(0)]] var<uniform> v : S;
+
+[[stage(compute)]]
+fn main() {
+ let p : ptr<uniform, i32> = &(v.a);
+ let use : i32 = (*(p) + 1);
+}
diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl
new file mode 100644
index 0000000..3327d40
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl
@@ -0,0 +1,8 @@
+var<workgroup> i : i32;
+
+[[stage(compute)]]
+fn main() {
+ i = 123;
+ let p : ptr<workgroup, i32> = &i;
+ let use : i32 = *p + 1;
+}
diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl
new file mode 100644
index 0000000..b3db42f
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl
@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL
diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl
new file mode 100644
index 0000000..1b2d3db
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.msl
@@ -0,0 +1 @@
+SKIP: TINT_UNIMPLEMENTED crbug.com/tint/726: module-scope private and workgroup variables not yet implemented
diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm
new file mode 100644
index 0000000..1cda015
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.spvasm
@@ -0,0 +1,25 @@
+; 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 %i "i"
+ OpName %main "main"
+ %int = OpTypeInt 32 1
+%_ptr_Workgroup_int = OpTypePointer Workgroup %int
+ %i = OpVariable %_ptr_Workgroup_int Workgroup
+ %void = OpTypeVoid
+ %4 = OpTypeFunction %void
+ %int_123 = OpConstant %int 123
+ %int_1 = OpConstant %int 1
+ %main = OpFunction %void None %4
+ %7 = OpLabel
+ OpStore %i %int_123
+ %11 = OpLoad %int %i
+ %13 = OpIAdd %int %11 %int_1
+ OpReturn
+ OpFunctionEnd
diff --git a/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.wgsl b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.wgsl
new file mode 100644
index 0000000..3f9fb08
--- /dev/null
+++ b/test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.wgsl
@@ -0,0 +1,8 @@
+var<workgroup> i : i32;
+
+[[stage(compute)]]
+fn main() {
+ i = 123;
+ let p : ptr<workgroup, i32> = &(i);
+ let use : i32 = (*(p) + 1);
+}
