[msl-out][spv-out][glsl-out][hlsl-out] Fix 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),
                 };
 

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()].access_id, 0)
+            }
             _ => return Err(Error::Validation("array length expression")),
         };
 

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

@@ -23,13 +23,15 @@ kernel void main_(
   threadgroup type_2& wg
 , threadgroup metal::atomic_uint& at_1
 , device Foo& alignment [[user(fake0)]]
+, device type_6& dummy [[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/spv/globals.spvasm

@@ -1,32 +1,32 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 61
+; Bound: 66
 OpCapability Shader
 OpExtension "SPV_KHR_storage_buffer_storage_class"
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint GLCompute %40 "main"
-OpExecutionMode %40 LocalSize 1 1 1
-OpDecorate %15 ArrayStride 4
-OpMemberDecorate %17 0 Offset 0
-OpMemberDecorate %17 1 Offset 12
-OpDecorate %19 ArrayStride 8
-OpDecorate %21 ArrayStride 16
-OpDecorate %26 NonWritable
-OpDecorate %26 DescriptorSet 0
-OpDecorate %26 Binding 1
-OpDecorate %27 Block
-OpMemberDecorate %27 0 Offset 0
-OpDecorate %29 NonWritable
-OpDecorate %29 DescriptorSet 0
-OpDecorate %29 Binding 2
-OpDecorate %30 Block
-OpMemberDecorate %30 0 Offset 0
-OpDecorate %32 DescriptorSet 0
-OpDecorate %32 Binding 3
-OpDecorate %33 Block
-OpMemberDecorate %33 0 Offset 0
+OpEntryPoint GLCompute %41 "main"
+OpExecutionMode %41 LocalSize 1 1 1
+OpDecorate %16 ArrayStride 4
+OpMemberDecorate %18 0 Offset 0
+OpMemberDecorate %18 1 Offset 12
+OpDecorate %20 ArrayStride 8
+OpDecorate %22 ArrayStride 16
+OpDecorate %27 NonWritable
+OpDecorate %27 DescriptorSet 0
+OpDecorate %27 Binding 1
+OpDecorate %28 Block
+OpMemberDecorate %28 0 Offset 0
+OpDecorate %30 NonWritable
+OpDecorate %30 DescriptorSet 0
+OpDecorate %30 Binding 2
+OpDecorate %31 Block
+OpMemberDecorate %31 0 Offset 0
+OpDecorate %33 DescriptorSet 0
+OpDecorate %33 Binding 3
+OpDecorate %34 Block
+OpMemberDecorate %34 0 Offset 0
 %2 = OpTypeVoid
 %4 = OpTypeBool
 %3 = OpConstantTrue  %4
@@ -36,59 +36,65 @@ OpMemberDecorate %33 0 Offset 0
 %7 = OpConstant  %8  20
 %9 = OpConstant  %8  3
 %10 = OpConstant  %8  2
-%11 = OpConstant  %6  2
-%13 = OpTypeFloat 32
-%12 = OpConstant  %13  1.0
-%14 = OpConstantTrue  %4
-%15 = OpTypeArray %13 %5
-%16 = OpTypeVector %13 3
-%17 = OpTypeStruct %16 %13
-%18 = OpTypeVector %13 2
-%19 = OpTypeRuntimeArray %18
-%20 = OpTypeVector %13 4
-%21 = OpTypeArray %20 %7
-%23 = OpTypePointer Workgroup %15
-%22 = OpVariable  %23  Workgroup
-%25 = OpTypePointer Workgroup %6
-%24 = OpVariable  %25  Workgroup
-%27 = OpTypeStruct %17
-%28 = OpTypePointer StorageBuffer %27
-%26 = OpVariable  %28  StorageBuffer
-%30 = OpTypeStruct %19
-%31 = OpTypePointer StorageBuffer %30
-%29 = OpVariable  %31  StorageBuffer
-%33 = OpTypeStruct %21
-%34 = OpTypePointer Uniform %33
-%32 = OpVariable  %34  Uniform
-%36 = OpTypePointer Function %13
-%38 = OpTypePointer Function %4
-%41 = OpTypeFunction %2
-%42 = OpTypePointer StorageBuffer %17
-%43 = OpConstant  %6  0
-%45 = OpTypePointer StorageBuffer %19
-%46 = OpTypePointer Uniform %21
-%48 = OpTypePointer Workgroup %13
-%49 = OpTypePointer StorageBuffer %13
-%50 = OpConstant  %6  1
-%53 = OpConstant  %6  3
-%55 = OpTypePointer StorageBuffer %16
-%56 = OpTypePointer StorageBuffer %13
-%60 = OpConstant  %6  256
-%40 = OpFunction  %2  None %41
-%39 = OpLabel
-%35 = OpVariable  %36  Function %12
-%37 = OpVariable  %38  Function %14
-%44 = OpAccessChain  %42  %26 %43
-OpBranch %47
-%47 = OpLabel
-%51 = OpAccessChain  %49  %44 %50
-%52 = OpLoad  %13  %51
-%54 = OpAccessChain  %48  %22 %53
-OpStore %54 %52
-%57 = OpAccessChain  %56  %44 %43 %43
-%58 = OpLoad  %13  %57
-%59 = OpAccessChain  %48  %22 %11
-OpStore %59 %58
-OpAtomicStore %24 %10 %60 %11
+%11 = OpConstant  %8  1
+%12 = OpConstant  %6  2
+%14 = OpTypeFloat 32
+%13 = OpConstant  %14  1.0
+%15 = OpConstantTrue  %4
+%16 = OpTypeArray %14 %5
+%17 = OpTypeVector %14 3
+%18 = OpTypeStruct %17 %14
+%19 = OpTypeVector %14 2
+%20 = OpTypeRuntimeArray %19
+%21 = OpTypeVector %14 4
+%22 = OpTypeArray %21 %7
+%24 = OpTypePointer Workgroup %16
+%23 = OpVariable  %24  Workgroup
+%26 = OpTypePointer Workgroup %6
+%25 = OpVariable  %26  Workgroup
+%28 = OpTypeStruct %18
+%29 = OpTypePointer StorageBuffer %28
+%27 = OpVariable  %29  StorageBuffer
+%31 = OpTypeStruct %20
+%32 = OpTypePointer StorageBuffer %31
+%30 = OpVariable  %32  StorageBuffer
+%34 = OpTypeStruct %22
+%35 = OpTypePointer Uniform %34
+%33 = OpVariable  %35  Uniform
+%37 = OpTypePointer Function %14
+%39 = OpTypePointer Function %4
+%42 = OpTypeFunction %2
+%43 = OpTypePointer StorageBuffer %18
+%44 = OpConstant  %6  0
+%46 = OpTypePointer StorageBuffer %20
+%48 = OpTypePointer Uniform %22
+%50 = OpTypePointer Workgroup %14
+%51 = OpTypePointer StorageBuffer %14
+%52 = OpConstant  %6  1
+%55 = OpConstant  %6  3
+%57 = OpTypePointer StorageBuffer %17
+%58 = OpTypePointer StorageBuffer %14
+%65 = OpConstant  %6  256
+%41 = OpFunction  %2  None %42
+%40 = OpLabel
+%36 = OpVariable  %37  Function %13
+%38 = OpVariable  %39  Function %15
+%45 = OpAccessChain  %43  %27 %44
+%47 = OpAccessChain  %46  %30 %44
+OpBranch %49
+%49 = OpLabel
+%53 = OpAccessChain  %51  %45 %52
+%54 = OpLoad  %14  %53
+%56 = OpAccessChain  %50  %23 %55
+OpStore %56 %54
+%59 = OpAccessChain  %58  %45 %44 %44
+%60 = OpLoad  %14  %59
+%61 = OpAccessChain  %50  %23 %12
+OpStore %61 %60
+%62 = OpArrayLength  %6  %47 0
+%63 = OpConvertUToF  %14  %62
+%64 = OpAccessChain  %50  %23 %52
+OpStore %64 %63
+OpAtomicStore %25 %10 %65 %12
 OpReturn
 OpFunctionEnd

tests/out/wgsl/globals.wgsl

@@ -23,6 +23,7 @@ fn main() {
     wg[3] = _e9;
     let _e14 = alignment.v3_.x;
     wg[2] = _e14;
+    wg[1] = f32(arrayLength((&dummy)));
     atomicStore((&at_1), 2u);
     return;
 }