diff --git a/src/back/msl/mod.rs b/src/back/msl/mod.rs index 8e26c7b16f..b886d1573b 100644 --- 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 index 13e570ba1d..26afaae64a 100644 --- 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 Writer { 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 index b717466f9b..9f60489978 100644 --- 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 index 16e4320085..93bbf910be 100644 --- 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 #include diff --git a/tests/out/msl/bounds-check-zero-atomic.msl b/tests/out/msl/bounds-check-zero-atomic.msl index d6911b1aa1..038bef6bfc 100644 --- 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 #include diff --git a/tests/out/msl/bounds-check-zero.msl b/tests/out/msl/bounds-check-zero.msl index 0662b33409..b61e523aab 100644 --- 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 #include diff --git a/tests/out/msl/collatz.msl b/tests/out/msl/collatz.msl index 22769a8a5e..540ddbbb59 100644 --- a/tests/out/msl/collatz.msl +++ b/tests/out/msl/collatz.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/control-flow.msl b/tests/out/msl/control-flow.msl index bed973ead0..e0d2f0cebc 100644 --- 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 #include diff --git a/tests/out/msl/empty-global-name.msl b/tests/out/msl/empty-global-name.msl index efae101f40..3b9dfec5fc 100644 --- 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 #include diff --git a/tests/out/msl/empty.msl b/tests/out/msl/empty.msl index ddc6b6bf1f..c3bb1a71b9 100644 --- a/tests/out/msl/empty.msl +++ b/tests/out/msl/empty.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/functions.msl b/tests/out/msl/functions.msl index 6ace92f749..11574c679e 100644 --- a/tests/out/msl/functions.msl +++ b/tests/out/msl/functions.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/globals.msl b/tests/out/msl/globals.msl index dde8160766..c8c58b3391 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include @@ -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 index 68d52340fa..7fe6776648 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/interface.msl b/tests/out/msl/interface.msl index f95620c375..6461aee75a 100644 --- a/tests/out/msl/interface.msl +++ b/tests/out/msl/interface.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/interpolate.msl b/tests/out/msl/interpolate.msl index b92adb3991..0f89a547ba 100644 --- a/tests/out/msl/interpolate.msl +++ b/tests/out/msl/interpolate.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/math-functions.msl b/tests/out/msl/math-functions.msl index 42f8be358f..ee0acbca84 100644 --- 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 #include diff --git a/tests/out/msl/operators.msl b/tests/out/msl/operators.msl index 7fa931624f..eebd1e73b8 100644 --- a/tests/out/msl/operators.msl +++ b/tests/out/msl/operators.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/policy-mix.msl b/tests/out/msl/policy-mix.msl index 06ed990be3..893c700ca0 100644 --- 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 #include @@ -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 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 index de1bfb598e..87114a8d74 100644 --- 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 #include diff --git a/tests/out/msl/quad.msl b/tests/out/msl/quad.msl index a99532b12c..a03d61bf57 100644 --- a/tests/out/msl/quad.msl +++ b/tests/out/msl/quad.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/shadow.msl b/tests/out/msl/shadow.msl index 54e918de49..d4868235ae 100644 --- a/tests/out/msl/shadow.msl +++ b/tests/out/msl/shadow.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include @@ -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 t_shadow [[user(fake0)]] , metal::sampler sampler_shadow [[user(fake0)]] ) { diff --git a/tests/out/msl/standard.msl b/tests/out/msl/standard.msl index cf7adc7d6b..0164252374 100644 --- a/tests/out/msl/standard.msl +++ b/tests/out/msl/standard.msl @@ -1,4 +1,4 @@ -// language: metal1.1 +// language: metal2.0 #include #include diff --git a/tests/out/msl/texture-arg.msl b/tests/out/msl/texture-arg.msl index 49533df9ef..20e76207d3 100644 --- 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 #include