From 0282d61f91c8fba865059dcb8ac1bfa4539de93c Mon Sep 17 00:00:00 2001 From: Jamie Nicol Date: Thu, 23 Jan 2025 15:56:18 +0000 Subject: [PATCH] [naga hlsl-out] Handle array types for function return values and calls (#6971) --- naga/src/back/hlsl/writer.rs | 60 ++++++++++------ .../in/array-in-function-return-type.wgsl | 8 ++- ...in-function-return-type.main.Fragment.glsl | 11 ++- .../hlsl/array-in-function-return-type.hlsl | 32 +++++++++ .../hlsl/array-in-function-return-type.ron | 12 ++++ .../out/msl/array-in-function-return-type.msl | 15 +++- .../spv/array-in-function-return-type.spvasm | 70 ++++++++++++------- .../wgsl/array-in-function-return-type.wgsl | 11 ++- naga/tests/snapshots.rs | 3 +- 9 files changed, 165 insertions(+), 57 deletions(-) create mode 100644 naga/tests/out/hlsl/array-in-function-return-type.hlsl create mode 100644 naga/tests/out/hlsl/array-in-function-return-type.ron diff --git a/naga/src/back/hlsl/writer.rs b/naga/src/back/hlsl/writer.rs index 2b516c1977..0e7350bade 100644 --- a/naga/src/back/hlsl/writer.rs +++ b/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)]; diff --git a/naga/tests/in/array-in-function-return-type.wgsl b/naga/tests/in/array-in-function-return-type.wgsl index 21e2012e78..02cc1d5313 100644 --- a/naga/tests/in/array-in-function-return-type.wgsl +++ b/naga/tests/in/array-in-function-return-type.wgsl @@ -2,8 +2,12 @@ fn ret_array() -> array { return array(1.0, 2.0); } +fn ret_array_array() -> array, 3> { + return array, 3>(ret_array(), ret_array(), ret_array()); +} + @fragment fn main() -> @location(0) vec4 { - let a = ret_array(); - return vec4(a[0], a[1], 0.0, 1.0); + let a = ret_array_array(); + return vec4(a[0][0], a[0][1], 0.0, 1.0); } diff --git a/naga/tests/out/glsl/array-in-function-return-type.main.Fragment.glsl b/naga/tests/out/glsl/array-in-function-return-type.main.Fragment.glsl index 45fc31a622..3b35f9e353 100644 --- a/naga/tests/out/glsl/array-in-function-return-type.main.Fragment.glsl +++ b/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; } diff --git a/naga/tests/out/hlsl/array-in-function-return-type.hlsl b/naga/tests/out/hlsl/array-in-function-return-type.hlsl new file mode 100644 index 0000000000..8d0eb7411c --- /dev/null +++ b/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); +} diff --git a/naga/tests/out/hlsl/array-in-function-return-type.ron b/naga/tests/out/hlsl/array-in-function-return-type.ron new file mode 100644 index 0000000000..341a4c528e --- /dev/null +++ b/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:[ + ], +) diff --git a/naga/tests/out/msl/array-in-function-return-type.msl b/naga/tests/out/msl/array-in-function-return-type.msl index 77399f6424..71e6862e91 100644 --- a/naga/tests/out/msl/array-in-function-return-type.msl +++ b/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) }; } diff --git a/naga/tests/out/spv/array-in-function-return-type.spvasm b/naga/tests/out/spv/array-in-function-return-type.spvasm index 79e94fba8a..146e032f35 100644 --- a/naga/tests/out/spv/array-in-function-return-type.spvasm +++ b/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 \ No newline at end of file diff --git a/naga/tests/out/wgsl/array-in-function-return-type.wgsl b/naga/tests/out/wgsl/array-in-function-return-type.wgsl index 2beacd3ff4..a39680093c 100644 --- a/naga/tests/out/wgsl/array-in-function-return-type.wgsl +++ b/naga/tests/out/wgsl/array-in-function-return-type.wgsl @@ -2,8 +2,15 @@ fn ret_array() -> array { return array(1f, 2f); } +fn ret_array_array() -> array, 3> { + let _e0 = ret_array(); + let _e1 = ret_array(); + let _e2 = ret_array(); + return array, 3>(_e0, _e1, _e2); +} + @fragment fn main() -> @location(0) vec4 { - let _e0 = ret_array(); - return vec4(_e0[0], _e0[1], 0f, 1f); + let _e0 = ret_array_array(); + return vec4(_e0[0][0], _e0[0][1], 0f, 1f); } diff --git a/naga/tests/snapshots.rs b/naga/tests/snapshots.rs index 2f6cfc5852..a93f7d7129 100644 --- a/naga/tests/snapshots.rs +++ b/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",