fix(hlsl): emit constructor functions for arrays (#2281)

* test(hlsl-out): add failing case for array as ctor arg

See issue #2184.

* fix(hlsl): emit constructor functions for arrays

src/back/hlsl/help.rs

@@ -882,6 +882,12 @@ impl<'a, W: Write> super::Writer<'a, W> {
                             }
                             crate::TypeInner::Array { base, .. } => {
                                 write_wrapped_constructor(writer, base, module)?;
+                                let constructor = WrappedConstructor { ty };
+                                if !writer.wrapped.constructors.contains(&constructor) {
+                                    writer
+                                        .write_wrapped_constructor_function(module, constructor)?;
+                                    writer.wrapped.constructors.insert(constructor);
+                                }
                             }
                             _ => {}
                         };

src/back/hlsl/storage.rs

@@ -147,12 +147,16 @@ impl<W: fmt::Write> super::Writer<'_, W> {
                 size: crate::ArraySize::Constant(const_handle),
                 ..
             } => {
-                write!(self.out, "{{")?;
+                let constructor = super::help::WrappedConstructor {
+                    ty: result_ty.handle().unwrap(),
+                };
+                self.write_wrapped_constructor_function_name(module, constructor)?;
+                write!(self.out, "(")?;
                 let count = module.constants[const_handle].to_array_length().unwrap();
                 let stride = module.types[base].inner.size(&module.constants);
                 let iter = (0..count).map(|i| (TypeResolution::Handle(base), stride * i));
                 self.write_storage_load_sequence(module, var_handle, iter, func_ctx)?;
-                write!(self.out, "}}")?;
+                write!(self.out, ")")?;
             }
             crate::TypeInner::Struct { ref members, .. } => {
                 let constructor = super::help::WrappedConstructor {

tests/in/array-in-ctor.param.ron

@@ -0,0 +1,2 @@
+(
+)

tests/in/array-in-ctor.wgsl

@@ -0,0 +1,10 @@
+struct Ah {
+    inner: array<f32, 2>,
+};
+@group(0) @binding(0)
+var<storage> ah: Ah;
+
+@compute @workgroup_size(1)
+fn cs_main() {
+    _ = ah;
+}

tests/out/glsl/array-in-ctor.cs_main.Compute.glsl

@@ -0,0 +1,17 @@
+#version 310 es
+
+precision highp float;
+precision highp int;
+
+layout(local_size_x = 1, local_size_y = 1, local_size_z = 1) in;
+
+struct Ah {
+    float inner[2];
+};
+layout(std430) readonly buffer Ah_block_0Compute { Ah _group_0_binding_0_cs; };
+
+
+void main() {
+    Ah unnamed = _group_0_binding_0_cs;
+}
+

tests/out/hlsl/access.hlsl

@@ -249,6 +249,12 @@ void assign_array_through_ptr_fn(inout float4 foo_2[2])
     return;
 }
 
+typedef uint2 ret_Constructarray2_uint2_[2];
+ret_Constructarray2_uint2_ Constructarray2_uint2_(uint2 arg0, uint2 arg1) {
+    uint2 ret[2] = { arg0, arg1 };
+    return ret;
+}
+
 uint NagaBufferLengthRW(RWByteAddressBuffer buffer)
 {
     uint ret;
@@ -273,7 +279,7 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
     test_matrix_within_struct_accesses();
     test_matrix_within_array_within_struct_accesses();
     float4x3 _matrix = float4x3(asfloat(bar.Load3(0+0)), asfloat(bar.Load3(0+16)), asfloat(bar.Load3(0+32)), asfloat(bar.Load3(0+48)));
-    uint2 arr_1[2] = {asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8))};
+    uint2 arr_1[2] = Constructarray2_uint2_(asuint(bar.Load2(144+0)), asuint(bar.Load2(144+8)));
     float b = asfloat(bar.Load(0+48+0));
     int a_1 = asint(bar.Load(0+(((NagaBufferLengthRW(bar) - 160) / 8) - 2u)*8+160));
     int2 c = asint(qux.Load2(0));
@@ -285,12 +291,6 @@ float4 foo_vert(uint vi : SV_VertexID) : SV_Position
     return float4(mul(float4((value).xxxx), _matrix), 2.0);
 }
 
