From 4cbef59dd60c21deb463a2275466734dd428df48 Mon Sep 17 00:00:00 2001 From: Lachlan Sneff Date: Sat, 17 Apr 2021 01:14:07 -0400 Subject: [PATCH] Add struct instantiation support to wgsl-in (#507) * Add struct instantiation support to wgsl-in * Readd type matching and move struct test * rebase onto master and update tests * Update tests to upstream * Fix merge errors * Rebase onto master * [wgsl-in] delay composite type creation * Use the new WGSL construction syntax in the tests Co-authored-by: Dzmitry Malyshau --- src/front/wgsl/mod.rs | 59 +++++++---- src/front/wgsl/tests.rs | 18 ++++ tests/in/quad.wgsl | 5 +- tests/in/skybox.wgsl | 5 +- tests/out/quad.Vertex.glsl | 7 +- tests/out/quad.dot | 51 +++------ tests/out/quad.msl | 5 +- tests/out/quad.spvasm | 138 ++++++++++++------------- tests/out/skybox.Vertex.glsl | 7 +- tests/out/skybox.msl | 5 +- tests/out/skybox.spvasm | 194 +++++++++++++++++------------------ 11 files changed, 238 insertions(+), 256 deletions(-) diff --git a/src/front/wgsl/mod.rs b/src/front/wgsl/mod.rs index e9f21a2919..8b48534066 100644 --- a/src/front/wgsl/mod.rs +++ b/src/front/wgsl/mod.rs @@ -10,7 +10,7 @@ mod tests; use crate::{ arena::{Arena, Handle}, - proc::{ensure_block_returns, ResolveContext, ResolveError}, + proc::{ensure_block_returns, ResolveContext, ResolveError, TypeResolution}, FastHashMap, }; @@ -1111,14 +1111,19 @@ impl Parser { //TODO: resolve the duplicate call in `parse_singular_expression` expr } else { - let inner = self.parse_type_decl_impl( - lexer, - TypeAttributes::default(), - word, - ctx.types, - ctx.constants, - )?; - let kind = inner.scalar_kind(); + let ty_resolution = match self.lookup_type.get(word) { + Some(&handle) => TypeResolution::Handle(handle), + None => { + let inner = self.parse_type_decl_impl( + lexer, + TypeAttributes::default(), + word, + ctx.types, + ctx.constants, + )?; + TypeResolution::Value(inner) + } + }; lexer.expect(Token::Paren('('))?; let mut components = Vec::new(); @@ -1130,8 +1135,15 @@ impl Parser { } lexer.expect(Token::Paren(')'))?; let expr = if components.is_empty() { - let last_component_inner = ctx.resolve_type(last_component)?; - match (&inner, last_component_inner) { + // We can't use the `TypeInner` returned by this because + // `resolve_type` borrows context mutably. + // Use it to insert into the right maps, + // and then grab it again immutably. + ctx.resolve_type(last_component)?; + match ( + ty_resolution.inner_with(ctx.types), + ctx.typifier.get(last_component, ctx.types), + ) { ( &crate::TypeInner::Vector { size, .. }, &crate::TypeInner::Scalar { .. }, @@ -1140,19 +1152,23 @@ impl Parser { value: last_component, }, ( - &crate::TypeInner::Scalar { .. }, + &crate::TypeInner::Scalar { kind, .. }, &crate::TypeInner::Scalar { .. }, ) | ( - &crate::TypeInner::Matrix { .. }, - &crate::TypeInner::Matrix { .. }, - ) - | ( - &crate::TypeInner::Vector { .. }, + &crate::TypeInner::Vector { kind, .. }, &crate::TypeInner::Vector { .. }, ) => crate::Expression::As { expr: last_component, - kind: kind.ok_or(Error::BadTypeCast(word))?, + kind, + convert: true, + }, + ( + &crate::TypeInner::Matrix { .. }, + &crate::TypeInner::Matrix { .. }, + ) => crate::Expression::As { + expr: last_component, + kind: crate::ScalarKind::Float, convert: true, }, _ => { @@ -1160,8 +1176,13 @@ impl Parser { } } } else { + let ty = match ty_resolution { + TypeResolution::Handle(handle) => handle, + TypeResolution::Value(inner) => { + ctx.types.fetch_or_append(crate::Type { name: None, inner }) + } + }; components.push(last_component); - let ty = ctx.types.fetch_or_append(crate::Type { name: None, inner }); crate::Expression::Compose { ty, components } }; ctx.expressions.append(expr) diff --git a/src/front/wgsl/tests.rs b/src/front/wgsl/tests.rs index 96b96a7187..d7428b65b0 100644 --- a/src/front/wgsl/tests.rs +++ b/src/front/wgsl/tests.rs @@ -286,3 +286,21 @@ fn parse_pointers() { ) .unwrap(); } + +#[test] +fn parse_struct_instantiation() { + parse_str( + " + struct Foo { + a: f32; + b: vec3; + }; + + [[stage(fragment)]] + fn fs_main() { + var foo: Foo = Foo(0.0, vec3(0.0, 1.0, 42.0)); + } + ", + ) + .unwrap(); +} diff --git a/tests/in/quad.wgsl b/tests/in/quad.wgsl index fdaa9600c0..0cd0187a8f 100644 --- a/tests/in/quad.wgsl +++ b/tests/in/quad.wgsl @@ -8,10 +8,7 @@ struct VertexOutput { [[stage(vertex)]] fn main([[location(0)]] pos : vec2, [[location(1)]] uv : vec2) -> VertexOutput { - var out: VertexOutput; - out.uv = uv; - out.position = vec4(c_scale * pos, 0.0, 1.0); - return out; + return VertexOutput(uv, vec4(c_scale * pos, 0.0, 1.0)); } // fragment diff --git a/tests/in/skybox.wgsl b/tests/in/skybox.wgsl index 8123737928..d5cb051262 100644 --- a/tests/in/skybox.wgsl +++ b/tests/in/skybox.wgsl @@ -25,10 +25,7 @@ fn vs_main([[builtin(vertex_index)]] vertex_index: u32) -> VertexOutput { let inv_model_view = transpose(mat3x3(r_data.view.x.xyz, r_data.view.y.xyz, r_data.view.z.xyz)); let unprojected = r_data.proj_inv * pos; - var out: VertexOutput; - out.uv = inv_model_view * unprojected.xyz; - out.position = pos; - return out; + return VertexOutput(pos, inv_model_view * unprojected.xyz); } [[group(0), binding(1)]] diff --git a/tests/out/quad.Vertex.glsl b/tests/out/quad.Vertex.glsl index b35cf16eb6..ccd93c498f 100644 --- a/tests/out/quad.Vertex.glsl +++ b/tests/out/quad.Vertex.glsl @@ -14,11 +14,8 @@ layout(location = 0) out vec2 _vs2fs_location0; void main() { vec2 pos = _p2vs_location0; vec2 uv1 = _p2vs_location1; - VertexOutput out1; - out1.uv = uv1; - out1.position = vec4((1.2 * pos), 0.0, 1.0); - _vs2fs_location0 = out1.uv; - gl_Position = out1.position; + _vs2fs_location0 = VertexOutput(uv1, vec4((1.2 * pos), 0.0, 1.0)).uv; + gl_Position = VertexOutput(uv1, vec4((1.2 * pos), 0.0, 1.0)).position; return; } diff --git a/tests/out/quad.dot b/tests/out/quad.dot index 6e8c541646..60587fc8b1 100644 --- a/tests/out/quad.dot +++ b/tests/out/quad.dot @@ -7,50 +7,29 @@ digraph Module { subgraph cluster_ep0 { label="Vertex/'main'" node [ style=filled ] - ep0_l0 [ shape=hexagon label="[1] 'out'" ] ep0_e0 [ fillcolor="#ffffb3" label="[1] Constant" ] ep0_e1 [ color="#8dd3c7" label="[2] Argument[0]" ] ep0_e2 [ color="#8dd3c7" label="[3] Argument[1]" ] - ep0_e3 [ color="#8dd3c7" label="[4] Local" ] - ep0_l0 -> ep0_e3 - ep0_e4 [ color="#8dd3c7" label="[5] AccessIndex[0]" ] - ep0_e3 -> ep0_e4 [ label="base" ] - ep0_e5 [ color="#8dd3c7" label="[6] AccessIndex[1]" ] - ep0_e3 -> ep0_e5 [ label="base" ] - ep0_e6 [ color="#fdb462" label="[7] Multiply" ] - ep0_e1 -> ep0_e6 [ label="right" ] - ep0_e0 -> ep0_e6 [ label="left" ] - ep0_e7 [ fillcolor="#ffffb3" label="[8] Constant" ] - ep0_e8 [ fillcolor="#ffffb3" label="[9] Constant" ] - ep0_e9 [ color="#bebada" label="[10] Compose" ] - { ep0_e6 ep0_e7 ep0_e8 } -> ep0_e9 - ep0_e10 [ color="#fb8072" label="[11] Load" ] - ep0_e3 -> ep0_e10 [ label="pointer" ] + ep0_e3 [ color="#fdb462" label="[4] Multiply" ] + ep0_e1 -> ep0_e3 [ label="right" ] + ep0_e0 -> ep0_e3 [ label="left" ] + ep0_e4 [ fillcolor="#ffffb3" label="[5] Constant" ] + ep0_e5 [ fillcolor="#ffffb3" label="[6] Constant" ] + ep0_e6 [ color="#bebada" label="[7] Compose" ] + { ep0_e3 ep0_e4 ep0_e5 } -> ep0_e6 + ep0_e7 [ color="#bebada" label="[8] Compose" ] + { ep0_e2 ep0_e6 } -> ep0_e7 ep0_s0 [ shape=square label="Root" ] ep0_s1 [ shape=square label="Emit" ] - ep0_s2 [ shape=square label="Store" ] - ep0_s3 [ shape=square label="Emit" ] - ep0_s4 [ shape=square label="Emit" ] - ep0_s5 [ shape=square label="Store" ] - ep0_s6 [ shape=square label="Emit" ] - ep0_s7 [ shape=square label="Return" ] + ep0_s2 [ shape=square label="Emit" ] + ep0_s3 [ shape=square label="Return" ] ep0_s0 -> ep0_s1 [ arrowhead=tee label="" ] ep0_s1 -> ep0_s2 [ arrowhead=tee label="" ] ep0_s2 -> ep0_s3 [ arrowhead=tee label="" ] - ep0_s3 -> ep0_s4 [ arrowhead=tee label="" ] - ep0_s4 -> ep0_s5 [ arrowhead=tee label="" ] - ep0_s5 -> ep0_s6 [ arrowhead=tee label="" ] - ep0_s6 -> ep0_s7 [ arrowhead=tee label="" ] - ep0_e2 -> ep0_s2 [ label="value" ] - ep0_e9 -> ep0_s5 [ label="value" ] - ep0_e10 -> ep0_s7 [ label="value" ] - ep0_s1 -> ep0_e4 [ style=dotted ] - ep0_s2 -> ep0_e4 [ style=dotted ] - ep0_s3 -> ep0_e5 [ style=dotted ] - ep0_s3 -> ep0_e6 [ style=dotted ] - ep0_s4 -> ep0_e9 [ style=dotted ] - ep0_s5 -> ep0_e5 [ style=dotted ] - ep0_s6 -> ep0_e10 [ style=dotted ] + ep0_e7 -> ep0_s3 [ label="value" ] + ep0_s1 -> ep0_e3 [ style=dotted ] + ep0_s2 -> ep0_e6 [ style=dotted ] + ep0_s2 -> ep0_e7 [ style=dotted ] } subgraph cluster_ep1 { label="Fragment/'main'" diff --git a/tests/out/quad.msl b/tests/out/quad.msl index 17f225b0dc..e4ea4e40bf 100644 --- a/tests/out/quad.msl +++ b/tests/out/quad.msl @@ -20,10 +20,7 @@ vertex main1Output main1( ) { const auto pos = varyings.pos; const auto uv1 = varyings.uv1; - VertexOutput out; - out.uv = uv1; - out.position = metal::float4(c_scale * pos, 0.0, 1.0); - const auto _tmp = out; + const auto _tmp = VertexOutput {uv1, metal::float4(c_scale * pos, 0.0, 1.0)}; return main1Output { _tmp.uv, _tmp.position }; } diff --git a/tests/out/quad.spvasm b/tests/out/quad.spvasm index 7f5af3c54e..d15def3de2 100644 --- a/tests/out/quad.spvasm +++ b/tests/out/quad.spvasm @@ -1,13 +1,13 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 65 +; Bound: 58 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %28 "main" %19 %22 %24 %26 -OpEntryPoint Fragment %51 "main" %48 %50 -OpExecutionMode %51 OriginUpperLeft +OpEntryPoint Vertex %26 "main" %17 %20 %22 %24 +OpEntryPoint Fragment %44 "main" %41 %43 +OpExecutionMode %44 OriginUpperLeft OpSource GLSL 450 OpName %3 "c_scale" OpName %9 "VertexOutput" @@ -15,26 +15,25 @@ OpMemberName %9 0 "uv" OpMemberName %9 1 "position" OpName %12 "u_texture" OpName %14 "u_sampler" -OpName %16 "out" -OpName %19 "pos" +OpName %17 "pos" +OpName %20 "uv" OpName %22 "uv" -OpName %24 "uv" -OpName %26 "position" -OpName %28 "main" -OpName %48 "uv" -OpName %51 "main" +OpName %24 "position" +OpName %26 "main" +OpName %41 "uv" +OpName %44 "main" OpMemberDecorate %9 0 Offset 0 OpMemberDecorate %9 1 Offset 16 OpDecorate %12 DescriptorSet 0 OpDecorate %12 Binding 0 OpDecorate %14 DescriptorSet 0 OpDecorate %14 Binding 1 -OpDecorate %19 Location 0 -OpDecorate %22 Location 1 -OpDecorate %24 Location 0 -OpDecorate %26 BuiltIn Position -OpDecorate %48 Location 0 -OpDecorate %50 Location 0 +OpDecorate %17 Location 0 +OpDecorate %20 Location 1 +OpDecorate %22 Location 0 +OpDecorate %24 BuiltIn Position +OpDecorate %41 Location 0 +OpDecorate %43 Location 0 %2 = OpTypeVoid %4 = OpTypeFloat 32 %3 = OpConstant %4 1.2 @@ -49,67 +48,58 @@ OpDecorate %50 Location 0 %12 = OpVariable %13 UniformConstant %15 = OpTypePointer UniformConstant %11 %14 = OpVariable %15 UniformConstant -%17 = OpTypePointer Function %9 -%20 = OpTypePointer Input %7 -%19 = OpVariable %20 Input -%22 = OpVariable %20 Input -%25 = OpTypePointer Output %7 +%18 = OpTypePointer Input %7 +%17 = OpVariable %18 Input +%20 = OpVariable %18 Input +%23 = OpTypePointer Output %7 +%22 = OpVariable %23 Output +%25 = OpTypePointer Output %8 %24 = OpVariable %25 Output -%27 = OpTypePointer Output %8 -%26 = OpVariable %27 Output -%29 = OpTypeFunction %2 -%31 = OpTypePointer Function %7 -%32 = OpTypeInt 32 1 -%33 = OpConstant %32 0 -%35 = OpTypePointer Function %8 -%38 = OpConstant %32 1 -%44 = OpTypePointer Output %4 -%48 = OpVariable %20 Input -%50 = OpVariable %27 Output -%55 = OpTypeSampledImage %10 -%59 = OpTypeBool -%28 = OpFunction %2 None %29 -%18 = OpLabel -%16 = OpVariable %17 Function -%21 = OpLoad %7 %19 -%23 = OpLoad %7 %22 -OpBranch %30 -%30 = OpLabel -%34 = OpAccessChain %31 %16 %33 -OpStore %34 %23 -%36 = OpVectorTimesScalar %7 %21 %3 -%37 = OpCompositeConstruct %8 %36 %5 %6 -%39 = OpAccessChain %35 %16 %38 -OpStore %39 %37 -%40 = OpLoad %9 %16 -%41 = OpCompositeExtract %7 %40 0 -OpStore %24 %41 -%42 = OpCompositeExtract %8 %40 1 -OpStore %26 %42 -%43 = OpAccessChain %44 %26 %38 -%45 = OpLoad %4 %43 -%46 = OpFNegate %4 %45 -OpStore %43 %46 +%27 = OpTypeFunction %2 +%35 = OpTypePointer Output %4 +%36 = OpTypeInt 32 1 +%37 = OpConstant %36 1 +%41 = OpVariable %18 Input +%43 = OpVariable %25 Output +%48 = OpTypeSampledImage %10 +%52 = OpTypeBool +%26 = OpFunction %2 None %27 +%16 = OpLabel +%19 = OpLoad %7 %17 +%21 = OpLoad %7 %20 +OpBranch %28 +%28 = OpLabel +%29 = OpVectorTimesScalar %7 %19 %3 +%30 = OpCompositeConstruct %8 %29 %5 %6 +%31 = OpCompositeConstruct %9 %21 %30 +%32 = OpCompositeExtract %7 %31 0 +OpStore %22 %32 +%33 = OpCompositeExtract %8 %31 1 +OpStore %24 %33 +%34 = OpAccessChain %35 %24 %37 +%38 = OpLoad %4 %34 +%39 = OpFNegate %4 %38 +OpStore %34 %39 OpReturn OpFunctionEnd -%51 = OpFunction %2 None %29 +%44 = OpFunction %2 None %27 +%40 = OpLabel +%42 = OpLoad %7 %41 +%45 = OpLoad %10 %12 +%46 = OpLoad %11 %14 +OpBranch %47 %47 = OpLabel -%49 = OpLoad %7 %48 -%52 = OpLoad %10 %12 -%53 = OpLoad %11 %14 -OpBranch %54 -%54 = OpLabel -%56 = OpSampledImage %55 %52 %53 -%57 = OpImageSampleImplicitLod %8 %56 %49 -%58 = OpCompositeExtract %4 %57 3 -%60 = OpFOrdEqual %59 %58 %5 -OpSelectionMerge %61 None -OpBranchConditional %60 %62 %61 -%62 = OpLabel +%49 = OpSampledImage %48 %45 %46 +%50 = OpImageSampleImplicitLod %8 %49 %42 +%51 = OpCompositeExtract %4 %50 3 +%53 = OpFOrdEqual %52 %51 %5 +OpSelectionMerge %54 None +OpBranchConditional %53 %55 %54 +%55 = OpLabel OpKill -%61 = OpLabel -%63 = OpCompositeExtract %4 %57 3 -%64 = OpVectorTimesScalar %8 %57 %63 -OpStore %50 %64 +%54 = OpLabel +%56 = OpCompositeExtract %4 %50 3 +%57 = OpVectorTimesScalar %8 %50 %56 +OpStore %43 %57 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/skybox.Vertex.glsl b/tests/out/skybox.Vertex.glsl index b30eced325..cc327f6a3f 100644 --- a/tests/out/skybox.Vertex.glsl +++ b/tests/out/skybox.Vertex.glsl @@ -18,15 +18,12 @@ void main() { uint vertex_index = uint(gl_VertexID); int tmp1_; int tmp2_; - VertexOutput out1; tmp1_ = (int(vertex_index) / 2); tmp2_ = (int(vertex_index) & 1); vec4 _expr24 = vec4(((float(tmp1_) * 4.0) - 1.0), ((float(tmp2_) * 4.0) - 1.0), 0.0, 1.0); vec4 _expr50 = (_group_0_binding_0.proj_inv * _expr24); - out1.uv = (transpose(mat3x3(vec3(_group_0_binding_0.view[0][0], _group_0_binding_0.view[0][1], _group_0_binding_0.view[0][2]), vec3(_group_0_binding_0.view[1][0], _group_0_binding_0.view[1][1], _group_0_binding_0.view[1][2]), vec3(_group_0_binding_0.view[2][0], _group_0_binding_0.view[2][1], _group_0_binding_0.view[2][2]))) * vec3(_expr50[0], _expr50[1], _expr50[2])); - out1.position = _expr24; - gl_Position = out1.position; - _vs2fs_location0 = out1.uv; + gl_Position = VertexOutput(_expr24, (transpose(mat3x3(vec3(_group_0_binding_0.view[0][0], _group_0_binding_0.view[0][1], _group_0_binding_0.view[0][2]), vec3(_group_0_binding_0.view[1][0], _group_0_binding_0.view[1][1], _group_0_binding_0.view[1][2]), vec3(_group_0_binding_0.view[2][0], _group_0_binding_0.view[2][1], _group_0_binding_0.view[2][2]))) * vec3(_expr50[0], _expr50[1], _expr50[2]))).position; + _vs2fs_location0 = VertexOutput(_expr24, (transpose(mat3x3(vec3(_group_0_binding_0.view[0][0], _group_0_binding_0.view[0][1], _group_0_binding_0.view[0][2]), vec3(_group_0_binding_0.view[1][0], _group_0_binding_0.view[1][1], _group_0_binding_0.view[1][2]), vec3(_group_0_binding_0.view[2][0], _group_0_binding_0.view[2][1], _group_0_binding_0.view[2][2]))) * vec3(_expr50[0], _expr50[1], _expr50[2]))).uv; return; } diff --git a/tests/out/skybox.msl b/tests/out/skybox.msl index bd43ac6f62..efdff96d50 100644 --- a/tests/out/skybox.msl +++ b/tests/out/skybox.msl @@ -22,14 +22,11 @@ vertex vs_mainOutput vs_main( ) { int tmp1_; int tmp2_; - VertexOutput out; tmp1_ = static_cast(vertex_index) / 2; tmp2_ = static_cast(vertex_index) & 1; metal::float4 _e24 = metal::float4((static_cast(tmp1_) * 4.0) - 1.0, (static_cast(tmp2_) * 4.0) - 1.0, 0.0, 1.0); metal::float4 _e50 = r_data.proj_inv * _e24; - out.uv = metal::transpose(metal::float3x3(metal::float3(r_data.view[0].x, r_data.view[0].y, r_data.view[0].z), metal::float3(r_data.view[1].x, r_data.view[1].y, r_data.view[1].z), metal::float3(r_data.view[2].x, r_data.view[2].y, r_data.view[2].z))) * metal::float3(_e50.x, _e50.y, _e50.z); - out.position = _e24; - const auto _tmp = out; + const auto _tmp = VertexOutput {_e24, metal::transpose(metal::float3x3(metal::float3(r_data.view[0].x, r_data.view[0].y, r_data.view[0].z), metal::float3(r_data.view[1].x, r_data.view[1].y, r_data.view[1].z), metal::float3(r_data.view[2].x, r_data.view[2].y, r_data.view[2].z))) * metal::float3(_e50.x, _e50.y, _e50.z)}; return vs_mainOutput { _tmp.position, _tmp.uv }; } diff --git a/tests/out/skybox.spvasm b/tests/out/skybox.spvasm index 7455b12941..c51900db33 100644 --- a/tests/out/skybox.spvasm +++ b/tests/out/skybox.spvasm @@ -1,13 +1,13 @@ ; SPIR-V ; Version: 1.0 ; Generator: rspirv -; Bound: 111 +; Bound: 105 OpCapability Shader %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint Vertex %38 "vs_main" %31 %34 %36 -OpEntryPoint Fragment %103 "fs_main" %96 %99 %102 -OpExecutionMode %103 OriginUpperLeft +OpEntryPoint Vertex %36 "vs_main" %29 %32 %34 +OpEntryPoint Fragment %97 "fs_main" %90 %93 %96 +OpExecutionMode %97 OriginUpperLeft OpMemberDecorate %12 0 Offset 0 OpMemberDecorate %12 1 Offset 16 OpDecorate %14 Block @@ -23,12 +23,12 @@ OpDecorate %21 DescriptorSet 0 OpDecorate %21 Binding 1 OpDecorate %23 DescriptorSet 0 OpDecorate %23 Binding 2 -OpDecorate %31 BuiltIn VertexIndex -OpDecorate %34 BuiltIn Position -OpDecorate %36 Location 0 -OpDecorate %96 BuiltIn FragCoord -OpDecorate %99 Location 0 -OpDecorate %102 Location 0 +OpDecorate %29 BuiltIn VertexIndex +OpDecorate %32 BuiltIn Position +OpDecorate %34 Location 0 +OpDecorate %90 BuiltIn FragCoord +OpDecorate %93 Location 0 +OpDecorate %96 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 2 @@ -53,101 +53,93 @@ OpDecorate %102 Location 0 %24 = OpTypePointer UniformConstant %18 %23 = OpVariable %24 UniformConstant %26 = OpTypePointer Function %4 -%29 = OpTypePointer Function %12 -%32 = OpTypePointer Input %15 -%31 = OpVariable %32 Input -%35 = OpTypePointer Output %10 +%30 = OpTypePointer Input %15 +%29 = OpVariable %30 Input +%33 = OpTypePointer Output %10 +%32 = OpVariable %33 Output +%35 = OpTypePointer Output %11 %34 = OpVariable %35 Output -%37 = OpTypePointer Output %11 -%36 = OpVariable %37 Output -%39 = OpTypeFunction %2 -%54 = OpTypePointer Uniform %13 -%78 = OpConstant %4 0 -%82 = OpTypePointer Function %11 -%89 = OpTypePointer Function %10 -%97 = OpTypePointer Input %10 -%96 = OpVariable %97 Input -%100 = OpTypePointer Input %11 -%99 = OpVariable %100 Input -%102 = OpVariable %35 Output -%108 = OpTypeSampledImage %17 -%38 = OpFunction %2 None %39 -%30 = OpLabel +%37 = OpTypeFunction %2 +%52 = OpTypePointer Uniform %13 +%76 = OpConstant %4 0 +%91 = OpTypePointer Input %10 +%90 = OpVariable %91 Input +%94 = OpTypePointer Input %11 +%93 = OpVariable %94 Input +%96 = OpVariable %33 Output +%102 = OpTypeSampledImage %17 +%36 = OpFunction %2 None %37 +%28 = OpLabel %25 = OpVariable %26 Function %27 = OpVariable %26 Function -%28 = OpVariable %29 Function -%33 = OpLoad %15 %31 -OpBranch %40 -%40 = OpLabel -%41 = OpBitcast %4 %33 -%42 = OpSDiv %4 %41 %3 -OpStore %25 %42 -%43 = OpBitcast %4 %33 -%44 = OpBitwiseAnd %4 %43 %5 -OpStore %27 %44 -%45 = OpLoad %4 %25 -%46 = OpConvertSToF %7 %45 -%47 = OpFMul %7 %46 %6 -%48 = OpFSub %7 %47 %8 -%49 = OpLoad %4 %27 -%50 = OpConvertSToF %7 %49 -%51 = OpFMul %7 %50 %6 -%52 = OpFSub %7 %51 %8 -%53 = OpCompositeConstruct %10 %48 %52 %9 %8 -%55 = OpAccessChain %54 %19 %5 -%56 = OpLoad %13 %55 -%57 = OpCompositeExtract %10 %56 0 -%58 = OpCompositeExtract %7 %57 0 -%59 = OpCompositeExtract %7 %57 1 -%60 = OpCompositeExtract %7 %57 2 -%61 = OpCompositeConstruct %11 %58 %59 %60 -%62 = OpAccessChain %54 %19 %5 -%63 = OpLoad %13 %62 -%64 = OpCompositeExtract %10 %63 1 -%65 = OpCompositeExtract %7 %64 0 -%66 = OpCompositeExtract %7 %64 1 -%67 = OpCompositeExtract %7 %64 2 -%68 = OpCompositeConstruct %11 %65 %66 %67 -%69 = OpAccessChain %54 %19 %5 -%70 = OpLoad %13 %69 -%71 = OpCompositeExtract %10 %70 2 -%72 = OpCompositeExtract %7 %71 0 -%73 = OpCompositeExtract %7 %71 1 -%74 = OpCompositeExtract %7 %71 2 -%75 = OpCompositeConstruct %11 %72 %73 %74 -%76 = OpCompositeConstruct %16 %61 %68 %75 -%77 = OpTranspose %16 %76 -%79 = OpAccessChain %54 %19 %78 -%80 = OpLoad %13 %79 -%81 = OpMatrixTimesVector %10 %80 %53 -%83 = OpCompositeExtract %7 %81 0 -%84 = OpCompositeExtract %7 %81 1 -%85 = OpCompositeExtract %7 %81 2 -%86 = OpCompositeConstruct %11 %83 %84 %85 -%87 = OpMatrixTimesVector %11 %77 %86 -%88 = OpAccessChain %82 %28 %5 -OpStore %88 %87 -%90 = OpAccessChain %89 %28 %78 -OpStore %90 %53 -%91 = OpLoad %12 %28 -%92 = OpCompositeExtract %10 %91 0 -OpStore %34 %92 -%93 = OpCompositeExtract %11 %91 1 -OpStore %36 %93 +%31 = OpLoad %15 %29 +OpBranch %38 +%38 = OpLabel +%39 = OpBitcast %4 %31 +%40 = OpSDiv %4 %39 %3 +OpStore %25 %40 +%41 = OpBitcast %4 %31 +%42 = OpBitwiseAnd %4 %41 %5 +OpStore %27 %42 +%43 = OpLoad %4 %25 +%44 = OpConvertSToF %7 %43 +%45 = OpFMul %7 %44 %6 +%46 = OpFSub %7 %45 %8 +%47 = OpLoad %4 %27 +%48 = OpConvertSToF %7 %47 +%49 = OpFMul %7 %48 %6 +%50 = OpFSub %7 %49 %8 +%51 = OpCompositeConstruct %10 %46 %50 %9 %8 +%53 = OpAccessChain %52 %19 %5 +%54 = OpLoad %13 %53 +%55 = OpCompositeExtract %10 %54 0 +%56 = OpCompositeExtract %7 %55 0 +%57 = OpCompositeExtract %7 %55 1 +%58 = OpCompositeExtract %7 %55 2 +%59 = OpCompositeConstruct %11 %56 %57 %58 +%60 = OpAccessChain %52 %19 %5 +%61 = OpLoad %13 %60 +%62 = OpCompositeExtract %10 %61 1 +%63 = OpCompositeExtract %7 %62 0 +%64 = OpCompositeExtract %7 %62 1 +%65 = OpCompositeExtract %7 %62 2 +%66 = OpCompositeConstruct %11 %63 %64 %65 +%67 = OpAccessChain %52 %19 %5 +%68 = OpLoad %13 %67 +%69 = OpCompositeExtract %10 %68 2 +%70 = OpCompositeExtract %7 %69 0 +%71 = OpCompositeExtract %7 %69 1 +%72 = OpCompositeExtract %7 %69 2 +%73 = OpCompositeConstruct %11 %70 %71 %72 +%74 = OpCompositeConstruct %16 %59 %66 %73 +%75 = OpTranspose %16 %74 +%77 = OpAccessChain %52 %19 %76 +%78 = OpLoad %13 %77 +%79 = OpMatrixTimesVector %10 %78 %51 +%80 = OpCompositeExtract %7 %79 0 +%81 = OpCompositeExtract %7 %79 1 +%82 = OpCompositeExtract %7 %79 2 +%83 = OpCompositeConstruct %11 %80 %81 %82 +%84 = OpMatrixTimesVector %11 %75 %83 +%85 = OpCompositeConstruct %12 %51 %84 +%86 = OpCompositeExtract %10 %85 0 +OpStore %32 %86 +%87 = OpCompositeExtract %11 %85 1 +OpStore %34 %87 OpReturn OpFunctionEnd -%103 = OpFunction %2 None %39 -%94 = OpLabel -%98 = OpLoad %10 %96 -%101 = OpLoad %11 %99 -%95 = OpCompositeConstruct %12 %98 %101 -%104 = OpLoad %17 %21 -%105 = OpLoad %18 %23 -OpBranch %106 -%106 = OpLabel -%107 = OpCompositeExtract %11 %95 1 -%109 = OpSampledImage %108 %104 %105 -%110 = OpImageSampleImplicitLod %10 %109 %107 -OpStore %102 %110 +%97 = OpFunction %2 None %37 +%88 = OpLabel +%92 = OpLoad %10 %90 +%95 = OpLoad %11 %93 +%89 = OpCompositeConstruct %12 %92 %95 +%98 = OpLoad %17 %21 +%99 = OpLoad %18 %23 +OpBranch %100 +%100 = OpLabel +%101 = OpCompositeExtract %11 %89 1 +%103 = OpSampledImage %102 %98 %99 +%104 = OpImageSampleImplicitLod %10 %103 %101 +OpStore %96 %104 OpReturn OpFunctionEnd \ No newline at end of file