[msl-out][spv-out][glsl-out] Implement ArraySize on globals.

src/back/glsl/mod.rs

@@ -129,6 +129,12 @@ impl<'a> GlobalTypeKind<'a> {
                 } => Self::Unsized(members),
                 _ => Self::WrappedStruct,
             },
+            // Naga IR permits globals to be dynamically sized arrays. Render
+            // these in GLSL as buffers.
+            crate::TypeInner::Array {
+                size: crate::ArraySize::Dynamic,
+                ..
+            } => Self::WrappedStruct,
             _ => Self::Other,
         }
     }

src/back/hlsl/help.rs

@@ -432,6 +432,7 @@ impl<'a, W: Write> super::Writer<'a, W> {
             match func_ctx.expressions[handle] {
                 crate::Expression::ArrayLength(expr) => {
                     let global_expr = match func_ctx.expressions[expr] {
+                        crate::Expression::GlobalVariable(_) => expr,
                         crate::Expression::AccessIndex { base, index: _ } => base,
                         ref other => unreachable!("Array length of {:?}", other),
                     };

src/back/msl/writer.rs

@@ -1397,6 +1397,7 @@ impl<W: Write> Writer<W> {
                             _ => return Err(Error::Validation),
                         }
                     }
+                    crate::Expression::GlobalVariable(handle) => handle,
                     _ => return Err(Error::Validation),
                 };
 
@@ -2859,14 +2860,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 +2902,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) {

src/back/spv/index.rs

@@ -2,7 +2,10 @@
 Bounds-checking for SPIR-V output.
 */
 
-use super::{selection::Selection, Block, BlockContext, Error, IdGenerator, Instruction, Word};
+use super::{
+    helpers::global_needs_wrapper, selection::Selection, Block, BlockContext, Error, IdGenerator,
+    Instruction, Word,
+};
 use crate::{arena::Handle, proc::BoundsCheckPolicy};
 
 /// The results of performing a bounds check.
@@ -32,16 +35,18 @@ pub(super) enum MaybeKnown<T> {
 impl<'w> BlockContext<'w> {
     /// Emit code to compute the length of a run-time array.
     ///
-    /// Given `array`, an expression referring to the final member of a struct,
-    /// where the member in question is a runtime-sized array, return the
+    /// Given `array`, an expression referring a runtime-sized array, return the
     /// instruction id for the array's length.
     pub(super) fn write_runtime_array_length(
         &mut self,
         array: Handle<crate::Expression>,
         block: &mut Block,
     ) -> Result<Word, Error> {
-        // Look into the expression to find the value and type of the struct
-        // holding the dynamically-sized array.
+        // Naga IR permits runtime-sized arrays as global variables or as the
+        // final member of a struct that is a global variable. SPIR-V permits
+        // only the latter, so this back end wraps bare runtime-sized arrays
+        // in a made-up struct; see `helpers::global_needs_wrapper` and its uses.
+        // This code must handle both cases.
         let (structure_id, last_member_index) = match self.ir_function.expressions[array] {
             crate::Expression::AccessIndex { base, index } => {
                 match self.ir_function.expressions[base] {
@@ -52,6 +57,14 @@ impl<'w> BlockContext<'w> {
                     _ => return Err(Error::Validation("array length expression")),
                 }
             }
+            crate::Expression::GlobalVariable(handle) => {
+                let global = &self.ir_module.global_variables[handle];
+                if !global_needs_wrapper(self.ir_module, global) {
+                    return Err(Error::Validation("array length expression"));
+                }
+
+                (self.writer.global_variables[handle.index()].var_id, 0)
+            }
             _ => return Err(Error::Validation("array length expression")),
         };
 

src/back/spv/mod.rs

@@ -433,12 +433,25 @@ impl recyclable::Recyclable for CachedExpressions {
 
 #[derive(Clone)]
 struct GlobalVariable {
-    /// ID of the variable. Not really used.
+    /// ID of the OpVariable that declares the global.
+    ///
+    /// If you need the variable's value, use [`access_id`] instead of this
+    /// field. If we wrapped the Naga IR `GlobalVariable`'s type in a struct to
+    /// comply with Vulkan's requirements, then this points to the `OpVariable`
+    /// with the synthesized struct type, whereas `access_id` points to the
+    /// field of said struct that holds the variable's actual value.
+    ///
+    /// This is used to compute the `access_id` pointer in function prologues,
+    /// and used for `ArrayLength` expressions, which do need the struct.
+    ///
+    /// [`access_id`]: GlobalVariable::access_id
     var_id: Word,
+
     /// For `AddressSpace::Handle` variables, this ID is recorded in the function
     /// prelude block (and reset before every function) as `OpLoad` of the variable.
     /// It is then used for all the global ops, such as `OpImageSample`.
     handle_id: Word,
+
     /// Actual ID used to access this variable.
     /// For wrapped buffer variables, this ID is `OpAccessChain` into the
     /// wrapper. Otherwise, the same as `var_id`.

src/lib.rs

@@ -652,7 +652,8 @@ pub enum TypeInner {
     /// An `Array` is [`SIZED`] unless its `size` is [`Dynamic`].
     /// Dynamically-sized arrays may only appear in a few situations:
     ///
-    /// -   They may appear as the last member of a [`Struct`].
+    /// -   They may appear as the type of a [`GlobalVariable`], or as the last
+    ///     member of a [`Struct`].
     ///
     /// -   They may appear as the base type of a [`Pointer`]. An
     ///     [`AccessIndex`] expression referring to a struct's final

tests/in/globals.wgsl

@@ -23,6 +23,7 @@ var<uniform> float_vecs: array<vec4<f32>, 20>;
 fn main() {
     wg[3] = alignment.v1;
     wg[2] = alignment.v3.x;
+    wg[1] = f32(arrayLength(&dummy));
     atomicStore(&at, 2u);
 
     // Valid, Foo and at is in function scope

tests/out/glsl/globals.main.Compute.glsl

@@ -15,6 +15,8 @@ shared uint at_1;
 
 layout(std430) readonly buffer Foo_block_0Compute { Foo _group_0_binding_1_cs; };
 
+layout(std430) readonly buffer type_6_block_1Compute { vec2 _group_0_binding_2_cs[]; };
+
 
 void main() {
     float Foo_1 = 1.0;
@@ -23,6 +25,7 @@ void main() {
     wg[3] = _e9;
     float _e14 = _group_0_binding_1_cs.v3_.x;
     wg[2] = _e14;
+    wg[1] = float(uint(_group_0_binding_2_cs.length()));
     at_1 = 2u;
     return;
 }

tests/out/hlsl/globals.hlsl

@@ -11,6 +11,13 @@ ByteAddressBuffer alignment : register(t1);
 ByteAddressBuffer dummy : register(t2);
 cbuffer float_vecs : register(b3) { float4 float_vecs[20]; }
 
+uint NagaBufferLength(ByteAddressBuffer buffer)
+{
+    uint ret;
+    buffer.GetDimensions(ret);
+    return ret;
+}
+
 [numthreads(1, 1, 1)]
 void main()
 {
@@ -21,6 +28,7 @@ void main()
     wg[3] = _expr9;
     float _expr14 = asfloat(alignment.Load(0+0));
     wg[2] = _expr14;
+    wg[1] = float(((NagaBufferLength(dummy) - 0) / 8));
     at_1 = 2u;
     return;
 }

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/globals.msl

@@ -23,13 +23,16 @@ kernel void main_(
   threadgroup type_2& wg
 , threadgroup metal::atomic_uint& at_1
 , device Foo& alignment [[user(fake0)]]
+, device type_6& dummy [[user(fake0)]]
+, constant _mslBufferSizes& _buffer_sizes [[user(fake0)]]
 ) {
     float Foo_1 = 1.0;
     bool at = true;
     float _e9 = alignment.v1_;
     wg.inner[3] = _e9;
     float _e14 = metal::float3(alignment.v3_).x;
     wg.inner[2] = _e14;
+    wg.inner[1] = static_cast<float>(1 + (_buffer_sizes.size3 - 0 - 8) / 8);
     metal::atomic_store_explicit(&at_1, 2u, metal::memory_order_relaxed);
     return;
 }

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;