msl: qualify read-only storage with const

src/back/msl/writer.rs

@@ -238,13 +238,12 @@ impl<'a> TypedGlobalVariable<'a> {
 
         let (space, access, reference) = match var.space.to_msl_name() {
             Some(space) if self.reference => {
-                let access = match var.space {
-                    crate::AddressSpace::Private | crate::AddressSpace::WorkGroup
-                        if !self.usage.contains(valid::GlobalUse::WRITE) =>
-                    {
-                        "const"
-                    }
-                    _ => "",
+                let access = if var.space.needs_access_qualifier()
+                    && !self.usage.contains(valid::GlobalUse::WRITE)
+                {
+                    "const"
+                } else {
+                    ""
                 };
                 (space, access, "&")
             }
@@ -401,13 +400,25 @@ impl crate::AddressSpace {
     /// passed through any functions called from the entry point.
     fn needs_pass_through(&self) -> bool {
         match *self {
-            crate::AddressSpace::Uniform
-            | crate::AddressSpace::Storage { .. }
-            | crate::AddressSpace::Private
-            | crate::AddressSpace::WorkGroup
-            | crate::AddressSpace::PushConstant
-            | crate::AddressSpace::Handle => true,
-            crate::AddressSpace::Function => false,
+            Self::Uniform
+            | Self::Storage { .. }
+            | Self::Private
+            | Self::WorkGroup
+            | Self::PushConstant
+            | Self::Handle => true,
+            Self::Function => false,
+        }
+    }
+
+    /// Returns true if the address space may need a "const" qualifier.
+    fn needs_access_qualifier(&self) -> bool {
+        match *self {
+            //Note: we are ignoring the storage access here, and instead
+            // 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::Uniform | Self::PushConstant | Self::Handle | Self::Function => false,
         }
     }
 

tests/in/access.param.ron

@@ -8,6 +8,12 @@
 		lang_version: (2, 0),
 		per_stage_map: (
 			vs: (
+				resources: {
+					(group: 0, binding: 0): (buffer: Some(0), mutable: false),
+				},
+				sizes_buffer: Some(24),
+			),
+			fs: (
 				resources: {
 					(group: 0, binding: 0): (buffer: Some(0), mutable: true),
 				},

tests/in/access.wgsl

@@ -20,7 +20,7 @@ fn read_from_private(foo: ptr<function, f32>) -> f32 {
 }
 
 @stage(vertex)
-fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
+fn foo_vert(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
     var foo: f32 = 0.0;
     // We should check that backed doesn't skip this expression
     let baz: f32 = foo;
@@ -37,12 +37,6 @@ fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
 	let data_pointer: ptr<storage, i32, read_write> = &bar.data[0].value;
 	let foo_value = read_from_private(&foo);
 
-	// test storage stores
-	bar.matrix[1].z = 1.0;
-	bar.matrix = mat4x4<f32>(vec4<f32>(0.0), vec4<f32>(1.0), vec4<f32>(2.0), vec4<f32>(3.0));
-	bar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
-	bar.data[1].value = 1;
-
 	// test array indexing
 	var c = array<i32, 5>(a, i32(b), 3, 4, 5);
 	c[vi + 1u] = 42;
@@ -51,6 +45,17 @@ fn foo(@builtin(vertex_index) vi: u32) -> @builtin(position) vec4<f32> {
 	return matrix * vec4<f32>(vec4<i32>(value));
 }
 
+@stage(fragment)
+fn foo_frag() -> @location(0) vec4<f32> {
+	// test storage stores
+	bar.matrix[1].z = 1.0;
+	bar.matrix = mat4x4<f32>(vec4<f32>(0.0), vec4<f32>(1.0), vec4<f32>(2.0), vec4<f32>(3.0));
+	bar.arr = array<vec2<u32>, 2>(vec2<u32>(0u), vec2<u32>(1u));
+	bar.data[1].value = 1;
+
+	return vec4<f32>(0.0);
+}
+
 @stage(compute) @workgroup_size(1)
 fn atomics() {
 	var tmp: i32;

tests/in/globals.wgsl

@@ -11,7 +11,7 @@ struct Foo {
     v1: f32;
 };
 @group(0) @binding(1)
-var<storage> alignment: Foo;
+var<storage, read_write> alignment: Foo;
 
 @group(0) @binding(2)
 var<storage> dummy: array<vec2<f32>>;
@@ -23,6 +23,7 @@ var<uniform> float_vecs: array<vec4<f32>, 20>;
 fn main() {
     wg[3] = alignment.v1;
     wg[2] = alignment.v3.x;
+    alignment.v1 = 4.0;
     wg[1] = f32(arrayLength(&dummy));
     atomicStore(&at, 2u);
 

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

@@ -17,8 +17,8 @@ layout(std430) buffer Bar_block_0Compute {
 } _group_0_binding_0_cs;
 
 
-float read_from_private(inout float foo_2) {
-    float _e2 = foo_2;
+float read_from_private(inout float foo_1) {
+    float _e2 = foo_1;
     return _e2;
 }
 

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

@@ -0,0 +1,32 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+struct AlignedWrapper {
+    int value;
+};
+layout(std430) buffer Bar_block_0Fragment {
+    mat4x4 matrix;
+    mat2x2 matrix_array[2];
+    int atom;
+    uvec2 arr[2];
+    AlignedWrapper data[];
+} _group_0_binding_0_fs;
+
+layout(location = 0) out vec4 _fs2p_location0;
+
+float read_from_private(inout float foo_1) {
+    float _e2 = foo_1;
+    return _e2;
+}
+
+void main() {
+    _group_0_binding_0_fs.matrix[1][2] = 1.0;
+    _group_0_binding_0_fs.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0));
+    _group_0_binding_0_fs.arr = uvec2[2](uvec2(0u), uvec2(1u));
+    _group_0_binding_0_fs.data[1].value = 1;
+    _fs2p_location0 = vec4(0.0);
+    return;
+}
+

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

@@ -15,26 +15,22 @@ layout(std430) buffer Bar_block_0Vertex {
 } _group_0_binding_0_vs;
 
 
-float read_from_private(inout float foo_2) {
-    float _e2 = foo_2;
+float read_from_private(inout float foo_1) {
+    float _e2 = foo_1;
     return _e2;
 }
 
 void main() {
     uint vi = uint(gl_VertexID);
-    float foo_1 = 0.0;
+    float foo = 0.0;
     int c[5] = int[5](0, 0, 0, 0, 0);
-    float baz = foo_1;
-    foo_1 = 1.0;
+    float baz = foo;
+    foo = 1.0;
     mat4x4 matrix = _group_0_binding_0_vs.matrix;
     uvec2 arr[2] = _group_0_binding_0_vs.arr;
     float b = _group_0_binding_0_vs.matrix[3][0];
     int a = _group_0_binding_0_vs.data[(uint(_group_0_binding_0_vs.data.length()) - 2u)].value;
-    float _e27 = read_from_private(foo_1);
-    _group_0_binding_0_vs.matrix[1][2] = 1.0;
-    _group_0_binding_0_vs.matrix = mat4x4(vec4(0.0), vec4(1.0), vec4(2.0), vec4(3.0));
-    _group_0_binding_0_vs.arr = uvec2[2](uvec2(0u), uvec2(1u));
-    _group_0_binding_0_vs.data[1].value = 1;
+    float _e27 = read_from_private(foo);
     c = int[5](a, int(b), 3, 4, 5);
     c[(vi + 1u)] = 42;
     int value = c[vi];

tests/out/glsl/globals.main.Compute.glsl

@@ -13,7 +13,7 @@ shared float wg[10];
 
 shared uint at_1;
 
-layout(std430) readonly buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; };
+layout(std430) buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; };
 
 layout(std430) readonly buffer type_6_block_1Compute { vec2 _group_0_binding_2_cs[]; };
 
@@ -25,6 +25,7 @@ void main() {
     wg[3] = _e9;
     float _e14 = _group_0_binding_1_cs.v3_.x;
     wg[2] = _e14;
+    _group_0_binding_1_cs.v1_ = 4.0;
     wg[1] = float(uint(_group_0_binding_2_cs.length()));
     at_1 = 2u;
     return;

tests/out/hlsl/access.hlsl

@@ -5,9 +5,9 @@ struct AlignedWrapper {
 
 RWByteAddressBuffer bar : register(u0);
 
-float read_from_private(inout float foo_2)
+float read_from_private(inout float foo_1)
 {
-    float _expr2 = foo_2;
+    float _expr2 = foo_1;
     return _expr2;
 }
 
@@ -18,18 +18,29 @@ uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
     return ret;
 }
 
-float4 foo(uint vi : SV_VertexID) : SV_Position
+float4 foo_vert(uint vi : SV_VertexID) : SV_Position
 {
-    float foo_1 = 0.0;
+    float foo = 0.0;
     int c[5] = {(int)0,(int)0,(int)0,(int)0,(int)0};
 
-    float baz = foo_1;
-    foo_1 = 1.0;
+    float baz = foo;
+    foo = 1.0;
     float4x4 matrix_ = float4x4(asfloat(bar.Load4(0+0)), asfloat(bar.Load4(0+16)), asfloat(bar.Load4(0+32)), asfloat(bar.Load4(0+48)));
     uint2 arr[2] = {asuint(bar.Load2(104+0)), asuint(bar.Load2(104+8))};
     float b = asfloat(bar.Load(0+48+0));
     int a = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 120) / 8) - 2u)*8+120));
-    const float _e27 = read_from_private(foo_1);
+    const float _e27 = read_from_private(foo);
+    {
+        int _result[5]={ a, int(b), 3, 4, 5 };
+        for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
+    }
+    c[(vi + 1u)] = 42;
+    int value = c[vi];
+    return mul(float4(int4(value.xxxx)), matrix_);
+}
+
+float4 foo_frag() : SV_Target0
+{
     bar.Store(8+16+0, asuint(1.0));
     {
         float4x4 _value2 = float4x4(float4(0.0.xxxx), float4(1.0.xxxx), float4(2.0.xxxx), float4(3.0.xxxx));
@@ -44,13 +55,7 @@ float4 foo(uint vi : SV_VertexID) : SV_Position
         bar.Store2(104+8, asuint(_value2[1]));
     }
     bar.Store(0+8+120, asuint(1));
-    {
-        int _result[5]={ a, int(b), 3, 4, 5 };
-        for(int _i=0; _i<5; ++_i) c[_i] = _result[_i];
-    }
-    c[(vi + 1u)] = 42;
-    int value = c[vi];
-    return mul(float4(int4(value.xxxx)), matrix_);
+    return float4(0.0.xxxx);
 }
 
 [numthreads(1, 1, 1)]

tests/out/hlsl/access.hlsl.config

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

tests/out/hlsl/globals.hlsl

@@ -7,7 +7,7 @@ struct Foo {
 
 groupshared float wg[10];
 groupshared uint at_1;
-ByteAddressBuffer alignment : register(t1);
+RWByteAddressBuffer alignment : register(u1);
 ByteAddressBuffer dummy : register(t2);
 cbuffer float_vecs : register(b3) { float4 float_vecs[20]; }
 
@@ -28,6 +28,7 @@ void main()
     wg[3] = _expr9;
     float _expr14 = asfloat(alignment.Load(0+0));
     wg[2] = _expr14;
+    alignment.Store(12, asuint(4.0));
     wg[1] = float(((NagaBufferLength(dummy) - 0) / 8));
     at_1 = 2u;
     return;

tests/out/msl/access.msl

@@ -31,39 +31,50 @@ struct type_13 {
 };
 
 float read_from_private(
-    thread float& foo_2
+    thread float& foo_1
 ) {
-    float _e2 = foo_2;
+    float _e2 = foo_1;
     return _e2;
 }
 
-struct fooInput {
+struct foo_vertInput {
 };
-struct fooOutput {
+struct foo_vertOutput {
     metal::float4 member [[position]];
 };
-vertex fooOutput foo(
+vertex foo_vertOutput foo_vert(
   uint vi [[vertex_id]]
-, device Bar& bar [[buffer(0)]]
+, device Bar const& bar [[buffer(0)]]
 , constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
 ) {
-    float foo_1 = 0.0;
+    float foo = 0.0;
     type_13 c;
-    float baz = foo_1;
-    foo_1 = 1.0;
+    float baz = foo;
+    foo = 1.0;
     metal::float4x4 matrix = bar.matrix;
     type_6 arr = bar.arr;
     float b = bar.matrix[3].x;
     int a = bar.data[(1 + (_buffer_sizes.size0 - 120 - 8) / 8) - 2u].value;
-    float _e27 = read_from_private(foo_1);
+    float _e27 = read_from_private(foo);
+    for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
+    c.inner[vi + 1u] = 42;
+    int value = c.inner[vi];
+    return foo_vertOutput { matrix * static_cast<metal::float4>(metal::int4(value)) };
+}
+
+
+struct foo_fragOutput {
+    metal::float4 member_1 [[color(0)]];
+};
+fragment foo_fragOutput foo_frag(
+  device Bar& bar [[buffer(0)]]
+, constant _mslBufferSizes& _buffer_sizes [[buffer(24)]]
+) {
     bar.matrix[1].z = 1.0;
     bar.matrix = metal::float4x4(metal::float4(0.0), metal::float4(1.0), metal::float4(2.0), metal::float4(3.0));
     for(int _i=0; _i<2; ++_i) bar.arr.inner[_i] = type_6 {metal::uint2(0u), metal::uint2(1u)}.inner[_i];
     bar.data[1].value = 1;
-    for(int _i=0; _i<5; ++_i) c.inner[_i] = type_13 {a, static_cast<int>(b), 3, 4, 5}.inner[_i];
-    c.inner[vi + 1u] = 42;
-    int value = c.inner[vi];
-    return fooOutput { matrix * static_cast<metal::float4>(metal::int4(value)) };
+    return foo_fragOutput { metal::float4(0.0) };
 }
 
 

tests/out/msl/boids.msl

@@ -33,7 +33,7 @@ struct main_Input {
 kernel void main_(
   metal::uint3 global_invocation_id [[thread_position_in_grid]]
 , constant SimParams& params [[buffer(0)]]
-, device Particles& particlesSrc [[buffer(1)]]
+, device Particles const& particlesSrc [[buffer(1)]]
 , device Particles& particlesDst [[buffer(2)]]
 , constant _mslBufferSizes& _buffer_sizes [[buffer(3)]]
 ) {