[msl-out] fix pointers to private or workgroup address spaces possibly being read only

src/back/msl/writer.rs

@@ -462,7 +462,9 @@ impl crate::AddressSpace {
             // rely on the actual use of a global by functions. This means we
             // may end up with "const" even if the binding is read-write,
             // and that should be OK.
-            Self::Storage { .. } | Self::Private | Self::WorkGroup => true,
+            Self::Storage { .. } => true,
+            // These should always be read-write.
+            Self::Private | Self::WorkGroup => false,
             // These translate to `constant` address space, no need for qualifiers.
             Self::Uniform | Self::PushConstant => false,
             // Not applicable.

tests/in/access.wgsl

@@ -120,3 +120,14 @@ fn atomics() {
 	// tmp = atomicCompareExchangeWeak(&bar.atom, 5, 5);
 	atomicStore(&bar.atom, value);
 }
+
+var<workgroup> val: u32;
+
+fn assign_through_ptr_fn(p: ptr<workgroup, u32>) {
+    *p = 42u;
+}
+
+@compute @workgroup_size(1)
+fn assign_through_ptr() {
+    assign_through_ptr_fn(&val);
+}

tests/out/glsl/access.assign_through_ptr.Compute.glsl

@@ -0,0 +1,35 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct AlignedWrapper {
+    int value;
+};
+struct Baz {
+    mat3x2 m;
+};
+shared uint val;
+
+
+float read_from_private(inout float foo_1) {
+    float _e4 = foo_1;
+    return _e4;
+}
+
+float test_arr_as_arg(float a[5][10]) {
+    return a[4][9];
+}
+
+void assign_through_ptr_fn(inout uint p) {
+    p = 42u;
+    return;
+}
+
+void main() {
+    assign_through_ptr_fn(val);
+    return;
+}
+

tests/out/glsl/access.atomics.Compute.glsl

@@ -29,6 +29,11 @@ float test_arr_as_arg(float a[5][10]) {
     return a[4][9];
 }
 
+void assign_through_ptr_fn(inout uint p) {
+    p = 42u;
+    return;
+}
+
 void main() {
     int tmp = 0;
     int value = _group_0_binding_0_cs.atom;

tests/out/glsl/access.foo_frag.Fragment.glsl

@@ -30,6 +30,11 @@ float test_arr_as_arg(float a[5][10]) {
     return a[4][9];
 }
 
+void assign_through_ptr_fn(inout uint p) {
+    p = 42u;
+    return;
+}
+
 void main() {
     _group_0_binding_0_fs._matrix[1][2] = 1.0;
     _group_0_binding_0_fs._matrix = mat4x3(vec3(0.0), vec3(1.0), vec3(2.0), vec3(3.0));

tests/out/glsl/access.foo_vert.Vertex.glsl

@@ -66,6 +66,11 @@ float test_arr_as_arg(float a[5][10]) {
     return a[4][9];
 }
 
+void assign_through_ptr_fn(inout uint p) {
+    p = 42u;
+    return;
+}
+
 void main() {
     uint vi = uint(gl_VertexID);
     float foo = 0.0;

tests/out/hlsl/access.hlsl

@@ -21,6 +21,7 @@ float Constructarray5_array10_float__(float arg0[10], float arg1[10], float arg2
 RWByteAddressBuffer bar : register(u0);
 cbuffer baz : register(b1) { Baz baz; }
 RWByteAddressBuffer qux : register(u2);
+groupshared uint val;
 
 float3x2 GetMatmOnBaz(Baz obj) {
     return float3x2(obj.m_0, obj.m_1, obj.m_2);
@@ -104,6 +105,12 @@ float test_arr_as_arg(float a[5][10])
     return a[4][9];
 }
 
+void assign_through_ptr_fn(inout uint p)
+{
+    p = 42u;
+    return;
+}
+
 uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
 {
     uint ret;
@@ -190,3 +197,10 @@ void atomics()
     bar.Store(96, asuint(value_1));
     return;
 }
+
+[numthreads(1, 1, 1)]
+void assign_through_ptr()
+{
+    assign_through_ptr_fn(val);
+    return;
+}

tests/out/hlsl/access.hlsl.config

@@ -1,3 +1,3 @@
 vertex=(foo_vert:vs_5_1 )
 fragment=(foo_frag:ps_5_1 )
-compute=(atomics:cs_5_1 )
+compute=(atomics:cs_5_1 assign_through_ptr:cs_5_1 )

tests/out/msl/access.msl

@@ -92,6 +92,13 @@ float test_arr_as_arg(
     return a.inner[4].inner[9];
 }
 
+void assign_through_ptr_fn(
+    threadgroup uint& p
+) {
+    p = 42u;
+    return;
+}
+
 struct foo_vertInput {
 };
 struct foo_vertOutput {
@@ -165,3 +172,11 @@ kernel void atomics(
     metal::atomic_store_explicit(&bar.atom, value_1, metal::memory_order_relaxed);
     return;
 }
+
+
+kernel void assign_through_ptr(
+  threadgroup uint& val
+) {
+    assign_through_ptr_fn(val);
+    return;
+}

tests/out/msl/policy-mix.msl

@@ -39,8 +39,8 @@ metal::float4 mock_function(
     device InStorage const& in_storage,
     constant InUniform& in_uniform,
     metal::texture2d_array<float, metal::access::sample> image_2d_array,
-    threadgroup type_5 const& in_workgroup,
-    thread type_6 const& in_private
+    threadgroup type_5& in_workgroup,
+    thread type_6& in_private
 ) {
     type_9 in_function = {};
     for(int _i=0; _i<2; ++_i) in_function.inner[_i] = type_9 {metal::float4(0.7070000171661377, 0.0, 0.0, 1.0), metal::float4(0.0, 0.7070000171661377, 0.0, 1.0)}.inner[_i];

tests/out/msl/quad-vert.msl

@@ -23,9 +23,9 @@ constant gl_PerVertex const_gl_PerVertex = {const_type_3_, 1.0, const_type_5_, c
 
 void main_1(
     thread metal::float2& v_uv,
-    thread metal::float2 const& a_uv_1,
+    thread metal::float2& a_uv_1,
     thread gl_PerVertex& perVertexStruct,
-    thread metal::float2 const& a_pos_1
+    thread metal::float2& a_pos_1
 ) {
     metal::float2 _e12 = a_uv_1;
     v_uv = _e12;

tests/out/spv/access.spvasm

@@ -1,77 +1,83 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 230
+; Bound: 242
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Vertex %149 "foo_vert" %144 %147
-OpEntryPoint Fragment %188 "foo_frag" %187
-OpEntryPoint GLCompute %207 "atomics"
-OpExecutionMode %188 OriginUpperLeft
-OpExecutionMode %207 LocalSize 1 1 1
+OpEntryPoint Vertex %157 "foo_vert" %152 %155
+OpEntryPoint Fragment %196 "foo_frag" %195
+OpEntryPoint GLCompute %215 "atomics"
+OpEntryPoint GLCompute %239 "assign_through_ptr"
+OpExecutionMode %196 OriginUpperLeft
+OpExecutionMode %215 LocalSize 1 1 1
+OpExecutionMode %239 LocalSize 1 1 1
 OpSource GLSL 450
-OpMemberName %32 0 "value"
-OpName %32 "AlignedWrapper"
-OpMemberName %41 0 "_matrix"
-OpMemberName %41 1 "matrix_array"
-OpMemberName %41 2 "atom"
-OpMemberName %41 3 "arr"
-OpMemberName %41 4 "data"
-OpName %41 "Bar"
-OpMemberName %43 0 "m"
-OpName %43 "Baz"
-OpName %54 "bar"
-OpName %56 "baz"
-OpName %59 "qux"
-OpName %62 "idx"
-OpName %64 "t"
-OpName %68 "test_matrix_within_struct_accesses"
-OpName %127 "foo"
-OpName %128 "read_from_private"
-OpName %133 "a"
-OpName %134 "test_arr_as_arg"
-OpName %139 "foo"
-OpName %140 "c"
-OpName %144 "vi"
-OpName %149 "foo_vert"
-OpName %188 "foo_frag"
-OpName %204 "tmp"
-OpName %207 "atomics"
-OpMemberDecorate %32 0 Offset 0
-OpDecorate %37 ArrayStride 16
-OpDecorate %39 ArrayStride 8
+OpMemberName %33 0 "value"
+OpName %33 "AlignedWrapper"
+OpMemberName %42 0 "_matrix"
+OpMemberName %42 1 "matrix_array"
+OpMemberName %42 2 "atom"
+OpMemberName %42 3 "arr"
+OpMemberName %42 4 "data"
+OpName %42 "Bar"
+OpMemberName %44 0 "m"
+OpName %44 "Baz"
+OpName %56 "bar"
+OpName %58 "baz"
+OpName %61 "qux"
+OpName %64 "val"
+OpName %65 "idx"
+OpName %67 "t"
+OpName %71 "test_matrix_within_struct_accesses"
+OpName %130 "foo"
+OpName %131 "read_from_private"
+OpName %136 "a"
+OpName %137 "test_arr_as_arg"
+OpName %143 "p"
+OpName %144 "assign_through_ptr_fn"
+OpName %147 "foo"
+OpName %148 "c"
+OpName %152 "vi"
+OpName %157 "foo_vert"
+OpName %196 "foo_frag"
+OpName %212 "tmp"
+OpName %215 "atomics"
+OpName %239 "assign_through_ptr"
+OpMemberDecorate %33 0 Offset 0
+OpDecorate %38 ArrayStride 16
 OpDecorate %40 ArrayStride 8
-OpMemberDecorate %41 0 Offset 0
-OpMemberDecorate %41 0 ColMajor
-OpMemberDecorate %41 0 MatrixStride 16
-OpMemberDecorate %41 1 Offset 64
-OpMemberDecorate %41 1 ColMajor
-OpMemberDecorate %41 1 MatrixStride 8
-OpMemberDecorate %41 2 Offset 96
-OpMemberDecorate %41 3 Offset 104
-OpMemberDecorate %41 4 Offset 120
-OpMemberDecorate %43 0 Offset 0
-OpMemberDecorate %43 0 ColMajor
-OpMemberDecorate %43 0 MatrixStride 8
-OpDecorate %46 ArrayStride 4
-OpDecorate %47 ArrayStride 40
-OpDecorate %50 ArrayStride 4
-OpDecorate %54 DescriptorSet 0
-OpDecorate %54 Binding 0
-OpDecorate %41 Block
+OpDecorate %41 ArrayStride 8
+OpMemberDecorate %42 0 Offset 0
+OpMemberDecorate %42 0 ColMajor
+OpMemberDecorate %42 0 MatrixStride 16
+OpMemberDecorate %42 1 Offset 64
+OpMemberDecorate %42 1 ColMajor
+OpMemberDecorate %42 1 MatrixStride 8
+OpMemberDecorate %42 2 Offset 96
+OpMemberDecorate %42 3 Offset 104
+OpMemberDecorate %42 4 Offset 120
+OpMemberDecorate %44 0 Offset 0
+OpMemberDecorate %44 0 ColMajor
+OpMemberDecorate %44 0 MatrixStride 8
+OpDecorate %47 ArrayStride 4
+OpDecorate %48 ArrayStride 40
+OpDecorate %51 ArrayStride 4
 OpDecorate %56 DescriptorSet 0
-OpDecorate %56 Binding 1
-OpDecorate %57 Block
-OpMemberDecorate %57 0 Offset 0
-OpDecorate %59 DescriptorSet 0
-OpDecorate %59 Binding 2
-OpDecorate %60 Block
-OpMemberDecorate %60 0 Offset 0
-OpDecorate %144 BuiltIn VertexIndex
-OpDecorate %147 BuiltIn Position
-OpDecorate %187 Location 0
+OpDecorate %56 Binding 0
+OpDecorate %42 Block
+OpDecorate %58 DescriptorSet 0
+OpDecorate %58 Binding 1
+OpDecorate %59 Block
+OpMemberDecorate %59 0 Offset 0
+OpDecorate %61 DescriptorSet 0
+OpDecorate %61 Binding 2
+OpDecorate %62 Block
+OpMemberDecorate %62 0 Offset 0
+OpDecorate %152 BuiltIn VertexIndex
+OpDecorate %155 BuiltIn Position
+OpDecorate %195 Location 0
 %2 = OpTypeVoid
 %4 = OpTypeInt 32 1
 %3 = OpConstant  %4  2
@@ -102,248 +108,267 @@ OpDecorate %187 Location 0
 %29 = OpConstant  %26  1
 %30 = OpConstant  %4  42
 %31 = OpConstant  %26  0
-%32 = OpTypeStruct %4
-%34 = OpTypeVector %9 3
-%33 = OpTypeMatrix %34 4
-%36 = OpTypeVector %9 2
-%35 = OpTypeMatrix %36 2
-%37 = OpTypeArray %35 %3
-%38 = OpTypeVector %26 2
-%39 = OpTypeArray %38 %3
-%40 = OpTypeRuntimeArray %32
-%41 = OpTypeStruct %33 %37 %4 %39 %40
-%42 = OpTypeMatrix %36 3
-%43 = OpTypeStruct %42
-%44 = OpTypeVector %4 2
-%45 = OpTypePointer Function %9
-%46 = OpTypeArray %9 %21
-%47 = OpTypeArray %46 %22
-%48 = OpTypeVector %9 4
-%49 = OpTypePointer StorageBuffer %4
-%50 = OpTypeArray %4 %22
-%51 = OpConstantComposite  %46  %24 %24 %24 %24 %24 %24 %24 %24 %24 %24
-%52 = OpConstantComposite  %47  %51 %51 %51 %51 %51
-%53 = OpConstantComposite  %44  %7 %7
-%55 = OpTypePointer StorageBuffer %41
-%54 = OpVariable  %55  StorageBuffer
-%57 = OpTypeStruct %43
-%58 = OpTypePointer Uniform %57
-%56 = OpVariable  %58  Uniform
-%60 = OpTypeStruct %44
-%61 = OpTypePointer StorageBuffer %60
-%59 = OpVariable  %61  StorageBuffer
-%63 = OpTypePointer Function %4
-%65 = OpTypePointer Function %43
-%66 = OpConstantNull  %43
-%69 = OpTypeFunction %2
-%70 = OpTypePointer Uniform %43
-%72 = OpTypePointer StorageBuffer %44
-%76 = OpTypePointer Uniform %42
-%79 = OpTypePointer Uniform %36
-%85 = OpTypePointer Uniform %9
-%105 = OpTypePointer Function %42
-%111 = OpTypePointer Function %36
-%117 = OpTypePointer Function %9
-%129 = OpTypeFunction %9 %45
-%135 = OpTypeFunction %9 %47
-%141 = OpTypePointer Function %50
-%142 = OpConstantNull  %50
-%145 = OpTypePointer Input %26
-%144 = OpVariable  %145  Input
-%148 = OpTypePointer Output %48
-%147 = OpVariable  %148  Output
-%155 = OpTypePointer StorageBuffer %33
-%158 = OpTypePointer StorageBuffer %39
-%161 = OpTypePointer StorageBuffer %34
-%162 = OpTypePointer StorageBuffer %9
-%165 = OpTypePointer StorageBuffer %40
-%168 = OpTypePointer StorageBuffer %32
-%169 = OpConstant  %26  4
-%181 = OpTypeVector %4 4
-%187 = OpVariable  %148  Output
-%205 = OpConstantNull  %4
-%209 = OpTypePointer StorageBuffer %4
-%212 = OpConstant  %26  64
-%68 = OpFunction  %2  None %69
-%67 = OpLabel
-%62 = OpVariable  %63  Function %5
-%64 = OpVariable  %65  Function %66
-%71 = OpAccessChain  %70  %56 %31
-OpBranch %73
-%73 = OpLabel
-%74 = OpLoad  %4  %62
-%75 = OpISub  %4  %74 %6
-OpStore %62 %75
-%77 = OpAccessChain  %76  %71 %31
-%78 = OpLoad  %42  %77
-%80 = OpAccessChain  %79  %71 %31 %31
-%81 = OpLoad  %36  %80
-%82 = OpLoad  %4  %62
-%83 = OpAccessChain  %79  %71 %31 %82
-%84 = OpLoad  %36  %83
-%86 = OpAccessChain  %85  %71 %31 %31 %29
-%87 = OpLoad  %9  %86
-%88 = OpLoad  %4  %62
-%89 = OpAccessChain  %85  %71 %31 %31 %88
+%32 = OpConstant  %26  42
+%33 = OpTypeStruct %4
+%35 = OpTypeVector %9 3
+%34 = OpTypeMatrix %35 4
+%37 = OpTypeVector %9 2
+%36 = OpTypeMatrix %37 2
+%38 = OpTypeArray %36 %3
+%39 = OpTypeVector %26 2
+%40 = OpTypeArray %39 %3
+%41 = OpTypeRuntimeArray %33
+%42 = OpTypeStruct %34 %38 %4 %40 %41
+%43 = OpTypeMatrix %37 3
+%44 = OpTypeStruct %43
+%45 = OpTypeVector %4 2
+%46 = OpTypePointer Function %9
+%47 = OpTypeArray %9 %21
+%48 = OpTypeArray %47 %22
+%49 = OpTypeVector %9 4
+%50 = OpTypePointer StorageBuffer %4
+%51 = OpTypeArray %4 %22
+%52 = OpTypePointer Workgroup %26
+%53 = OpConstantComposite  %47  %24 %24 %24 %24 %24 %24 %24 %24 %24 %24
+%54 = OpConstantComposite  %48  %53 %53 %53 %53 %53
+%55 = OpConstantComposite  %45  %7 %7
+%57 = OpTypePointer StorageBuffer %42
+%56 = OpVariable  %57  StorageBuffer
+%59 = OpTypeStruct %44
+%60 = OpTypePointer Uniform %59
+%58 = OpVariable  %60  Uniform
+%62 = OpTypeStruct %45
+%63 = OpTypePointer StorageBuffer %62
+%61 = OpVariable  %63  StorageBuffer
+%64 = OpVariable  %52  Workgroup
+%66 = OpTypePointer Function %4
+%68 = OpTypePointer Function %44
+%69 = OpConstantNull  %44
+%72 = OpTypeFunction %2
+%73 = OpTypePointer Uniform %44
+%75 = OpTypePointer StorageBuffer %45
+%79 = OpTypePointer Uniform %43
+%82 = OpTypePointer Uniform %37
+%88 = OpTypePointer Uniform %9
+%108 = OpTypePointer Function %43
+%114 = OpTypePointer Function %37
+%120 = OpTypePointer Function %9
+%132 = OpTypeFunction %9 %46
+%138 = OpTypeFunction %9 %48
+%145 = OpTypeFunction %2 %52
+%149 = OpTypePointer Function %51
+%150 = OpConstantNull  %51
+%153 = OpTypePointer Input %26
+%152 = OpVariable  %153  Input
+%156 = OpTypePointer Output %49
+%155 = OpVariable  %156  Output
+%163 = OpTypePointer StorageBuffer %34
+%166 = OpTypePointer StorageBuffer %40
+%169 = OpTypePointer StorageBuffer %35
+%170 = OpTypePointer StorageBuffer %9
+%173 = OpTypePointer StorageBuffer %41
+%176 = OpTypePointer StorageBuffer %33
+%177 = OpConstant  %26  4
+%189 = OpTypeVector %4 4
+%195 = OpVariable  %156  Output
+%213 = OpConstantNull  %4
+%217 = OpTypePointer StorageBuffer %4
+%220 = OpConstant  %26  64
+%71 = OpFunction  %2  None %72
+%70 = OpLabel
+%65 = OpVariable  %66  Function %5
+%67 = OpVariable  %68  Function %69
+%74 = OpAccessChain  %73  %58 %31
+OpBranch %76
+%76 = OpLabel
+%77 = OpLoad  %4  %65
+%78 = OpISub  %4  %77 %6
+OpStore %65 %78
+%80 = OpAccessChain  %79  %74 %31
+%81 = OpLoad  %43  %80
+%83 = OpAccessChain  %82  %74 %31 %31
+%84 = OpLoad  %37  %83
+%85 = OpLoad  %4  %65
+%86 = OpAccessChain  %82  %74 %31 %85
+%87 = OpLoad  %37  %86
+%89 = OpAccessChain  %88  %74 %31 %31 %29
 %90 = OpLoad  %9  %89
-%91 = OpLoad  %4  %62
-%92 = OpAccessChain  %85  %71 %31 %91 %29
+%91 = OpLoad  %4  %65
+%92 = OpAccessChain  %88  %74 %31 %31 %91
 %93 = OpLoad  %9  %92
-%94 = OpLoad  %4  %62
-%95 = OpLoad  %4  %62
-%96 = OpAccessChain  %85  %71 %31 %94 %95
-%97 = OpLoad  %9  %96
-%98 = OpCompositeConstruct  %36  %8 %8
-%99 = OpCompositeConstruct  %36  %10 %10
-%100 = OpCompositeConstruct  %36  %11 %11
-%101 = OpCompositeConstruct  %42  %98 %99 %100
-%102 = OpCompositeConstruct  %43  %101
-OpStore %64 %102
-%103 = OpLoad  %4  %62
-%104 = OpIAdd  %4  %103 %6
-OpStore %62 %104
-%106 = OpCompositeConstruct  %36  %12 %12
-%107 = OpCompositeConstruct  %36  %13 %13
-%108 = OpCompositeConstruct  %36  %14 %14
-%109 = OpCompositeConstruct  %42  %106 %107 %108
-%110 = OpAccessChain  %105  %64 %31
-OpStore %110 %109
-%112 = OpCompositeConstruct  %36  %15 %15
-%113 = OpAccessChain  %111  %64 %31 %31
+%94 = OpLoad  %4  %65
+%95 = OpAccessChain  %88  %74 %31 %94 %29
+%96 = OpLoad  %9  %95
+%97 = OpLoad  %4  %65
+%98 = OpLoad  %4  %65
+%99 = OpAccessChain  %88  %74 %31 %97 %98
+%100 = OpLoad  %9  %99
+%101 = OpCompositeConstruct  %37  %8 %8
+%102 = OpCompositeConstruct  %37  %10 %10
+%103 = OpCompositeConstruct  %37  %11 %11
+%104 = OpCompositeConstruct  %43  %101 %102 %103
+%105 = OpCompositeConstruct  %44  %104
+OpStore %67 %105
+%106 = OpLoad  %4  %65
+%107 = OpIAdd  %4  %106 %6
+OpStore %65 %107
+%109 = OpCompositeConstruct  %37  %12 %12
+%110 = OpCompositeConstruct  %37  %13 %13
+%111 = OpCompositeConstruct  %37  %14 %14
+%112 = OpCompositeConstruct  %43  %109 %110 %111
+%113 = OpAccessChain  %108  %67 %31
 OpStore %113 %112
-%114 = OpLoad  %4  %62
-%115 = OpCompositeConstruct  %36  %16 %16
-%116 = OpAccessChain  %111  %64 %31 %114
+%115 = OpCompositeConstruct  %37  %15 %15
+%116 = OpAccessChain  %114  %67 %31 %31
 OpStore %116 %115
-%118 = OpAccessChain  %117  %64 %31 %31 %29
-OpStore %118 %17
-%119 = OpLoad  %4  %62
-%120 = OpAccessChain  %117  %64 %31 %31 %119
-OpStore %120 %18
-%121 = OpLoad  %4  %62
-%122 = OpAccessChain  %117  %64 %31 %121 %29
-OpStore %122 %19
-%123 = OpLoad  %4  %62
-%124 = OpLoad  %4  %62
-%125 = OpAccessChain  %117  %64 %31 %123 %124
-OpStore %125 %20
+%117 = OpLoad  %4  %65
+%118 = OpCompositeConstruct  %37  %16 %16
+%119 = OpAccessChain  %114  %67 %31 %117
+OpStore %119 %118
+%121 = OpAccessChain  %120  %67 %31 %31 %29
+OpStore %121 %17
+%122 = OpLoad  %4  %65
+%123 = OpAccessChain  %120  %67 %31 %31 %122
+OpStore %123 %18
+%124 = OpLoad  %4  %65
+%125 = OpAccessChain  %120  %67 %31 %124 %29
+OpStore %125 %19
+%126 = OpLoad  %4  %65
+%127 = OpLoad  %4  %65
+%128 = OpAccessChain  %120  %67 %31 %126 %127
+OpStore %128 %20
 OpReturn
 OpFunctionEnd
-%128 = OpFunction  %9  None %129
-%127 = OpFunctionParameter  %45
-%126 = OpLabel
-OpBranch %130
-%130 = OpLabel
-%131 = OpLoad  %9  %127
-OpReturnValue %131
+%131 = OpFunction  %9  None %132
+%130 = OpFunctionParameter  %46
+%129 = OpLabel
+OpBranch %133
+%133 = OpLabel
+%134 = OpLoad  %9  %130
+OpReturnValue %134
 OpFunctionEnd
-%134 = OpFunction  %9  None %135
-%133 = OpFunctionParameter  %47
-%132 = OpLabel
-OpBranch %136
-%136 = OpLabel
-%137 = OpCompositeExtract  %46  %133 4
-%138 = OpCompositeExtract  %9  %137 9
-OpReturnValue %138
+%137 = OpFunction  %9  None %138
+%136 = OpFunctionParameter  %48
+%135 = OpLabel
+OpBranch %139
+%139 = OpLabel
+%140 = OpCompositeExtract  %47  %136 4
+%141 = OpCompositeExtract  %9  %140 9
+OpReturnValue %141
 OpFunctionEnd
-%149 = OpFunction  %2  None %69
-%143 = OpLabel
-%139 = OpVariable  %45  Function %24
-%140 = OpVariable  %141  Function %142
-%146 = OpLoad  %26  %144
-%150 = OpAccessChain  %70  %56 %31
-%151 = OpAccessChain  %72  %59 %31
-OpBranch %152
-%152 = OpLabel
-%153 = OpLoad  %9  %139
-OpStore %139 %8
-%154 = OpFunctionCall  %2  %68
-%156 = OpAccessChain  %155  %54 %31
-%157 = OpLoad  %33  %156
-%159 = OpAccessChain  %158  %54 %25
-%160 = OpLoad  %39  %159
-%163 = OpAccessChain  %162  %54 %31 %25 %31
-%164 = OpLoad  %9  %163
-%166 = OpArrayLength  %26  %54 4
-%167 = OpISub  %26  %166 %27
-%170 = OpAccessChain  %49  %54 %169 %167 %31
-%171 = OpLoad  %4  %170
-%172 = OpLoad  %44  %151
-%173 = OpFunctionCall  %9  %128 %139
-%174 = OpConvertFToS  %4  %164
-%175 = OpCompositeConstruct  %50  %171 %174 %28 %23 %22
-OpStore %140 %175
-%176 = OpIAdd  %26  %146 %29
-%177 = OpAccessChain  %63  %140 %176
-OpStore %177 %30
-%178 = OpAccessChain  %63  %140 %146
+%144 = OpFunction  %2  None %145
+%143 = OpFunctionParameter  %52
+%142 = OpLabel
+OpBranch %146
+%146 = OpLabel
+OpStore %143 %32
+OpReturn
+OpFunctionEnd
+%157 = OpFunction  %2  None %72
+%151 = OpLabel
+%147 = OpVariable  %46  Function %24
+%148 = OpVariable  %149  Function %150
+%154 = OpLoad  %26  %152
+%158 = OpAccessChain  %73  %58 %31
+%159 = OpAccessChain  %75  %61 %31
+OpBranch %160
+%160 = OpLabel
+%161 = OpLoad  %9  %147
+OpStore %147 %8
+%162 = OpFunctionCall  %2  %71
+%164 = OpAccessChain  %163  %56 %31
+%165 = OpLoad  %34  %164
+%167 = OpAccessChain  %166  %56 %25
+%168 = OpLoad  %40  %167
+%171 = OpAccessChain  %170  %56 %31 %25 %31
+%172 = OpLoad  %9  %171
+%174 = OpArrayLength  %26  %56 4
+%175 = OpISub  %26  %174 %27
+%178 = OpAccessChain  %50  %56 %177 %175 %31
 %179 = OpLoad  %4  %178
-%180 = OpFunctionCall  %9  %134 %52
-%182 = OpCompositeConstruct  %181  %179 %179 %179 %179
-%183 = OpConvertSToF  %48  %182
-%184 = OpMatrixTimesVector  %34  %157 %183
-%185 = OpCompositeConstruct  %48  %184 %10
-OpStore %147 %185
+%180 = OpLoad  %45  %159
+%181 = OpFunctionCall  %9  %131 %147
+%182 = OpConvertFToS  %4  %172
+%183 = OpCompositeConstruct  %51  %179 %182 %28 %23 %22
+OpStore %148 %183
+%184 = OpIAdd  %26  %154 %29
+%185 = OpAccessChain  %66  %148 %184
+OpStore %185 %30
+%186 = OpAccessChain  %66  %148 %154
+%187 = OpLoad  %4  %186
+%188 = OpFunctionCall  %9  %137 %54
+%190 = OpCompositeConstruct  %189  %187 %187 %187 %187
+%191 = OpConvertSToF  %49  %190
+%192 = OpMatrixTimesVector  %35  %165 %191
+%193 = OpCompositeConstruct  %49  %192 %10
+OpStore %155 %193
 OpReturn
 OpFunctionEnd
-%188 = OpFunction  %2  None %69
-%186 = OpLabel
-%189 = OpAccessChain  %72  %59 %31
-OpBranch %190
-%190 = OpLabel
-%191 = OpAccessChain  %162  %54 %31 %29 %27
-OpStore %191 %8
-%192 = OpCompositeConstruct  %34  %24 %24 %24
-%193 = OpCompositeConstruct  %34  %8 %8 %8
-%194 = OpCompositeConstruct  %34  %10 %10 %10
-%195 = OpCompositeConstruct  %34  %11 %11 %11
-%196 = OpCompositeConstruct  %33  %192 %193 %194 %195
-%197 = OpAccessChain  %155  %54 %31
-OpStore %197 %196
-%198 = OpCompositeConstruct  %38  %31 %31
-%199 = OpCompositeConstruct  %38  %29 %29
-%200 = OpCompositeConstruct  %39  %198 %199
-%201 = OpAccessChain  %158  %54 %25
-OpStore %201 %200
-%202 = OpAccessChain  %49  %54 %169 %29 %31
-OpStore %202 %6
-OpStore %189 %53
-%203 = OpCompositeConstruct  %48  %24 %24 %24 %24
-OpStore %187 %203
+%196 = OpFunction  %2  None %72
+%194 = OpLabel
+%197 = OpAccessChain  %75  %61 %31
+OpBranch %198
+%198 = OpLabel
+%199 = OpAccessChain  %170  %56 %31 %29 %27
+OpStore %199 %8
+%200 = OpCompositeConstruct  %35  %24 %24 %24
+%201 = OpCompositeConstruct  %35  %8 %8 %8
+%202 = OpCompositeConstruct  %35  %10 %10 %10
+%203 = OpCompositeConstruct  %35  %11 %11 %11
+%204 = OpCompositeConstruct  %34  %200 %201 %202 %203
+%205 = OpAccessChain  %163  %56 %31
+OpStore %205 %204
+%206 = OpCompositeConstruct  %39  %31 %31
+%207 = OpCompositeConstruct  %39  %29 %29
+%208 = OpCompositeConstruct  %40  %206 %207
+%209 = OpAccessChain  %166  %56 %25
+OpStore %209 %208
+%210 = OpAccessChain  %50  %56 %177 %29 %31
+OpStore %210 %6
+OpStore %197 %55
+%211 = OpCompositeConstruct  %49  %24 %24 %24 %24
+OpStore %195 %211
 OpReturn
 OpFunctionEnd
-%207 = OpFunction  %2  None %69
-%206 = OpLabel
-%204 = OpVariable  %63  Function %205
-OpBranch %208
-%208 = OpLabel
-%210 = OpAccessChain  %209  %54 %27
-%211 = OpAtomicLoad  %4  %210 %6 %212
-%214 = OpAccessChain  %209  %54 %27
-%213 = OpAtomicIAdd  %4  %214 %6 %212 %22
-OpStore %204 %213
-%216 = OpAccessChain  %209  %54 %27
-%215 = OpAtomicISub  %4  %216 %6 %212 %22
-OpStore %204 %215
-%218 = OpAccessChain  %209  %54 %27
-%217 = OpAtomicAnd  %4  %218 %6 %212 %22
-OpStore %204 %217
-%220 = OpAccessChain  %209  %54 %27
-%219 = OpAtomicOr  %4  %220 %6 %212 %22
-OpStore %204 %219
-%222 = OpAccessChain  %209  %54 %27
-%221 = OpAtomicXor  %4  %222 %6 %212 %22
-OpStore %204 %221
-%224 = OpAccessChain  %209  %54 %27
-%223 = OpAtomicSMin  %4  %224 %6 %212 %22
-OpStore %204 %223
-%226 = OpAccessChain  %209  %54 %27
-%225 = OpAtomicSMax  %4  %226 %6 %212 %22
-OpStore %204 %225
-%228 = OpAccessChain  %209  %54 %27
-%227 = OpAtomicExchange  %4  %228 %6 %212 %22
-OpStore %204 %227
-%229 = OpAccessChain  %209  %54 %27
-OpAtomicStore %229 %6 %212 %211
+%215 = OpFunction  %2  None %72
+%214 = OpLabel
+%212 = OpVariable  %66  Function %213
+OpBranch %216
+%216 = OpLabel
+%218 = OpAccessChain  %217  %56 %27
+%219 = OpAtomicLoad  %4  %218 %6 %220
+%222 = OpAccessChain  %217  %56 %27
+%221 = OpAtomicIAdd  %4  %222 %6 %220 %22
+OpStore %212 %221
+%224 = OpAccessChain  %217  %56 %27
+%223 = OpAtomicISub  %4  %224 %6 %220 %22
+OpStore %212 %223
+%226 = OpAccessChain  %217  %56 %27
+%225 = OpAtomicAnd  %4  %226 %6 %220 %22
+OpStore %212 %225
+%228 = OpAccessChain  %217  %56 %27
+%227 = OpAtomicOr  %4  %228 %6 %220 %22
+OpStore %212 %227
+%230 = OpAccessChain  %217  %56 %27
+%229 = OpAtomicXor  %4  %230 %6 %220 %22
+OpStore %212 %229
+%232 = OpAccessChain  %217  %56 %27
+%231 = OpAtomicSMin  %4  %232 %6 %220 %22
+OpStore %212 %231
+%234 = OpAccessChain  %217  %56 %27
+%233 = OpAtomicSMax  %4  %234 %6 %220 %22
+OpStore %212 %233
+%236 = OpAccessChain  %217  %56 %27
+%235 = OpAtomicExchange  %4  %236 %6 %220 %22
+OpStore %212 %235
+%237 = OpAccessChain  %217  %56 %27
+OpAtomicStore %237 %6 %220 %219
+OpReturn
+OpFunctionEnd
+%239 = OpFunction  %2  None %72
+%238 = OpLabel
+OpBranch %240
+%240 = OpLabel
+%241 = OpFunctionCall  %2  %144 %64
 OpReturn
 OpFunctionEnd

tests/out/wgsl/access.wgsl

@@ -20,6 +20,7 @@ var<storage, read_write> bar: Bar;
 var<uniform> baz: Baz;
 @group(0) @binding(2) 
 var<storage, read_write> qux: vec2<i32>;
+var<workgroup> val: u32;
 
 fn test_matrix_within_struct_accesses() {
     var idx: i32 = 9;
@@ -66,6 +67,11 @@ fn test_arr_as_arg(a: array<array<f32,10>,5>) -> f32 {
     return a[4][9];
 }
 
+fn assign_through_ptr_fn(p: ptr<workgroup, u32>) {
+    (*p) = 42u;
+    return;
+}
+
 @vertex 
 fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
     var foo: f32 = 0.0;
@@ -122,3 +128,9 @@ fn atomics() {
     atomicStore((&bar.atom), value_1);
     return;
 }
+
+@compute @workgroup_size(1, 1, 1) 
+fn assign_through_ptr() {
+    assign_through_ptr_fn((&val));
+    return;
+}