From 8412b2912e92e657ae3ee953a3201ccb87f323eb Mon Sep 17 00:00:00 2001 From: Taylor Holliday Date: Wed, 2 Feb 2022 14:55:12 -0800 Subject: [PATCH 1/4] WGSL storage address space should always correspond to MSL device address space. See MSL spec section 4.2. --- src/back/msl/writer.rs | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 13e570ba1d..6a4bed42a1 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"), } From ffb3bd26058e7d3ce5bc563d7897efcee4b8c8d1 Mon Sep 17 00:00:00 2001 From: Taylor Holliday Date: Wed, 2 Feb 2022 16:19:28 -0800 Subject: [PATCH 2/4] Update baselines --- tests/out/msl/boids.msl | 2 +- tests/out/msl/globals.msl | 2 +- tests/out/msl/policy-mix.msl | 2 +- tests/out/msl/shadow.msl | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) 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/globals.msl b/tests/out/msl/globals.msl index dde8160766..5eaca7543e 100644 --- a/tests/out/msl/globals.msl +++ b/tests/out/msl/globals.msl @@ -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/policy-mix.msl b/tests/out/msl/policy-mix.msl index 06ed990be3..534981736c 100644 --- a/tests/out/msl/policy-mix.msl +++ b/tests/out/msl/policy-mix.msl @@ -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/shadow.msl b/tests/out/msl/shadow.msl index 54e918de49..56098ca07f 100644 --- a/tests/out/msl/shadow.msl +++ b/tests/out/msl/shadow.msl @@ -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)]] ) { From 773f11266c3b66e8c25251aac25b3611eda9d6bf Mon Sep 17 00:00:00 2001 From: Taylor Holliday Date: Thu, 3 Feb 2022 09:17:07 -0800 Subject: [PATCH 3/4] Add error for storage address space if MSL version < 2. --- src/back/msl/mod.rs | 2 ++ src/back/msl/writer.rs | 3 +++ 2 files changed, 5 insertions(+) diff --git a/src/back/msl/mod.rs b/src/back/msl/mod.rs index 8e26c7b16f..ec6b311f47 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)] diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index 6a4bed42a1..26afaae64a 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -3030,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(), From f4be3863d8f0a4721d3cd960e8891d46291f3e71 Mon Sep 17 00:00:00 2001 From: Taylor Holliday Date: Thu, 3 Feb 2022 09:24:02 -0800 Subject: [PATCH 4/4] Update default MSL version to 2.0. --- src/back/msl/mod.rs | 2 +- tests/out/msl/bounds-check-restrict.msl | 2 +- tests/out/msl/bounds-check-zero-atomic.msl | 2 +- tests/out/msl/bounds-check-zero.msl | 2 +- tests/out/msl/collatz.msl | 2 +- tests/out/msl/control-flow.msl | 2 +- tests/out/msl/empty-global-name.msl | 2 +- tests/out/msl/empty.msl | 2 +- tests/out/msl/functions.msl | 2 +- tests/out/msl/globals.msl | 2 +- tests/out/msl/image.msl | 2 +- tests/out/msl/interface.msl | 2 +- tests/out/msl/interpolate.msl | 2 +- tests/out/msl/math-functions.msl | 2 +- tests/out/msl/operators.msl | 2 +- tests/out/msl/policy-mix.msl | 2 +- tests/out/msl/quad-vert.msl | 2 +- tests/out/msl/quad.msl | 2 +- tests/out/msl/shadow.msl | 2 +- tests/out/msl/standard.msl | 2 +- tests/out/msl/texture-arg.msl | 2 +- 21 files changed, 21 insertions(+), 21 deletions(-) diff --git a/src/back/msl/mod.rs b/src/back/msl/mod.rs index ec6b311f47..b886d1573b 100644 --- a/src/back/msl/mod.rs +++ b/src/back/msl/mod.rs @@ -187,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/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 5eaca7543e..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 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 534981736c..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 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 56098ca07f..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 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