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.
gfx-rs · Feb 3, 2022 · f63003c · f63003c
1 parent 9b89c5f
commit f63003c
Show file tree

Hide file tree

Showing 23 changed files with 31 additions and 29 deletions.
diff --git a/src/back/msl/mod.rs b/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,

diff --git a/src/back/msl/writer.rs b/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(),

diff --git a/tests/out/msl/boids.msl b/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)]]
 ) {

diff --git a/tests/out/msl/bounds-check-restrict.msl b/tests/out/msl/bounds-check-restrict.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/bounds-check-zero-atomic.msl b/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>
 

diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/collatz.msl b/tests/out/msl/collatz.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/control-flow.msl b/tests/out/msl/control-flow.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/empty-global-name.msl b/tests/out/msl/empty-global-name.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/empty.msl b/tests/out/msl/empty.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/functions.msl b/tests/out/msl/functions.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/globals.msl b/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;

diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/interface.msl b/tests/out/msl/interface.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/interpolate.msl b/tests/out/msl/interpolate.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/policy-mix.msl b/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,

diff --git a/tests/out/msl/quad-vert.msl b/tests/out/msl/quad-vert.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/quad.msl b/tests/out/msl/quad.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/shadow.msl b/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)]]
 ) {

diff --git a/tests/out/msl/standard.msl b/tests/out/msl/standard.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>
 

diff --git a/tests/out/msl/texture-arg.msl b/tests/out/msl/texture-arg.msl
@@ -1,4 +1,4 @@
-// language: metal1.1
+// language: metal2.0
 #include <metal_stdlib>
 #include <simd/simd.h>