WGSL storage address space should always correspond to MSL device add… (#1711)

* WGSL storage address space should always correspond to MSL device address space.

See MSL spec section 4.2.

* Update baselines

* Add error for storage address space if MSL version < 2.

* Update default MSL version to 2.0.

src/back/msl/mod.rs

@@ -141,6 +141,8 @@ pub enum Error {
     UnsupportedBuiltIn(crate::BuiltIn),
     #[error("capability {0:?} is not supported")]
     CapabilityNotSupported(crate::valid::Capabilities),
+    #[error("address space {0:?} is not supported for target MSL version")]
+    UnsupportedAddressSpace(crate::AddressSpace),
 }
 
 #[derive(Clone, Debug, PartialEq, thiserror::Error)]
@@ -185,7 +187,7 @@ pub struct Options {
 impl Default for Options {
     fn default() -> Self {
         Options {
-            lang_version: (1, 1),
+            lang_version: (2, 0),
             per_stage_map: PerStageMap::default(),
             inline_samplers: Vec::new(),
             spirv_cross_compatibility: false,

src/back/msl/writer.rs

@@ -410,10 +410,7 @@ impl crate::AddressSpace {
         match self {
             Self::Handle => None,
             Self::Uniform | Self::PushConstant => Some("constant"),
-            Self::Storage { access } if access.contains(crate::StorageAccess::STORE) => {
-                Some("device")
-            }
-            Self::Storage { .. } => Some("constant"),
+            Self::Storage { .. } => Some("device"),
             Self::Private | Self::Function => Some("thread"),
             Self::WorkGroup => Some("threadgroup"),
         }
@@ -3033,6 +3030,9 @@ impl<W: Write> Writer<W> {
                         options.resolve_push_constants(ep.stage).ok()
                     }
                     crate::AddressSpace::WorkGroup => None,
+                    crate::AddressSpace::Storage { .. } if options.lang_version < (2, 0) => {
+                        return Err(Error::UnsupportedAddressSpace(var.space))
+                    }
                     _ => options
                         .resolve_resource_binding(ep.stage, var.binding.as_ref().unwrap())
                         .ok(),

tests/out/msl/boids.msl

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

tests/out/msl/bounds-check-restrict.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/bounds-check-zero-atomic.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/bounds-check-zero.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/collatz.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/control-flow.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/empty-global-name.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/empty.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/functions.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/globals.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -22,7 +22,7 @@ struct type_8 {
 kernel void main_(
   threadgroup type_2& wg
 , threadgroup metal::atomic_uint& at_1
-, constant Foo& alignment [[user(fake0)]]
+, device Foo& alignment [[user(fake0)]]
 ) {
     float Foo_1 = 1.0;
     bool at = true;

tests/out/msl/image.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/interface.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/interpolate.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/math-functions.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/operators.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/policy-mix.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -34,7 +34,7 @@ metal::float4 mock_function(
     metal::int2 c,
     int i,
     int l,
-    constant InStorage& in_storage,
+    device InStorage& in_storage,
     constant InUniform& in_uniform,
     metal::texture2d_array<float, metal::access::sample> image_2d_array,
     threadgroup type_5 const& in_workgroup,

tests/out/msl/quad-vert.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/quad.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/shadow.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 
@@ -46,7 +46,7 @@ struct fs_mainOutput {
 fragment fs_mainOutput fs_main(
   fs_mainInput varyings [[stage_in]]
 , constant Globals& u_globals [[user(fake0)]]
-, constant Lights& s_lights [[user(fake0)]]
+, device Lights& s_lights [[user(fake0)]]
 , metal::depth2d_array<float, metal::access::sample> t_shadow [[user(fake0)]]
 , metal::sampler sampler_shadow [[user(fake0)]]
 ) {

tests/out/msl/standard.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

tests/out/msl/texture-arg.msl

@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>