From c84aa775790682f2e4ed28440bc681b09808d0cb Mon Sep 17 00:00:00 2001 From: Jim Blandy Date: Tue, 8 Mar 2022 07:07:30 -0800 Subject: [PATCH] [msl-out] Properly rename entry point arguments for struct members. (#1766) --- src/back/msl/writer.rs | 92 ++++++++++--- tests/in/interface.wgsl | 14 ++ tests/out/hlsl/interface.hlsl | 16 +++ tests/out/hlsl/interface.hlsl.config | 2 +- tests/out/msl/interface.msl | 24 ++++ tests/out/spv/interface.compute.spvasm | 103 ++++++++------- tests/out/spv/interface.fragment.spvasm | 123 +++++++++--------- tests/out/spv/interface.vertex.spvasm | 101 +++++++------- .../spv/interface.vertex_two_structs.spvasm | 68 ++++++++++ tests/out/wgsl/interface.wgsl | 16 +++ 10 files changed, 383 insertions(+), 176 deletions(-) create mode 100644 tests/out/spv/interface.vertex_two_structs.spvasm diff --git a/src/back/msl/writer.rs b/src/back/msl/writer.rs index c04abe01a9..076a37bd15 100644 --- a/src/back/msl/writer.rs +++ b/src/back/msl/writer.rs @@ -3288,9 +3288,6 @@ impl Writer { 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 Writer { } }; - 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 `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 Writer { writeln!(self.out, "}};")?; } + // Define a struct type named for the return value, if any, named + // `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 Writer { } 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 Writer { )?; 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 Writer { 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 Writer { 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 Writer { } } - // 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 Writer { 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, ", ")?; } diff --git a/tests/in/interface.wgsl b/tests/in/interface.wgsl index f5946c13a2..7375295521 100644 --- a/tests/in/interface.wgsl +++ b/tests/in/interface.wgsl @@ -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 { + var index = 2u; + return vec4(f32(in1.index), f32(in2.index), f32(index), 0.0); +} diff --git a/tests/out/hlsl/interface.hlsl b/tests/out/hlsl/interface.hlsl index 99ed31acd5..9ccafc3405 100644 --- a/tests/out/hlsl/interface.hlsl +++ b/tests/out/hlsl/interface.hlsl @@ -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); +} diff --git a/tests/out/hlsl/interface.hlsl.config b/tests/out/hlsl/interface.hlsl.config index cec1fc2b98..d7e51aecf7 100644 --- a/tests/out/hlsl/interface.hlsl.config +++ b/tests/out/hlsl/interface.hlsl.config @@ -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 ) diff --git a/tests/out/msl/interface.msl b/tests/out/msl/interface.msl index dbc6871984..00789692cc 100644 --- a/tests/out/msl/interface.msl +++ b/tests/out/msl/interface.msl @@ -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(in1_.index), static_cast(in2_.index), static_cast(_e9), 0.0), 1.0 }; +} diff --git a/tests/out/spv/interface.compute.spvasm b/tests/out/spv/interface.compute.spvasm index 91397f1593..6c1dac372d 100644 --- a/tests/out/spv/interface.compute.spvasm +++ b/tests/out/spv/interface.compute.spvasm @@ -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 \ No newline at end of file diff --git a/tests/out/spv/interface.fragment.spvasm b/tests/out/spv/interface.fragment.spvasm index 5e7b97af1f..60e270e1c6 100644 --- a/tests/out/spv/interface.fragment.spvasm +++ b/tests/out/spv/interface.fragment.spvasm @@ -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 \ No newline at end of file diff --git a/tests/out/spv/interface.vertex.spvasm b/tests/out/spv/interface.vertex.spvasm index c51a4d563b..e42d243e9c 100644 --- a/tests/out/spv/interface.vertex.spvasm +++ b/tests/out/spv/interface.vertex.spvasm @@ -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 \ No newline at end of file diff --git a/tests/out/spv/interface.vertex_two_structs.spvasm b/tests/out/spv/interface.vertex_two_structs.spvasm new file mode 100644 index 0000000000..919bbeef6d --- /dev/null +++ b/tests/out/spv/interface.vertex_two_structs.spvasm @@ -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 \ No newline at end of file diff --git a/tests/out/wgsl/interface.wgsl b/tests/out/wgsl/interface.wgsl index 7d40bb321b..ad24717f36 100644 --- a/tests/out/wgsl/interface.wgsl +++ b/tests/out/wgsl/interface.wgsl @@ -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 output: array; @stage(vertex) @@ -29,3 +37,11 @@ fn compute(@builtin(global_invocation_id) global_id: vec3, @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 { + var index: u32 = 2u; + + let _e9: u32 = index; + return vec4(f32(in1_.index), f32(in2_.index), f32(_e9), 0.0); +}