-typedef uint2 ret_Constructarray2_uint2_[2];
-ret_Constructarray2_uint2_ Constructarray2_uint2_(uint2 arg0, uint2 arg1) {
-    uint2 ret[2] = { arg0, arg1 };
-    return ret;
-}
-
 float4 foo_frag() : SV_Target0
 {
     bar.Store(8+16+0, asuint(1.0));

tests/out/hlsl/array-in-ctor.hlsl

@@ -0,0 +1,24 @@
+
+struct Ah {
+    float inner[2];
+};
+
+ByteAddressBuffer ah : register(t0);
+
+typedef float ret_Constructarray2_float_[2];
+ret_Constructarray2_float_ Constructarray2_float_(float arg0, float arg1) {
+    float ret[2] = { arg0, arg1 };
+    return ret;
+}
+
+Ah ConstructAh(float arg0[2]) {
+    Ah ret = (Ah)0;
+    ret.inner = arg0;
+    return ret;
+}
+
+[numthreads(1, 1, 1)]
+void cs_main()
+{
+    Ah unnamed = ConstructAh(Constructarray2_float_(asfloat(ah.Load(0+0)), asfloat(ah.Load(0+4))));
+}

tests/out/hlsl/array-in-ctor.hlsl.config

@@ -0,0 +1,3 @@
+vertex=()
+fragment=()
+compute=(cs_main:cs_5_1 )

tests/out/msl/array-in-ctor.msl

@@ -0,0 +1,18 @@
+// language: metal2.0
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+using metal::uint;
+
+struct type_1 {
+    float inner[2];
+};
+struct Ah {
+    type_1 inner;
+};
+
+kernel void cs_main(
+  device Ah const& ah [[user(fake0)]]
+) {
+    Ah unnamed = ah;
+}

tests/out/spv/array-in-ctor.spvasm

@@ -0,0 +1,38 @@
+; SPIR-V
+; Version: 1.1
+; Generator: rspirv
+; Bound: 20
+OpCapability Shader
+OpExtension "SPV_KHR_storage_buffer_storage_class"
+%1 = OpExtInstImport "GLSL.std.450"
+OpMemoryModel Logical GLSL450
+OpEntryPoint GLCompute %12 "cs_main"
+OpExecutionMode %12 LocalSize 1 1 1
+OpDecorate %6 ArrayStride 4
+OpMemberDecorate %7 0 Offset 0
+OpDecorate %8 NonWritable
+OpDecorate %8 DescriptorSet 0
+OpDecorate %8 Binding 0
+OpDecorate %9 Block
+OpMemberDecorate %9 0 Offset 0
+%2 = OpTypeVoid
+%4 = OpTypeInt 32 1
+%3 = OpConstant  %4  2
+%5 = OpTypeFloat 32
+%6 = OpTypeArray %5 %3
+%7 = OpTypeStruct %6
+%9 = OpTypeStruct %7
+%10 = OpTypePointer StorageBuffer %9
+%8 = OpVariable  %10  StorageBuffer
+%13 = OpTypeFunction %2
+%14 = OpTypePointer StorageBuffer %7
+%16 = OpTypeInt 32 0
+%15 = OpConstant  %16  0
+%12 = OpFunction  %2  None %13
+%11 = OpLabel
+%17 = OpAccessChain  %14  %8 %15
+OpBranch %18
+%18 = OpLabel
+%19 = OpLoad  %7  %17
+OpReturn
+OpFunctionEnd

tests/out/wgsl/array-in-ctor.wgsl

@@ -0,0 +1,11 @@
+struct Ah {
+    inner: array<f32,2>,
+}
+
+@group(0) @binding(0) 
+var<storage> ah: Ah;
+
+@compute @workgroup_size(1, 1, 1) 
+fn cs_main() {
+    _ = ah;
+}

tests/snapshots.rs

@@ -426,6 +426,10 @@ fn convert_wgsl() {
 
     let root = env!("CARGO_MANIFEST_DIR");
     let inputs = [
+        (
+            "array-in-ctor",
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
+        ),
         (
             "empty",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,