[naga hlsl-out] Handle array types for function return values and cal…

…ls (gfx-rs#6971)

naga/src/back/hlsl/writer.rs

@@ -1336,25 +1336,37 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
 
         self.update_expressions_to_bake(module, func, info);
 
-        // Write modifier
-        if let Some(crate::FunctionResult {
-            binding:
-                Some(
-                    ref binding @ crate::Binding::BuiltIn(crate::BuiltIn::Position {
-                        invariant: true,
-                    }),
-                ),
-            ..
-        }) = func.result
-        {
-            self.write_modifier(binding)?;
-        }
-
-        // Write return type
         if let Some(ref result) = func.result {
+            // Write typedef if return type is an array
+            let array_return_type = match module.types[result.ty].inner {
+                TypeInner::Array { base, size, .. } => {
+                    let array_return_type = self.namer.call(&format!("ret_{name}"));
+                    write!(self.out, "typedef ")?;
+                    self.write_type(module, result.ty)?;
+                    write!(self.out, " {}", array_return_type)?;
+                    self.write_array_size(module, base, size)?;
+                    writeln!(self.out, ";")?;
+                    Some(array_return_type)
+                }
+                _ => None,
+            };
+
+            // Write modifier
+            if let Some(
+                ref binding @ crate::Binding::BuiltIn(crate::BuiltIn::Position { invariant: true }),
+            ) = result.binding
+            {
+                self.write_modifier(binding)?;
+            }
+
+            // Write return type
             match func_ctx.ty {
                 back::FunctionType::Function(_) => {
-                    self.write_type(module, result.ty)?;
+                    if let Some(array_return_type) = array_return_type {
+                        write!(self.out, "{array_return_type}")?;
+                    } else {
+                        self.write_type(module, result.ty)?;
+                    }
                 }
                 back::FunctionType::EntryPoint(index) => {
                     if let Some(ref ep_output) = self.entry_point_io[index as usize].output {
@@ -2212,13 +2224,21 @@ impl<'a, W: fmt::Write> super::Writer<'a, W> {
                     write!(self.out, "const ")?;
                     let name = Baked(expr).to_string();
                     let expr_ty = &func_ctx.info[expr].ty;
-                    match *expr_ty {
-                        proc::TypeResolution::Handle(handle) => self.write_type(module, handle)?,
+                    let ty_inner = match *expr_ty {
+                        proc::TypeResolution::Handle(handle) => {
+                            self.write_type(module, handle)?;
+                            &module.types[handle].inner
+                        }
                         proc::TypeResolution::Value(ref value) => {
-                            self.write_value_type(module, value)?
+                            self.write_value_type(module, value)?;
+                            value
                         }
                     };
-                    write!(self.out, " {name} = ")?;
+                    write!(self.out, " {name}")?;
+                    if let TypeInner::Array { base, size, .. } = *ty_inner {
+                        self.write_array_size(module, base, size)?;
+                    }
+                    write!(self.out, " = ")?;
                     self.named_expressions.insert(expr, name);
                 }
                 let func_name = &self.names[&NameKey::Function(function)];

naga/tests/in/array-in-function-return-type.wgsl

@@ -2,8 +2,12 @@ fn ret_array() -> array<f32, 2> {
     return array<f32, 2>(1.0, 2.0);
 }
 
+fn ret_array_array() -> array<array<f32, 2>, 3> {
+    return array<array<f32, 2>, 3>(ret_array(), ret_array(), ret_array());
+}
+
 @fragment
 fn main() -> @location(0) vec4<f32> {
-    let a = ret_array();
-    return vec4<f32>(a[0], a[1], 0.0, 1.0);
+    let a = ret_array_array();
+    return vec4<f32>(a[0][0], a[0][1], 0.0, 1.0);
 }

naga/tests/out/glsl/array-in-function-return-type.main.Fragment.glsl

@@ -9,9 +9,16 @@ float[2] ret_array() {
     return float[2](1.0, 2.0);
 }
 
-void main() {
+float[3][2] ret_array_array() {
     float _e0[2] = ret_array();
-    _fs2p_location0 = vec4(_e0[0], _e0[1], 0.0, 1.0);
+    float _e1[2] = ret_array();
+    float _e2[2] = ret_array();
+    return float[3][2](_e0, _e1, _e2);
+}
+
+void main() {
+    float _e0[3][2] = ret_array_array();
+    _fs2p_location0 = vec4(_e0[0][0], _e0[0][1], 0.0, 1.0);
     return;
 }
 

naga/tests/out/hlsl/array-in-function-return-type.hlsl

@@ -0,0 +1,32 @@
+typedef float ret_Constructarray2_float_[2];
+ret_Constructarray2_float_ Constructarray2_float_(float arg0, float arg1) {
+    float ret[2] = { arg0, arg1 };
+    return ret;
+}
+
+typedef float ret_ret_array[2];
+ret_ret_array ret_array()
+{
+    return Constructarray2_float_(1.0, 2.0);
+}
+
+typedef float ret_Constructarray3_array2_float__[3][2];
+ret_Constructarray3_array2_float__ Constructarray3_array2_float__(float arg0[2], float arg1[2], float arg2[2]) {
+    float ret[3][2] = { arg0, arg1, arg2 };
+    return ret;
+}
+
+typedef float ret_ret_array_array[3][2];
+ret_ret_array_array ret_array_array()
+{
+    const float _e0[2] = ret_array();
+    const float _e1[2] = ret_array();
+    const float _e2[2] = ret_array();
+    return Constructarray3_array2_float__(_e0, _e1, _e2);
+}
+
+float4 main() : SV_Target0
+{
+    const float _e0[3][2] = ret_array_array();
+    return float4(_e0[0][0], _e0[0][1], 0.0, 1.0);
+}

naga/tests/out/hlsl/array-in-function-return-type.ron

@@ -0,0 +1,12 @@
+(
+    vertex:[
+    ],
+    fragment:[
+        (
+            entry_point:"main",
+            target_profile:"ps_5_1",
+        ),
+    ],
+    compute:[
+    ],
+)

naga/tests/out/msl/array-in-function-return-type.msl

@@ -7,17 +7,28 @@ using metal::uint;
 struct type_1 {
     float inner[2];
 };
+struct type_2 {
+    type_1 inner[3];
+};
 
 type_1 ret_array(
 ) {
     return type_1 {1.0, 2.0};
 }
 
+type_2 ret_array_array(
+) {
+    type_1 _e0 = ret_array();
+    type_1 _e1 = ret_array();
+    type_1 _e2 = ret_array();
+    return type_2 {_e0, _e1, _e2};
+}
+
 struct main_Output {
     metal::float4 member [[color(0)]];
 };
 fragment main_Output main_(
 ) {
-    type_1 _e0 = ret_array();
-    return main_Output { metal::float4(_e0.inner[0], _e0.inner[1], 0.0, 1.0) };
+    type_2 _e0 = ret_array_array();
+    return main_Output { metal::float4(_e0.inner[0].inner[0], _e0.inner[0].inner[1], 0.0, 1.0) };
 }

naga/tests/out/spv/array-in-function-return-type.spvasm

@@ -1,42 +1,58 @@
 ; SPIR-V
 ; Version: 1.1
 ; Generator: rspirv
-; Bound: 26
+; Bound: 38
 OpCapability Shader
 %1 = OpExtInstImport "GLSL.std.450"
 OpMemoryModel Logical GLSL450
-OpEntryPoint Fragment %18 "main" %16
-OpExecutionMode %18 OriginUpperLeft
+OpEntryPoint Fragment %28 "main" %26
+OpExecutionMode %28 OriginUpperLeft
 OpDecorate %4 ArrayStride 4
-OpDecorate %16 Location 0
+OpDecorate %7 ArrayStride 8
+OpDecorate %26 Location 0
 %2 = OpTypeVoid
 %3 = OpTypeFloat 32
 %6 = OpTypeInt 32 0
 %5 = OpConstant  %6  2
 %4 = OpTypeArray %3 %5
-%7 = OpTypeVector %3 4
-%10 = OpTypeFunction %4
-%11 = OpConstant  %3  1.0
-%12 = OpConstant  %3  2.0
-%13 = OpConstantComposite  %4  %11 %12
-%17 = OpTypePointer Output %7
-%16 = OpVariable  %17  Output
-%19 = OpTypeFunction %2
-%20 = OpConstant  %3  0.0
-%9 = OpFunction  %4  None %10
-%8 = OpLabel
-OpBranch %14
-%14 = OpLabel
-OpReturnValue %13
+%8 = OpConstant  %6  3
+%7 = OpTypeArray %4 %8
+%9 = OpTypeVector %3 4
+%12 = OpTypeFunction %4
+%13 = OpConstant  %3  1.0
+%14 = OpConstant  %3  2.0
+%15 = OpConstantComposite  %4  %13 %14
+%19 = OpTypeFunction %7
+%27 = OpTypePointer Output %9
+%26 = OpVariable  %27  Output
+%29 = OpTypeFunction %2
+%30 = OpConstant  %3  0.0
+%11 = OpFunction  %4  None %12
+%10 = OpLabel
+OpBranch %16
+%16 = OpLabel
+OpReturnValue %15
 OpFunctionEnd
-%18 = OpFunction  %2  None %19
-%15 = OpLabel
-OpBranch %21
-%21 = OpLabel
-%22 = OpFunctionCall  %4  %9
-%23 = OpCompositeExtract  %3  %22 0
-%24 = OpCompositeExtract  %3  %22 1
-%25 = OpCompositeConstruct  %7  %23 %24 %20 %11
-OpStore %16 %25
+%18 = OpFunction  %7  None %19
+%17 = OpLabel
+OpBranch %20
+%20 = OpLabel
+%21 = OpFunctionCall  %4  %11
+%22 = OpFunctionCall  %4  %11
+%23 = OpFunctionCall  %4  %11
+%24 = OpCompositeConstruct  %7  %21 %22 %23
+OpReturnValue %24
+OpFunctionEnd
+%28 = OpFunction  %2  None %29
+%25 = OpLabel
+OpBranch %31
+%31 = OpLabel
+%32 = OpFunctionCall  %7  %18
+%33 = OpCompositeExtract  %4  %32 0
+%34 = OpCompositeExtract  %3  %33 0
+%35 = OpCompositeExtract  %4  %32 0
+%36 = OpCompositeExtract  %3  %35 1
+%37 = OpCompositeConstruct  %9  %34 %36 %30 %13
+OpStore %26 %37
 OpReturn
 OpFunctionEnd

naga/tests/out/wgsl/array-in-function-return-type.wgsl

@@ -2,8 +2,15 @@ fn ret_array() -> array<f32, 2> {
     return array<f32, 2>(1f, 2f);
 }
 
+fn ret_array_array() -> array<array<f32, 2>, 3> {
+    let _e0 = ret_array();
+    let _e1 = ret_array();
+    let _e2 = ret_array();
+    return array<array<f32, 2>, 3>(_e0, _e1, _e2);
+}
+
 @fragment 
 fn main() -> @location(0) vec4<f32> {
-    let _e0 = ret_array();
-    return vec4<f32>(_e0[0], _e0[1], 0f, 1f);
+    let _e0 = ret_array_array();
+    return vec4<f32>(_e0[0][0], _e0[0][1], 0f, 1f);
 }

naga/tests/snapshots.rs

@@ -676,14 +676,13 @@ fn convert_wgsl() {
     let _ = env_logger::try_init();
 
     let inputs = [
-        // TODO: merge array-in-ctor and array-in-function-return-type tests after fix HLSL issue https://github.com/gfx-rs/naga/issues/1930
         (
             "array-in-ctor",
             Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
         ),
         (
             "array-in-function-return-type",
-            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::WGSL,
+            Targets::SPIRV | Targets::METAL | Targets::GLSL | Targets::HLSL | Targets::WGSL,
         ),
         (
             "empty",