github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

[naga hlsl-out] Handle array types for function return values and calls (#6971)

Browse Source
This commit is contained in:
Jamie Nicol
2025-01-23 15:56:18 +00:00
committed by GitHub
parent 5e2bcc9191
commit 0282d61f91
9 changed files with 165 additions and 57 deletions
Download Patch File Download Diff File Expand all files Collapse all files

60
naga/src/back/hlsl/writer.rs
View File

@@ -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)];

8
naga/tests/in/array-in-function-return-type.wgsl
View File

@@ -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);
}

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

@@ -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;
}

32
naga/tests/out/hlsl/array-in-function-return-type.hlsl Normal file
View File

@@ -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);
}

12
naga/tests/out/hlsl/array-in-function-return-type.ron Normal file
View File

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

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

@@ -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) };
}

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

@@ -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

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

@@ -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);
}

3
naga/tests/snapshots.rs
View File

@@ -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",