nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2025-02-17 01:12:41 +00:00
Code Issues Packages Projects Releases Wiki Activity

[msl-out] Properly rename entry point arguments for struct members. (#1766)

Browse Source
This commit is contained in:
Jim Blandy 2022-03-08 07:07:30 -08:00 committed by GitHub
parent 79845371d3
commit c84aa77579
No known key found for this signature in database
GPG Key ID: 4AEE18F83AFDEB23
10 changed files with 383 additions and 176 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

92
src/back/msl/writer.rs
View File

@ -3288,9 +3288,6 @@ impl<W: Write> Writer<W> {
writeln!(self.out)?;
let stage_out_name = format!("{}Output", fun_name);
let stage_in_name = format!("{}Input", fun_name);
let (em_str, in_mode, out_mode) = match ep.stage {
crate::ShaderStage::Vertex => (
"vertex",
@ -3307,35 +3304,44 @@ impl<W: Write> Writer<W> {
}
};
let mut argument_members = Vec::new();
// List all the Naga `EntryPoint`'s `Function`'s arguments,
// flattening structs into their members. In Metal, we will pass
// each of these values to the entry point as a separate argument—
// except for the varyings, handled next.
let mut flattened_arguments = Vec::new();
for (arg_index, arg) in fun.arguments.iter().enumerate() {
match module.types[arg.ty].inner {
crate::TypeInner::Struct { ref members, .. } => {
for (member_index, member) in members.iter().enumerate() {
argument_members.push((
NameKey::StructMember(arg.ty, member_index as u32),
let member_index = member_index as u32;
flattened_arguments.push((
NameKey::StructMember(arg.ty, member_index),
member.ty,
member.binding.as_ref(),
))
));
}
}
_ => argument_members.push((
_ => flattened_arguments.push((
NameKey::EntryPointArgument(ep_index as _, arg_index as u32),
arg.ty,
arg.binding.as_ref(),
)),
}
}
// Identify the varyings among the argument values, and emit a
// struct type named `<fun>Input` to hold them.
let stage_in_name = format!("{}Input", fun_name);
let varyings_member_name = self.namer.call("varyings");
let mut varying_count = 0;
if !argument_members.is_empty() {
let mut has_varyings = false;
if !flattened_arguments.is_empty() {
writeln!(self.out, "struct {} {{", stage_in_name)?;
for &(ref name_key, ty, binding) in argument_members.iter() {
for &(ref name_key, ty, binding) in flattened_arguments.iter() {
let binding = match binding {
Some(ref binding @ &crate::Binding::Location { .. }) => binding,
_ => continue,
};
varying_count += 1;
has_varyings = true;
let name = &self.names[name_key];
let ty_name = TypeContext {
handle: ty,
@ -3352,6 +3358,9 @@ impl<W: Write> Writer<W> {
writeln!(self.out, "}};")?;
}
// Define a struct type named for the return value, if any, named
// `<fun>Output`.
let stage_out_name = format!("{}Output", fun_name);
let result_member_name = self.namer.call("member");
let result_type_name = match fun.result {
Some(ref result) => {
@ -3444,10 +3453,14 @@ impl<W: Write> Writer<W> {
}
None => "void",
};
writeln!(self.out, "{} {} {}(", em_str, result_type_name, fun_name)?;
// Write the entry point function's name, and begin its argument list.
writeln!(self.out, "{} {} {}(", em_str, result_type_name, fun_name)?;
let mut is_first_argument = true;
if varying_count != 0 {
// If we have produced a struct holding the `EntryPoint`'s
// `Function`'s arguments' varyings, pass that struct first.
if has_varyings {
writeln!(
self.out,
" {} {} [[stage_in]]",
@ -3455,12 +3468,31 @@ impl<W: Write> Writer<W> {
)?;
is_first_argument = false;
}
for &(ref name_key, ty, binding) in argument_members.iter() {
// Then pass the remaining arguments not included in the varyings
// struct.
//
// Since `Namer.reset` wasn't expecting struct members to be
// suddenly injected into the normal namespace like this,
// `self.names` doesn't keep them distinct from other variables.
// Generate fresh names for these arguments, and remember the
// mapping.
let mut flattened_member_names = FastHashMap::default();
for &(ref name_key, ty, binding) in flattened_arguments.iter() {
let binding = match binding {
Some(ref binding @ &crate::Binding::BuiltIn(..)) => binding,
_ => continue,
};
let name = &self.names[name_key];
let name = if let NameKey::StructMember(ty, index) = *name_key {
// We should always insert a fresh entry here, but use
// `or_insert` to get a reference to the `String` we just
// inserted.
flattened_member_names
.entry(NameKey::StructMember(ty, index))
.or_insert_with(|| self.namer.call(&self.names[name_key]))
} else {
&self.names[name_key]
};
let ty_name = TypeContext {
handle: ty,
arena: &module.types,
@ -3479,6 +3511,11 @@ impl<W: Write> Writer<W> {
resolved.try_fmt_decorated(&mut self.out)?;
writeln!(self.out)?;
}
// Those global variables used by this entry point and its callees
// get passed as arguments. `Private` globals are an exception, they
// don't outlive this invocation, so we declare them below as locals
// within the entry point.
for (handle, var) in module.global_variables.iter() {
let usage = fun_info[handle];
if usage.is_empty() || var.space == crate::AddressSpace::Private {
@ -3534,6 +3571,8 @@ impl<W: Write> Writer<W> {
writeln!(self.out)?;
}
// If this entry uses any variable-length arrays, their sizes are
// passed as a final struct-typed argument.
if supports_array_length {
// this is checked earlier
let resolved = options.resolve_sizes_buffer(ep.stage).unwrap();
@ -3603,7 +3642,16 @@ impl<W: Write> Writer<W> {
}
}
// Now refactor the inputs in a way that the rest of the code expects
// Now take the arguments that we gathered into structs, and the
// structs that we flattened into arguments, and emit local
// variables with initializers that put everything back the way the
// body code expects.
//
// If we had to generate fresh names for struct members passed as
// arguments, be sure to use those names when rebuilding the struct.
//
// "Each day, I change some zeros to ones, and some ones to zeros.
// The rest, I leave alone."
for (arg_index, arg) in fun.arguments.iter().enumerate() {
let arg_name =
&self.names[&NameKey::EntryPointArgument(ep_index as _, arg_index as u32)];
@ -3618,8 +3666,14 @@ impl<W: Write> Writer<W> {
arg_name
)?;
for (member_index, member) in members.iter().enumerate() {
let name =
&self.names[&NameKey::StructMember(arg.ty, member_index as u32)];
let key = NameKey::StructMember(arg.ty, member_index as u32);
// If it's not in the varying struct, then we should
// have passed it as its own argument and assigned
// it a new name.
let name = match member.binding {
Some(crate::Binding::BuiltIn(_)) => &flattened_member_names[&key],
_ => &self.names[&key],
};
if member_index != 0 {
write!(self.out, ", ")?;
}

14
tests/in/interface.wgsl
View File

@ -45,3 +45,17 @@ fn compute(
) {
output[0] = global_id.x + local_id.x + local_index + wg_id.x + num_wgs.x;
}
struct Input1 {
@builtin(vertex_index) index: u32;
};
struct Input2 {
@builtin(instance_index) index: u32;
};
@stage(vertex)
fn vertex_two_structs(in1: Input1, in2: Input2) -> @builtin(position) vec4<f32> {
var index = 2u;
return vec4<f32>(f32(in1.index), f32(in2.index), f32(index), 0.0);
}

16
tests/out/hlsl/interface.hlsl
View File

@ -16,6 +16,14 @@ struct FragmentOutput {
linear float color : SV_Target0;
};
struct Input1_ {
uint index : SV_VertexID;
};
struct Input2_ {
uint index : SV_InstanceID;
};
groupshared uint output[1];
struct VertexOutput_vertex {
@ -72,3 +80,11 @@ void compute(uint3 global_id : SV_DispatchThreadID, uint3 local_id : SV_GroupThr
output[0] = ((((global_id.x + local_id.x) + local_index) + wg_id.x) + uint3(_NagaConstants.base_vertex, _NagaConstants.base_instance, _NagaConstants.other).x);
return;
}
float4 vertex_two_structs(Input1_ in1_, Input2_ in2_) : SV_Position
{
uint index = 2u;
uint _expr9 = index;
return float4(float((_NagaConstants.base_vertex + in1_.index)), float((_NagaConstants.base_instance + in2_.index)), float(_expr9), 0.0);
}

2
tests/out/hlsl/interface.hlsl.config
View File

@ -1,3 +1,3 @@
vertex=(vertex:vs_5_1 )
vertex=(vertex:vs_5_1 vertex_two_structs:vs_5_1 )
fragment=(fragment:ps_5_1 )
compute=(compute:cs_5_1 )

24
tests/out/msl/interface.msl
View File

@ -16,6 +16,12 @@ struct FragmentOutput {
struct type_4 {
uint inner[1];
};
struct Input1_ {
uint index;
};
struct Input2_ {
uint index;
};
struct vertex_Input {
uint color [[attribute(10)]];
@ -73,3 +79,21 @@ kernel void compute_(
output.inner[0] = (((global_id.x + local_id.x) + local_index) + wg_id.x) + num_wgs.x;
return;
}
struct vertex_two_structsInput {
};
struct vertex_two_structsOutput {
metal::float4 member_3 [[position]];
float _point_size [[point_size]];
};
vertex vertex_two_structsOutput vertex_two_structs(
uint index_1 [[vertex_id]]
, uint index_2 [[instance_id]]
) {
const Input1_ in1_ = { index_1 };
const Input2_ in2_ = { index_2 };
uint index = 2u;
uint _e9 = index;
return vertex_two_structsOutput { metal::float4(static_cast<float>(in1_.index), static_cast<float>(in2_.index), static_cast<float>(_e9), 0.0), 1.0 };
}

103
tests/out/spv/interface.compute.spvasm
View File

@ -1,23 +1,25 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 46
; Bound: 49
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %32 "compute" %20 %23 %25 %28 %30
OpExecutionMode %32 LocalSize 1 1 1
OpMemberDecorate %12 0 Offset 0
OpMemberDecorate %12 1 Offset 16
OpEntryPoint GLCompute %35 "compute" %23 %26 %28 %31 %33
OpExecutionMode %35 LocalSize 1 1 1
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 4
OpMemberDecorate %13 2 Offset 8
OpDecorate %15 ArrayStride 4
OpDecorate %20 BuiltIn GlobalInvocationId
OpDecorate %23 BuiltIn LocalInvocationId
OpDecorate %25 BuiltIn LocalInvocationIndex
OpDecorate %28 BuiltIn WorkgroupId
OpDecorate %30 BuiltIn NumWorkgroups
OpMemberDecorate %13 1 Offset 16
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 1 Offset 4
OpMemberDecorate %14 2 Offset 8
OpDecorate %16 ArrayStride 4
OpMemberDecorate %18 0 Offset 0
OpMemberDecorate %19 0 Offset 0
OpDecorate %23 BuiltIn GlobalInvocationId
OpDecorate %26 BuiltIn LocalInvocationId
OpDecorate %28 BuiltIn LocalInvocationIndex
OpDecorate %31 BuiltIn WorkgroupId
OpDecorate %33 BuiltIn NumWorkgroups
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpConstant %4 1.0
@ -27,42 +29,45 @@ OpDecorate %30 BuiltIn NumWorkgroups
%9 = OpTypeInt 32 1
%8 = OpConstant %9 1
%10 = OpConstant %9 0
%11 = OpTypeVector %4 4
%12 = OpTypeStruct %11 %4
%13 = OpTypeStruct %4 %6 %4
%14 = OpTypeBool
%15 = OpTypeArray %6 %8
%16 = OpTypeVector %6 3
%18 = OpTypePointer Workgroup %15
%17 = OpVariable %18 Workgroup
%21 = OpTypePointer Input %16
%20 = OpVariable %21 Input
%23 = OpVariable %21 Input
%26 = OpTypePointer Input %6
%25 = OpVariable %26 Input
%28 = OpVariable %21 Input
%30 = OpVariable %21 Input
%33 = OpTypeFunction %2
%35 = OpTypePointer Workgroup %6
%44 = OpConstant %6 0
%32 = OpFunction %2 None %33
%19 = OpLabel
%22 = OpLoad %16 %20
%24 = OpLoad %16 %23
%27 = OpLoad %6 %25
%29 = OpLoad %16 %28
%31 = OpLoad %16 %30
OpBranch %34
%34 = OpLabel
%36 = OpCompositeExtract %6 %22 0
%37 = OpCompositeExtract %6 %24 0
%38 = OpIAdd %6 %36 %37
%39 = OpIAdd %6 %38 %27
%40 = OpCompositeExtract %6 %29 0
%11 = OpConstant %6 2
%12 = OpTypeVector %4 4
%13 = OpTypeStruct %12 %4
%14 = OpTypeStruct %4 %6 %4
%15 = OpTypeBool
%16 = OpTypeArray %6 %8
%17 = OpTypeVector %6 3
%18 = OpTypeStruct %6
%19 = OpTypeStruct %6
%21 = OpTypePointer Workgroup %16
%20 = OpVariable %21 Workgroup
%24 = OpTypePointer Input %17
%23 = OpVariable %24 Input
%26 = OpVariable %24 Input
%29 = OpTypePointer Input %6
%28 = OpVariable %29 Input
%31 = OpVariable %24 Input
%33 = OpVariable %24 Input
%36 = OpTypeFunction %2
%38 = OpTypePointer Workgroup %6
%47 = OpConstant %6 0
%35 = OpFunction %2 None %36
%22 = OpLabel
%25 = OpLoad %17 %23
%27 = OpLoad %17 %26
%30 = OpLoad %6 %28
%32 = OpLoad %17 %31
%34 = OpLoad %17 %33
OpBranch %37
%37 = OpLabel
%39 = OpCompositeExtract %6 %25 0
%40 = OpCompositeExtract %6 %27 0
%41 = OpIAdd %6 %39 %40
%42 = OpCompositeExtract %6 %31 0
%43 = OpIAdd %6 %41 %42
%45 = OpAccessChain %35 %17 %44
OpStore %45 %43
%42 = OpIAdd %6 %41 %30
%43 = OpCompositeExtract %6 %32 0
%44 = OpIAdd %6 %42 %43
%45 = OpCompositeExtract %6 %34 0
%46 = OpIAdd %6 %44 %45
%48 = OpAccessChain %38 %20 %47
OpStore %48 %46
OpReturn
OpFunctionEnd

123
tests/out/spv/interface.fragment.spvasm
View File

@ -1,28 +1,30 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 51
; Bound: 54
OpCapability Shader
OpCapability SampleRateShading
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Fragment %38 "fragment" %19 %22 %25 %28 %31 %33 %35 %37
OpExecutionMode %38 OriginUpperLeft
OpExecutionMode %38 DepthReplacing
OpMemberDecorate %12 0 Offset 0
OpMemberDecorate %12 1 Offset 16
OpEntryPoint Fragment %41 "fragment" %22 %25 %28 %31 %34 %36 %38 %40
OpExecutionMode %41 OriginUpperLeft
OpExecutionMode %41 DepthReplacing
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 4
OpMemberDecorate %13 2 Offset 8
OpDecorate %15 ArrayStride 4
OpDecorate %19 BuiltIn FragCoord
OpDecorate %22 Location 1
OpDecorate %25 BuiltIn FrontFacing
OpDecorate %28 BuiltIn SampleId
OpDecorate %31 BuiltIn SampleMask
OpDecorate %33 BuiltIn FragDepth
OpDecorate %35 BuiltIn SampleMask
OpDecorate %37 Location 0
OpMemberDecorate %13 1 Offset 16
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 1 Offset 4
OpMemberDecorate %14 2 Offset 8
OpDecorate %16 ArrayStride 4
OpMemberDecorate %18 0 Offset 0
OpMemberDecorate %19 0 Offset 0
OpDecorate %22 BuiltIn FragCoord
OpDecorate %25 Location 1
OpDecorate %28 BuiltIn FrontFacing
OpDecorate %31 BuiltIn SampleId
OpDecorate %34 BuiltIn SampleMask
OpDecorate %36 BuiltIn FragDepth
OpDecorate %38 BuiltIn SampleMask
OpDecorate %40 Location 0
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpConstant %4 1.0
@ -32,50 +34,53 @@ OpDecorate %37 Location 0
%9 = OpTypeInt 32 1
%8 = OpConstant %9 1
%10 = OpConstant %9 0
%11 = OpTypeVector %4 4
%12 = OpTypeStruct %11 %4
%13 = OpTypeStruct %4 %6 %4
%14 = OpTypeBool
%15 = OpTypeArray %6 %8
%16 = OpTypeVector %6 3
%20 = OpTypePointer Input %11
%19 = OpVariable %20 Input
%23 = OpTypePointer Input %4
%11 = OpConstant %6 2
%12 = OpTypeVector %4 4
%13 = OpTypeStruct %12 %4
%14 = OpTypeStruct %4 %6 %4
%15 = OpTypeBool
%16 = OpTypeArray %6 %8
%17 = OpTypeVector %6 3
%18 = OpTypeStruct %6
%19 = OpTypeStruct %6
%23 = OpTypePointer Input %12
%22 = OpVariable %23 Input
%26 = OpTypePointer Input %14
%26 = OpTypePointer Input %4
%25 = OpVariable %26 Input
%29 = OpTypePointer Input %6
%29 = OpTypePointer Input %15
%28 = OpVariable %29 Input
%31 = OpVariable %29 Input
%34 = OpTypePointer Output %4
%33 = OpVariable %34 Output
%36 = OpTypePointer Output %6
%35 = OpVariable %36 Output
%37 = OpVariable %34 Output
%39 = OpTypeFunction %2
%38 = OpFunction %2 None %39
%17 = OpLabel
%21 = OpLoad %11 %19
%24 = OpLoad %4 %22
%18 = OpCompositeConstruct %12 %21 %24
%27 = OpLoad %14 %25
%30 = OpLoad %6 %28
%32 = OpLoad %6 %31
OpBranch %40
%40 = OpLabel
%41 = OpShiftLeftLogical %6 %5 %30
%42 = OpBitwiseAnd %6 %32 %41
%43 = OpSelect %4 %27 %3 %7
%44 = OpCompositeExtract %4 %18 1
%45 = OpCompositeConstruct %13 %44 %42 %43
%46 = OpCompositeExtract %4 %45 0
OpStore %33 %46
%47 = OpLoad %4 %33
%48 = OpExtInst %4 %1 FClamp %47 %7 %3
OpStore %33 %48
%49 = OpCompositeExtract %6 %45 1
OpStore %35 %49
%50 = OpCompositeExtract %4 %45 2
OpStore %37 %50
%32 = OpTypePointer Input %6
%31 = OpVariable %32 Input
%34 = OpVariable %32 Input
%37 = OpTypePointer Output %4
%36 = OpVariable %37 Output
%39 = OpTypePointer Output %6
%38 = OpVariable %39 Output
%40 = OpVariable %37 Output
%42 = OpTypeFunction %2
%41 = OpFunction %2 None %42
%20 = OpLabel
%24 = OpLoad %12 %22
%27 = OpLoad %4 %25
%21 = OpCompositeConstruct %13 %24 %27
%30 = OpLoad %15 %28
%33 = OpLoad %6 %31
%35 = OpLoad %6 %34
OpBranch %43
%43 = OpLabel
%44 = OpShiftLeftLogical %6 %5 %33
%45 = OpBitwiseAnd %6 %35 %44
%46 = OpSelect %4 %30 %3 %7
%47 = OpCompositeExtract %4 %21 1
%48 = OpCompositeConstruct %14 %47 %45 %46
%49 = OpCompositeExtract %4 %48 0
OpStore %36 %49
%50 = OpLoad %4 %36
%51 = OpExtInst %4 %1 FClamp %50 %7 %3
OpStore %36 %51
%52 = OpCompositeExtract %6 %48 1
OpStore %38 %52
%53 = OpCompositeExtract %4 %48 2
OpStore %40 %53
OpReturn
OpFunctionEnd

101
tests/out/spv/interface.vertex.spvasm
View File

@ -1,24 +1,26 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 41
; Bound: 44
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %31 "vertex" %18 %21 %23 %25 %27 %29
OpMemberDecorate %12 0 Offset 0
OpMemberDecorate %12 1 Offset 16
OpEntryPoint Vertex %34 "vertex" %21 %24 %26 %28 %30 %32
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 4
OpMemberDecorate %13 2 Offset 8
OpDecorate %15 ArrayStride 4
OpDecorate %18 BuiltIn VertexIndex
OpDecorate %21 BuiltIn InstanceIndex
OpDecorate %23 Location 10
OpDecorate %23 Flat
OpDecorate %25 BuiltIn Position
OpDecorate %27 Location 1
OpDecorate %29 BuiltIn PointSize
OpMemberDecorate %13 1 Offset 16
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 1 Offset 4
OpMemberDecorate %14 2 Offset 8
OpDecorate %16 ArrayStride 4
OpMemberDecorate %18 0 Offset 0
OpMemberDecorate %19 0 Offset 0
OpDecorate %21 BuiltIn VertexIndex
OpDecorate %24 BuiltIn InstanceIndex
OpDecorate %26 Location 10
OpDecorate %26 Flat
OpDecorate %28 BuiltIn Position
OpDecorate %30 Location 1
OpDecorate %32 BuiltIn PointSize
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpConstant %4 1.0
@ -28,39 +30,42 @@ OpDecorate %29 BuiltIn PointSize
%9 = OpTypeInt 32 1
%8 = OpConstant %9 1
%10 = OpConstant %9 0
%11 = OpTypeVector %4 4
%12 = OpTypeStruct %11 %4
%13 = OpTypeStruct %4 %6 %4
%14 = OpTypeBool
%15 = OpTypeArray %6 %8
%16 = OpTypeVector %6 3
%19 = OpTypePointer Input %6
%18 = OpVariable %19 Input
%21 = OpVariable %19 Input
%23 = OpVariable %19 Input
%26 = OpTypePointer Output %11
%25 = OpVariable %26 Output
%28 = OpTypePointer Output %4
%27 = OpVariable %28 Output
%30 = OpTypePointer Output %4
%29 = OpVariable %30 Output
%32 = OpTypeFunction %2
%31 = OpFunction %2 None %32
%17 = OpLabel
%20 = OpLoad %6 %18
%22 = OpLoad %6 %21
%24 = OpLoad %6 %23
OpStore %29 %3
OpBranch %33
%33 = OpLabel
%34 = OpIAdd %6 %20 %22
%35 = OpIAdd %6 %34 %24
%36 = OpCompositeConstruct %11 %3 %3 %3 %3
%37 = OpConvertUToF %4 %35
%38 = OpCompositeConstruct %12 %36 %37
%39 = OpCompositeExtract %11 %38 0
OpStore %25 %39
%40 = OpCompositeExtract %4 %38 1
OpStore %27 %40
%11 = OpConstant %6 2
%12 = OpTypeVector %4 4
%13 = OpTypeStruct %12 %4
%14 = OpTypeStruct %4 %6 %4
%15 = OpTypeBool
%16 = OpTypeArray %6 %8
%17 = OpTypeVector %6 3
%18 = OpTypeStruct %6
%19 = OpTypeStruct %6
%22 = OpTypePointer Input %6
%21 = OpVariable %22 Input
%24 = OpVariable %22 Input
%26 = OpVariable %22 Input
%29 = OpTypePointer Output %12
%28 = OpVariable %29 Output
%31 = OpTypePointer Output %4
%30 = OpVariable %31 Output
%33 = OpTypePointer Output %4
%32 = OpVariable %33 Output
%35 = OpTypeFunction %2
%34 = OpFunction %2 None %35
%20 = OpLabel
%23 = OpLoad %6 %21
%25 = OpLoad %6 %24
%27 = OpLoad %6 %26
OpStore %32 %3
OpBranch %36
%36 = OpLabel
%37 = OpIAdd %6 %23 %25
%38 = OpIAdd %6 %37 %27
%39 = OpCompositeConstruct %12 %3 %3 %3 %3
%40 = OpConvertUToF %4 %38
%41 = OpCompositeConstruct %13 %39 %40
%42 = OpCompositeExtract %12 %41 0
OpStore %28 %42
%43 = OpCompositeExtract %4 %41 1
OpStore %30 %43
OpReturn
OpFunctionEnd

68
tests/out/spv/interface.vertex_two_structs.spvasm Normal file
View File

@ -0,0 +1,68 @@
; SPIR-V
; Version: 1.0
; Generator: rspirv
; Bound: 45
OpCapability Shader
%1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450
OpEntryPoint Vertex %34 "vertex_two_structs" %24 %28 %30 %32
OpMemberDecorate %13 0 Offset 0
OpMemberDecorate %13 1 Offset 16
OpMemberDecorate %14 0 Offset 0
OpMemberDecorate %14 1 Offset 4
OpMemberDecorate %14 2 Offset 8
OpDecorate %16 ArrayStride 4
OpMemberDecorate %18 0 Offset 0
OpMemberDecorate %19 0 Offset 0
OpDecorate %24 BuiltIn VertexIndex
OpDecorate %28 BuiltIn InstanceIndex
OpDecorate %30 BuiltIn Position
OpDecorate %32 BuiltIn PointSize
%2 = OpTypeVoid
%4 = OpTypeFloat 32
%3 = OpConstant %4 1.0
%6 = OpTypeInt 32 0
%5 = OpConstant %6 1
%7 = OpConstant %4 0.0
%9 = OpTypeInt 32 1
%8 = OpConstant %9 1
%10 = OpConstant %9 0
%11 = OpConstant %6 2
%12 = OpTypeVector %4 4
%13 = OpTypeStruct %12 %4
%14 = OpTypeStruct %4 %6 %4
%15 = OpTypeBool
%16 = OpTypeArray %6 %8
%17 = OpTypeVector %6 3
%18 = OpTypeStruct %6
%19 = OpTypeStruct %6
%21 = OpTypePointer Function %6
%25 = OpTypePointer Input %6
%24 = OpVariable %25 Input
%28 = OpVariable %25 Input
%31 = OpTypePointer Output %12
%30 = OpVariable %31 Output
%33 = OpTypePointer Output %4
%32 = OpVariable %33 Output
%35 = OpTypeFunction %2
%36 = OpTypePointer Workgroup %16
%34 = OpFunction %2 None %35
%22 = OpLabel
%20 = OpVariable %21 Function %11
%26 = OpLoad %6 %24
%23 = OpCompositeConstruct %18 %26
%29 = OpLoad %6 %28
%27 = OpCompositeConstruct %19 %29
OpStore %32 %3
OpBranch %37
%37 = OpLabel
%38 = OpCompositeExtract %6 %23 0
%39 = OpConvertUToF %4 %38
%40 = OpCompositeExtract %6 %27 0
%41 = OpConvertUToF %4 %40
%42 = OpLoad %6 %20
%43 = OpConvertUToF %4 %42
%44 = OpCompositeConstruct %12 %39 %41 %43 %7
OpStore %30 %44
OpReturn
OpFunctionEnd

16
tests/out/wgsl/interface.wgsl
View File

@ -9,6 +9,14 @@ struct FragmentOutput {
@location(0) color: f32;
};
struct Input1_ {
@builtin(vertex_index) index: u32;
};
struct Input2_ {
@builtin(instance_index) index: u32;
};
var<workgroup> output: array<u32,1>;
@stage(vertex)
@ -29,3 +37,11 @@ fn compute(@builtin(global_invocation_id) global_id: vec3<u32>, @builtin(local_i
output[0] = ((((global_id.x + local_id.x) + local_index) + wg_id.x) + num_wgs.x);
return;
}
@stage(vertex)
fn vertex_two_structs(in1_: Input1_, in2_: Input2_) -> @builtin(position) vec4<f32> {
var index: u32 = 2u;
let _e9: u32 = index;
return vec4<f32>(f32(in1_.index), f32(in2_.index), f32(_e9), 0.0);
}