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/src/writer/msl/generator_impl.cc b/src/writer/msl/generator_impl.cc
index c4d25cf..ee976d0 100644
--- a/src/writer/msl/generator_impl.cc
+++ b/src/writer/msl/generator_impl.cc
@@ -1913,7 +1913,23 @@
}
out_ << mat->columns() << "x" << mat->rows();
} else if (auto* ptr = type->As<sem::Pointer>()) {
- // TODO(dsinclair): Storage class?
+ switch (ptr->StorageClass()) {
+ case ast::StorageClass::kFunction:
+ case ast::StorageClass::kPrivate:
+ out_ << "thread ";
+ break;
+ case ast::StorageClass::kWorkgroup:
+ out_ << "threadgroup ";
+ break;
+ case ast::StorageClass::kStorage:
+ out_ << "device ";
+ break;
+ case ast::StorageClass::kUniform:
+ out_ << "constant ";
+ break;
+ default:
+ TINT_ICE(diagnostics_) << "unhandled storage class for pointer";
+ }
if (!EmitType(ptr->StoreType(), "")) {
return false;
}
diff --git a/test/ptr_ref/access/matrix.wgsl.expected.msl b/test/ptr_ref/access/matrix.wgsl.expected.msl
index 5bf9bb9..1a95911 100644
--- a/test/ptr_ref/access/matrix.wgsl.expected.msl
+++ b/test/ptr_ref/access/matrix.wgsl.expected.msl
@@ -3,7 +3,7 @@
using namespace metal;
kernel void tint_symbol() {
float3x3 m = float3x3(float3(1.0f, 2.0f, 3.0f), float3(4.0f, 5.0f, 6.0f), float3(7.0f, 8.0f, 9.0f));
- float3* const v = &(m[1]);
+ thread float3* const v = &(m[1]);
*(v) = float3(5.0f, 5.0f, 5.0f);
return;
}
diff --git a/test/ptr_ref/access/vector.wgsl.expected.msl b/test/ptr_ref/access/vector.wgsl.expected.msl
index 88a1b99..fc61851 100644
--- a/test/ptr_ref/access/vector.wgsl.expected.msl
+++ b/test/ptr_ref/access/vector.wgsl.expected.msl
@@ -3,7 +3,7 @@
using namespace metal;
kernel void tint_symbol() {
float3 v = float3(1.0f, 2.0f, 3.0f);
- float* const f = &(v.y);
+ thread float* const f = &(v.y);
*(f) = 5.0f;
return;
}
diff --git a/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl b/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl
index ec58031..6577255 100644
--- a/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl
+++ b/test/ptr_ref/copy/ptr_copy.spvasm.expected.msl
@@ -3,8 +3,8 @@
using namespace metal;
kernel void tint_symbol() {
uint x_10 = 0u;
- uint* const x_1 = &(x_10);
- uint* const x_2 = x_1;
+ thread uint* const x_1 = &(x_10);
+ thread uint* const x_2 = x_1;
return;
}
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);
+}
diff --git a/test/ptr_ref/load/param/ptr.spvasm.expected.msl b/test/ptr_ref/load/param/ptr.spvasm.expected.msl
index cbe0a55..cf1067f 100644
--- a/test/ptr_ref/load/param/ptr.spvasm.expected.msl
+++ b/test/ptr_ref/load/param/ptr.spvasm.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-int func(int value, int* pointer) {
+int func(int value, thread int* pointer) {
int const x_9 = *(pointer);
return (value + x_9);
}
diff --git a/test/ptr_ref/load/param/ptr.wgsl.expected.msl b/test/ptr_ref/load/param/ptr.wgsl.expected.msl
index 48556e1..20a0cc9 100644
--- a/test/ptr_ref/load/param/ptr.wgsl.expected.msl
+++ b/test/ptr_ref/load/param/ptr.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-int func(int value, int* pointer) {
+int func(int value, thread int* pointer) {
return (value + *(pointer));
}
diff --git a/test/ptr_ref/store/local/i32.wgsl.expected.msl b/test/ptr_ref/store/local/i32.wgsl.expected.msl
index 6674f82..3c243ed 100644
--- a/test/ptr_ref/store/local/i32.wgsl.expected.msl
+++ b/test/ptr_ref/store/local/i32.wgsl.expected.msl
@@ -3,7 +3,7 @@
using namespace metal;
kernel void tint_symbol() {
int i = 123;
- int* const p = &(i);
+ thread int* const p = &(i);
*(p) = 123;
*(p) = ((100 + 20) + 3);
return;
diff --git a/test/ptr_ref/store/param/ptr.spvasm.expected.msl b/test/ptr_ref/store/param/ptr.spvasm.expected.msl
index 5d459da..0633ba6 100644
--- a/test/ptr_ref/store/param/ptr.spvasm.expected.msl
+++ b/test/ptr_ref/store/param/ptr.spvasm.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-void func(int value, int* pointer) {
+void func(int value, thread int* pointer) {
*(pointer) = value;
return;
}
diff --git a/test/ptr_ref/store/param/ptr.wgsl.expected.msl b/test/ptr_ref/store/param/ptr.wgsl.expected.msl
index 517183a..9a64bcb 100644
--- a/test/ptr_ref/store/param/ptr.wgsl.expected.msl
+++ b/test/ptr_ref/store/param/ptr.wgsl.expected.msl
@@ -1,7 +1,7 @@
#include <metal_stdlib>
using namespace metal;
-void func(int value, int* pointer) {
+void func(int value, thread int* pointer) {
*(pointer) = value;
}
