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

[teoxoy]

Changes pointers to private or workgroup address spaces to always be read-write (according to the WGSL spec).

fixes #1262

The added test used to fail with

error: no matching function for call to 'assign_through_ptr_fn'
    assign_through_ptr_fn(val);
    ^~~~~~~~~~~~~~~~~~~~~
note: candidate function not viable: 1st argument ('const threadgroup metal::uint' (aka 'const threadgroup unsigned int')) would lose const qualifier
void assign_through_ptr_fn(
     ^
1 error generated.

References
https://gpuweb.github.io/gpuweb/wgsl/#access-mode-defaults
https://gpuweb.github.io/gpuweb/wgsl/#address-space-table