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>

src/writer/msl/generator_impl.cc

@@ -1913,7 +1913,23 @@ bool GeneratorImpl::EmitType(const sem::Type* type, const std::string& name) {
     }
     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;
     }

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;
 }

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;
 }

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);
+  thread uint* const x_1 = &(x_10);
-  uint* const x_2 = x_1;
+  thread uint* const x_2 = x_1;
   return;
 }
 

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;
+}

test/ptr_ref/load/local/ptr_function.wgsl.expected.hlsl

@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL

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;
+}
+

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

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);
+}

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;
+}

test/ptr_ref/load/local/ptr_private.wgsl.expected.hlsl

@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL

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

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

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);
+}

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;
+}

test/ptr_ref/load/local/ptr_storage.wgsl.expected.hlsl

@@ -0,0 +1 @@
+SKIP: error: cannot take the address of expression

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;
+}
+

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

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);
+}

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;
+}

test/ptr_ref/load/local/ptr_uniform.wgsl.expected.hlsl

@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL

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;
+}
+

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

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);
+}

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;
+}

test/ptr_ref/load/local/ptr_workgroup.wgsl.expected.hlsl

@@ -0,0 +1 @@
+SKIP: Failed to generate: error: pointers not supported in HLSL

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

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

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);
+}

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);
 }

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));
 }
 

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;

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;
 }

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;
 }