[msl-out] Always check whether _buffer_sizes arg is needed.

The _buffer_sizes argument should be inserted regardless of whether or not `
`options.fake_missing_bindings` is set, so lift the computation of
`supports_array_length` out of that conditional.

src/back/msl/writer.rs

@@ -2859,14 +2859,20 @@ impl<W: Write> Writer<W> {
             let fun = &ep.function;
             let fun_info = mod_info.get_entry_point(ep_index);
             let mut ep_error = None;
-            let mut supports_array_length = false;
 
             log::trace!(
                 "entry point {:?}, index {:?}",
                 fun.name.as_deref().unwrap_or("(anonymous)"),
                 ep_index
             );
 
+            // Is any global variable used by this entry point dynamically sized?
+            let supports_array_length = module
+                .global_variables
+                .iter()
+                .filter(|&(handle, _)| !fun_info[handle].is_empty())
+                .any(|(_, var)| needs_array_length(var.ty, &module.types));
+
             // skip this entry point if any global bindings are missing,
             // or their types are incompatible.
             if !options.fake_missing_bindings {
@@ -2895,7 +2901,6 @@ impl<W: Write> Writer<W> {
                             break;
                         }
                     }
-                    supports_array_length |= needs_array_length(var.ty, &module.types);
                 }
                 if supports_array_length {
                     if let Err(err) = options.resolve_sizes_buffer(ep.stage) {

tests/out/msl/collatz.msl

@@ -42,6 +42,7 @@ struct main_Input {
 kernel void main_(
   metal::uint3 global_id [[thread_position_in_grid]]
 , device PrimeIndices& v_indices [[user(fake0)]]
+, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
 ) {
     metal::uint _e8 = v_indices.data[global_id.x];
     metal::uint _e9 = collatz_iterations(_e8);

tests/out/msl/shadow.msl

@@ -49,6 +49,7 @@ fragment fs_mainOutput fs_main(
 , device Lights& s_lights [[user(fake0)]]
 , metal::depth2d_array<float, metal::access::sample> t_shadow [[user(fake0)]]
 , metal::sampler sampler_shadow [[user(fake0)]]
+, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
 ) {
     const auto raw_normal = varyings.raw_normal;
     const auto position = varyings.position;