From bf4e62b1ac78759705f544f71bf0b43b654a163c Mon Sep 17 00:00:00 2001 From: JMS55 <47158642+JMS55@users.noreply.github.com> Date: Mon, 5 Dec 2022 05:03:11 -0500 Subject: [PATCH] Fix textureGather on texture_2d (#2138) * Fix textureGather on texture_2d * Add textureGather u32/i32 tests * Update src/valid/expression.rs Co-authored-by: Teodor Tanasoaia <28601907+teoxoy@users.noreply.github.com> * Fix formatting * undo analyzer change Co-authored-by: Teodor Tanasoaia <28601907+teoxoy@users.noreply.github.com> --- src/valid/expression.rs | 4 + tests/in/image.wgsl | 21 +- tests/out/glsl/image.gather.Fragment.glsl | 9 +- tests/out/glsl/image.queries.Vertex.glsl | 28 +- tests/out/hlsl/image.hlsl | 17 +- tests/out/msl/image.msl | 7 +- tests/out/spv/image.spvasm | 741 +++++++++++----------- tests/out/wgsl/image.wgsl | 17 +- 8 files changed, 452 insertions(+), 392 deletions(-) diff --git a/src/valid/expression.rs b/src/valid/expression.rs index eb95c0e095..20a6237d97 100644 --- a/src/valid/expression.rs +++ b/src/valid/expression.rs @@ -384,6 +384,10 @@ impl super::Validator { kind: crate::ScalarKind::Float, multi: false, } => false, + crate::ImageClass::Sampled { + kind: crate::ScalarKind::Uint | crate::ScalarKind::Sint, + multi: false, + } if gather.is_some() => false, crate::ImageClass::Depth { multi: false } => true, _ => return Err(ExpressionError::InvalidImageClass(class)), }; diff --git a/tests/in/image.wgsl b/tests/in/image.wgsl index e1fb11a40f..8ece6d7a0e 100644 --- a/tests/in/image.wgsl +++ b/tests/in/image.wgsl @@ -45,14 +45,18 @@ var image_1d: texture_1d; @group(0) @binding(1) var image_2d: texture_2d; @group(0) @binding(2) -var image_2d_array: texture_2d_array; +var image_2d_u32: texture_2d; @group(0) @binding(3) -var image_cube: texture_cube; +var image_2d_i32: texture_2d; @group(0) @binding(4) -var image_cube_array: texture_cube_array; +var image_2d_array: texture_2d_array; @group(0) @binding(5) -var image_3d: texture_3d; +var image_cube: texture_cube; @group(0) @binding(6) +var image_cube_array: texture_cube_array; +@group(0) @binding(7) +var image_3d: texture_3d; +@group(0) @binding(8) var image_aa: texture_multisampled_2d; @vertex @@ -71,7 +75,7 @@ fn queries() -> @builtin(position) vec4 { let dim_3d_lod = textureDimensions(image_3d, 1); let dim_2s_ms = textureDimensions(image_aa); - let sum = dim_1d + dim_2d.y + dim_2d_lod.y + dim_2d_array.y + dim_2d_array_lod.y + + let sum = dim_1d + dim_2d.y + dim_2d_lod.y + dim_2d_array.y + dim_2d_array_lod.y + dim_cube.y + dim_cube_lod.y + dim_cube_array.y + dim_cube_array_lod.y + dim_3d.z + dim_3d_lod.z; return vec4(f32(sum)); @@ -134,7 +138,12 @@ fn gather() -> @location(0) vec4 { let s2d_offset = textureGather(3, image_2d, sampler_reg, tc, vec2(3, 1)); let s2d_depth = textureGatherCompare(image_2d_depth, sampler_cmp, tc, dref); let s2d_depth_offset = textureGatherCompare(image_2d_depth, sampler_cmp, tc, dref, vec2(3, 1)); - return s2d + s2d_offset + s2d_depth + s2d_depth_offset; + + let u = textureGather(0, image_2d_u32, sampler_reg, tc); + let i = textureGather(0, image_2d_i32, sampler_reg, tc); + let f = vec4(u) + vec4(i); + + return s2d + s2d_offset + s2d_depth + s2d_depth_offset + f; } @fragment diff --git a/tests/out/glsl/image.gather.Fragment.glsl b/tests/out/glsl/image.gather.Fragment.glsl index 5a44544a27..c7c2fc5348 100644 --- a/tests/out/glsl/image.gather.Fragment.glsl +++ b/tests/out/glsl/image.gather.Fragment.glsl @@ -6,6 +6,10 @@ precision highp int; uniform highp sampler2D _group_0_binding_1_fs; +uniform highp usampler2D _group_0_binding_2_fs; + +uniform highp isampler2D _group_0_binding_3_fs; + uniform highp sampler2DShadow _group_1_binding_2_fs; layout(location = 0) out vec4 _fs2p_location0; @@ -16,7 +20,10 @@ void main() { vec4 s2d_offset = textureGatherOffset(_group_0_binding_1_fs, vec2(tc), ivec2(3, 1), 3); vec4 s2d_depth = textureGather(_group_1_binding_2_fs, vec2(tc), 0.5); vec4 s2d_depth_offset = textureGatherOffset(_group_1_binding_2_fs, vec2(tc), 0.5, ivec2(3, 1)); - _fs2p_location0 = (((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset); + uvec4 u = textureGather(_group_0_binding_2_fs, vec2(tc), 0); + ivec4 i = textureGather(_group_0_binding_3_fs, vec2(tc), 0); + vec4 f = (vec4(u) + vec4(i)); + _fs2p_location0 = ((((s2d + s2d_offset) + s2d_depth) + s2d_depth_offset) + f); return; } diff --git a/tests/out/glsl/image.queries.Vertex.glsl b/tests/out/glsl/image.queries.Vertex.glsl index dfafa041b6..d416c73cf8 100644 --- a/tests/out/glsl/image.queries.Vertex.glsl +++ b/tests/out/glsl/image.queries.Vertex.glsl @@ -8,15 +8,15 @@ uniform highp sampler2D _group_0_binding_0_vs; uniform highp sampler2D _group_0_binding_1_vs; -uniform highp sampler2DArray _group_0_binding_2_vs; +uniform highp sampler2DArray _group_0_binding_4_vs; -uniform highp samplerCube _group_0_binding_3_vs; +uniform highp samplerCube _group_0_binding_5_vs; -uniform highp samplerCubeArray _group_0_binding_4_vs; +uniform highp samplerCubeArray _group_0_binding_6_vs; -uniform highp sampler3D _group_0_binding_5_vs; +uniform highp sampler3D _group_0_binding_7_vs; -uniform highp sampler2DMS _group_0_binding_6_vs; +uniform highp sampler2DMS _group_0_binding_8_vs; void main() { @@ -24,15 +24,15 @@ void main() { int dim_1d_lod = textureSize(_group_0_binding_0_vs, int(dim_1d)).x; ivec2 dim_2d = textureSize(_group_0_binding_1_vs, 0).xy; ivec2 dim_2d_lod = textureSize(_group_0_binding_1_vs, 1).xy; - ivec2 dim_2d_array = textureSize(_group_0_binding_2_vs, 0).xy; - ivec2 dim_2d_array_lod = textureSize(_group_0_binding_2_vs, 1).xy; - ivec2 dim_cube = textureSize(_group_0_binding_3_vs, 0).xy; - ivec2 dim_cube_lod = textureSize(_group_0_binding_3_vs, 1).xy; - ivec2 dim_cube_array = textureSize(_group_0_binding_4_vs, 0).xy; - ivec2 dim_cube_array_lod = textureSize(_group_0_binding_4_vs, 1).xy; - ivec3 dim_3d = textureSize(_group_0_binding_5_vs, 0).xyz; - ivec3 dim_3d_lod = textureSize(_group_0_binding_5_vs, 1).xyz; - ivec2 dim_2s_ms = textureSize(_group_0_binding_6_vs).xy; + ivec2 dim_2d_array = textureSize(_group_0_binding_4_vs, 0).xy; + ivec2 dim_2d_array_lod = textureSize(_group_0_binding_4_vs, 1).xy; + ivec2 dim_cube = textureSize(_group_0_binding_5_vs, 0).xy; + ivec2 dim_cube_lod = textureSize(_group_0_binding_5_vs, 1).xy; + ivec2 dim_cube_array = textureSize(_group_0_binding_6_vs, 0).xy; + ivec2 dim_cube_array_lod = textureSize(_group_0_binding_6_vs, 1).xy; + ivec3 dim_3d = textureSize(_group_0_binding_7_vs, 0).xyz; + ivec3 dim_3d_lod = textureSize(_group_0_binding_7_vs, 1).xyz; + ivec2 dim_2s_ms = textureSize(_group_0_binding_8_vs).xy; int sum = ((((((((((dim_1d + dim_2d.y) + dim_2d_lod.y) + dim_2d_array.y) + dim_2d_array_lod.y) + dim_cube.y) + dim_cube_lod.y) + dim_cube_array.y) + dim_cube_array_lod.y) + dim_3d.z) + dim_3d_lod.z); gl_Position = vec4(float(sum)); gl_Position.yz = vec2(-gl_Position.y, gl_Position.z * 2.0 - gl_Position.w); diff --git a/tests/out/hlsl/image.hlsl b/tests/out/hlsl/image.hlsl index a21e308dac..10a3c38c03 100644 --- a/tests/out/hlsl/image.hlsl +++ b/tests/out/hlsl/image.hlsl @@ -9,11 +9,13 @@ Texture1D image_1d_src : register(t7); RWTexture1D image_dst : register(u2); Texture1D image_1d : register(t0); Texture2D image_2d : register(t1); -Texture2DArray image_2d_array : register(t2); -TextureCube image_cube : register(t3); -TextureCubeArray image_cube_array : register(t4); -Texture3D image_3d : register(t5); -Texture2DMS image_aa : register(t6); +Texture2D image_2d_u32_ : register(t2); +Texture2D image_2d_i32_ : register(t3); +Texture2DArray image_2d_array : register(t4); +TextureCube image_cube : register(t5); +TextureCubeArray image_cube_array : register(t6); +Texture3D image_3d : register(t7); +Texture2DMS image_aa : register(t8); SamplerState sampler_reg : register(s0, space1); SamplerComparisonState sampler_cmp : register(s1, space1); Texture2D image_2d_depth : register(t2, space1); @@ -258,7 +260,10 @@ float4 gather() : SV_Target0 float4 s2d_offset_1 = image_2d.GatherAlpha(sampler_reg, tc_2, int2(3, 1)); float4 s2d_depth_1 = image_2d_depth.GatherCmp(sampler_cmp, tc_2, 0.5); float4 s2d_depth_offset = image_2d_depth.GatherCmp(sampler_cmp, tc_2, 0.5, int2(3, 1)); - return (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset); + uint4 u = image_2d_u32_.Gather(sampler_reg, tc_2); + int4 i = image_2d_i32_.Gather(sampler_reg, tc_2); + float4 f = (float4(u) + float4(i)); + return ((((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset) + f); } float4 depth_no_comparison() : SV_Target0 diff --git a/tests/out/msl/image.msl b/tests/out/msl/image.msl index df247be4f1..958e147a31 100644 --- a/tests/out/msl/image.msl +++ b/tests/out/msl/image.msl @@ -139,6 +139,8 @@ struct gatherOutput { }; fragment gatherOutput gather( metal::texture2d image_2d [[user(fake0)]] +, metal::texture2d image_2d_u32_ [[user(fake0)]] +, metal::texture2d image_2d_i32_ [[user(fake0)]] , metal::sampler sampler_reg [[user(fake0)]] , metal::sampler sampler_cmp [[user(fake0)]] , metal::depth2d image_2d_depth [[user(fake0)]] @@ -148,7 +150,10 @@ fragment gatherOutput gather( metal::float4 s2d_offset_1 = image_2d.gather(sampler_reg, tc_2, const_type_9_, metal::component::w); metal::float4 s2d_depth_1 = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5); metal::float4 s2d_depth_offset = image_2d_depth.gather_compare(sampler_cmp, tc_2, 0.5, const_type_9_); - return gatherOutput { ((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset }; + metal::uint4 u = image_2d_u32_.gather(sampler_reg, tc_2); + metal::int4 i = image_2d_i32_.gather(sampler_reg, tc_2); + metal::float4 f = static_cast(u) + static_cast(i); + return gatherOutput { (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset) + f }; } diff --git a/tests/out/spv/image.spvasm b/tests/out/spv/image.spvasm index f0556d35d4..23afaea87c 100644 --- a/tests/out/spv/image.spvasm +++ b/tests/out/spv/image.spvasm @@ -1,7 +1,7 @@ ; SPIR-V ; Version: 1.1 ; Generator: rspirv -; Bound: 306 +; Bound: 323 OpCapability SampledCubeArray OpCapability ImageQuery OpCapability Image1D @@ -9,75 +9,77 @@ OpCapability Shader OpCapability Sampled1D %1 = OpExtInstImport "GLSL.std.450" OpMemoryModel Logical GLSL450 -OpEntryPoint GLCompute %76 "main" %73 -OpEntryPoint GLCompute %121 "depth_load" %119 -OpEntryPoint Vertex %142 "queries" %140 -OpEntryPoint Vertex %194 "levels_queries" %193 -OpEntryPoint Fragment %223 "texture_sample" %222 -OpEntryPoint Fragment %252 "texture_sample_comparison" %250 -OpEntryPoint Fragment %272 "gather" %271 -OpEntryPoint Fragment %294 "depth_no_comparison" %293 -OpExecutionMode %76 LocalSize 16 1 1 -OpExecutionMode %121 LocalSize 16 1 1 -OpExecutionMode %223 OriginUpperLeft -OpExecutionMode %252 OriginUpperLeft -OpExecutionMode %272 OriginUpperLeft -OpExecutionMode %294 OriginUpperLeft +OpEntryPoint GLCompute %80 "main" %77 +OpEntryPoint GLCompute %125 "depth_load" %123 +OpEntryPoint Vertex %146 "queries" %144 +OpEntryPoint Vertex %198 "levels_queries" %197 +OpEntryPoint Fragment %227 "texture_sample" %226 +OpEntryPoint Fragment %256 "texture_sample_comparison" %254 +OpEntryPoint Fragment %276 "gather" %275 +OpEntryPoint Fragment %311 "depth_no_comparison" %310 +OpExecutionMode %80 LocalSize 16 1 1 +OpExecutionMode %125 LocalSize 16 1 1 +OpExecutionMode %227 OriginUpperLeft +OpExecutionMode %256 OriginUpperLeft +OpExecutionMode %276 OriginUpperLeft +OpExecutionMode %311 OriginUpperLeft OpSource GLSL 450 -OpName %34 "image_mipmapped_src" -OpName %36 "image_multisampled_src" -OpName %38 "image_depth_multisampled_src" -OpName %40 "image_storage_src" -OpName %42 "image_array_src" -OpName %44 "image_dup_src" -OpName %46 "image_1d_src" -OpName %48 "image_dst" -OpName %50 "image_1d" -OpName %52 "image_2d" -OpName %54 "image_2d_array" -OpName %56 "image_cube" -OpName %58 "image_cube_array" -OpName %60 "image_3d" -OpName %62 "image_aa" -OpName %64 "sampler_reg" -OpName %66 "sampler_cmp" -OpName %68 "image_2d_depth" -OpName %70 "image_cube_depth" -OpName %73 "local_id" -OpName %76 "main" -OpName %119 "local_id" -OpName %121 "depth_load" -OpName %142 "queries" -OpName %194 "levels_queries" -OpName %223 "texture_sample" -OpName %252 "texture_sample_comparison" -OpName %272 "gather" -OpName %294 "depth_no_comparison" -OpDecorate %34 DescriptorSet 0 -OpDecorate %34 Binding 0 -OpDecorate %36 DescriptorSet 0 -OpDecorate %36 Binding 3 -OpDecorate %38 DescriptorSet 0 -OpDecorate %38 Binding 4 -OpDecorate %40 NonWritable -OpDecorate %40 DescriptorSet 0 -OpDecorate %40 Binding 1 -OpDecorate %42 DescriptorSet 0 -OpDecorate %42 Binding 5 -OpDecorate %44 NonWritable -OpDecorate %44 DescriptorSet 0 -OpDecorate %44 Binding 6 -OpDecorate %46 DescriptorSet 0 -OpDecorate %46 Binding 7 -OpDecorate %48 NonReadable -OpDecorate %48 DescriptorSet 0 -OpDecorate %48 Binding 2 -OpDecorate %50 DescriptorSet 0 -OpDecorate %50 Binding 0 -OpDecorate %52 DescriptorSet 0 -OpDecorate %52 Binding 1 -OpDecorate %54 DescriptorSet 0 -OpDecorate %54 Binding 2 +OpName %35 "image_mipmapped_src" +OpName %37 "image_multisampled_src" +OpName %39 "image_depth_multisampled_src" +OpName %41 "image_storage_src" +OpName %43 "image_array_src" +OpName %45 "image_dup_src" +OpName %47 "image_1d_src" +OpName %49 "image_dst" +OpName %51 "image_1d" +OpName %53 "image_2d" +OpName %55 "image_2d_u32" +OpName %56 "image_2d_i32" +OpName %58 "image_2d_array" +OpName %60 "image_cube" +OpName %62 "image_cube_array" +OpName %64 "image_3d" +OpName %66 "image_aa" +OpName %68 "sampler_reg" +OpName %70 "sampler_cmp" +OpName %72 "image_2d_depth" +OpName %74 "image_cube_depth" +OpName %77 "local_id" +OpName %80 "main" +OpName %123 "local_id" +OpName %125 "depth_load" +OpName %146 "queries" +OpName %198 "levels_queries" +OpName %227 "texture_sample" +OpName %256 "texture_sample_comparison" +OpName %276 "gather" +OpName %311 "depth_no_comparison" +OpDecorate %35 DescriptorSet 0 +OpDecorate %35 Binding 0 +OpDecorate %37 DescriptorSet 0 +OpDecorate %37 Binding 3 +OpDecorate %39 DescriptorSet 0 +OpDecorate %39 Binding 4 +OpDecorate %41 NonWritable +OpDecorate %41 DescriptorSet 0 +OpDecorate %41 Binding 1 +OpDecorate %43 DescriptorSet 0 +OpDecorate %43 Binding 5 +OpDecorate %45 NonWritable +OpDecorate %45 DescriptorSet 0 +OpDecorate %45 Binding 6 +OpDecorate %47 DescriptorSet 0 +OpDecorate %47 Binding 7 +OpDecorate %49 NonReadable +OpDecorate %49 DescriptorSet 0 +OpDecorate %49 Binding 2 +OpDecorate %51 DescriptorSet 0 +OpDecorate %51 Binding 0 +OpDecorate %53 DescriptorSet 0 +OpDecorate %53 Binding 1 +OpDecorate %55 DescriptorSet 0 +OpDecorate %55 Binding 2 OpDecorate %56 DescriptorSet 0 OpDecorate %56 Binding 3 OpDecorate %58 DescriptorSet 0 @@ -86,22 +88,26 @@ OpDecorate %60 DescriptorSet 0 OpDecorate %60 Binding 5 OpDecorate %62 DescriptorSet 0 OpDecorate %62 Binding 6 -OpDecorate %64 DescriptorSet 1 -OpDecorate %64 Binding 0 -OpDecorate %66 DescriptorSet 1 -OpDecorate %66 Binding 1 +OpDecorate %64 DescriptorSet 0 +OpDecorate %64 Binding 7 +OpDecorate %66 DescriptorSet 0 +OpDecorate %66 Binding 8 OpDecorate %68 DescriptorSet 1 -OpDecorate %68 Binding 2 +OpDecorate %68 Binding 0 OpDecorate %70 DescriptorSet 1 -OpDecorate %70 Binding 3 -OpDecorate %73 BuiltIn LocalInvocationId -OpDecorate %119 BuiltIn LocalInvocationId -OpDecorate %140 BuiltIn Position -OpDecorate %193 BuiltIn Position -OpDecorate %222 Location 0 -OpDecorate %250 Location 0 -OpDecorate %271 Location 0 -OpDecorate %293 Location 0 +OpDecorate %70 Binding 1 +OpDecorate %72 DescriptorSet 1 +OpDecorate %72 Binding 2 +OpDecorate %74 DescriptorSet 1 +OpDecorate %74 Binding 3 +OpDecorate %77 BuiltIn LocalInvocationId +OpDecorate %123 BuiltIn LocalInvocationId +OpDecorate %144 BuiltIn Position +OpDecorate %197 BuiltIn Position +OpDecorate %226 Location 0 +OpDecorate %254 Location 0 +OpDecorate %275 Location 0 +OpDecorate %310 Location 0 %2 = OpTypeVoid %4 = OpTypeInt 32 1 %3 = OpConstant %4 10 @@ -124,316 +130,333 @@ OpDecorate %293 Location 0 %21 = OpTypeVector %4 2 %22 = OpTypeImage %8 1D 0 0 0 1 Unknown %23 = OpTypeImage %8 2D 0 0 0 1 Unknown -%24 = OpTypeImage %8 2D 0 1 0 1 Unknown -%25 = OpTypeImage %8 Cube 0 0 0 1 Unknown -%26 = OpTypeImage %8 Cube 0 1 0 1 Unknown -%27 = OpTypeImage %8 3D 0 0 0 1 Unknown -%28 = OpTypeImage %8 2D 0 0 1 1 Unknown -%29 = OpTypeVector %8 4 -%30 = OpTypeSampler -%31 = OpTypeImage %8 2D 1 0 0 1 Unknown -%32 = OpTypeImage %8 Cube 1 0 0 1 Unknown -%33 = OpConstantComposite %21 %10 %6 -%35 = OpTypePointer UniformConstant %12 -%34 = OpVariable %35 UniformConstant -%37 = OpTypePointer UniformConstant %14 -%36 = OpVariable %37 UniformConstant -%39 = OpTypePointer UniformConstant %15 -%38 = OpVariable %39 UniformConstant -%41 = OpTypePointer UniformConstant %16 -%40 = OpVariable %41 UniformConstant -%43 = OpTypePointer UniformConstant %17 -%42 = OpVariable %43 UniformConstant -%45 = OpTypePointer UniformConstant %18 -%44 = OpVariable %45 UniformConstant -%47 = OpTypePointer UniformConstant %19 -%46 = OpVariable %47 UniformConstant -%49 = OpTypePointer UniformConstant %18 -%48 = OpVariable %49 UniformConstant -%51 = OpTypePointer UniformConstant %22 -%50 = OpVariable %51 UniformConstant -%53 = OpTypePointer UniformConstant %23 -%52 = OpVariable %53 UniformConstant -%55 = OpTypePointer UniformConstant %24 -%54 = OpVariable %55 UniformConstant -%57 = OpTypePointer UniformConstant %25 +%24 = OpTypeImage %4 2D 0 0 0 1 Unknown +%25 = OpTypeImage %8 2D 0 1 0 1 Unknown +%26 = OpTypeImage %8 Cube 0 0 0 1 Unknown +%27 = OpTypeImage %8 Cube 0 1 0 1 Unknown +%28 = OpTypeImage %8 3D 0 0 0 1 Unknown +%29 = OpTypeImage %8 2D 0 0 1 1 Unknown +%30 = OpTypeVector %8 4 +%31 = OpTypeSampler +%32 = OpTypeImage %8 2D 1 0 0 1 Unknown +%33 = OpTypeImage %8 Cube 1 0 0 1 Unknown +%34 = OpConstantComposite %21 %10 %6 +%36 = OpTypePointer UniformConstant %12 +%35 = OpVariable %36 UniformConstant +%38 = OpTypePointer UniformConstant %14 +%37 = OpVariable %38 UniformConstant +%40 = OpTypePointer UniformConstant %15 +%39 = OpVariable %40 UniformConstant +%42 = OpTypePointer UniformConstant %16 +%41 = OpVariable %42 UniformConstant +%44 = OpTypePointer UniformConstant %17 +%43 = OpVariable %44 UniformConstant +%46 = OpTypePointer UniformConstant %18 +%45 = OpVariable %46 UniformConstant +%48 = OpTypePointer UniformConstant %19 +%47 = OpVariable %48 UniformConstant +%50 = OpTypePointer UniformConstant %18 +%49 = OpVariable %50 UniformConstant +%52 = OpTypePointer UniformConstant %22 +%51 = OpVariable %52 UniformConstant +%54 = OpTypePointer UniformConstant %23 +%53 = OpVariable %54 UniformConstant +%55 = OpVariable %36 UniformConstant +%57 = OpTypePointer UniformConstant %24 %56 = OpVariable %57 UniformConstant -%59 = OpTypePointer UniformConstant %26 +%59 = OpTypePointer UniformConstant %25 %58 = OpVariable %59 UniformConstant -%61 = OpTypePointer UniformConstant %27 +%61 = OpTypePointer UniformConstant %26 %60 = OpVariable %61 UniformConstant -%63 = OpTypePointer UniformConstant %28 +%63 = OpTypePointer UniformConstant %27 %62 = OpVariable %63 UniformConstant -%65 = OpTypePointer UniformConstant %30 +%65 = OpTypePointer UniformConstant %28 %64 = OpVariable %65 UniformConstant -%67 = OpTypePointer UniformConstant %30 +%67 = OpTypePointer UniformConstant %29 %66 = OpVariable %67 UniformConstant %69 = OpTypePointer UniformConstant %31 %68 = OpVariable %69 UniformConstant -%71 = OpTypePointer UniformConstant %32 +%71 = OpTypePointer UniformConstant %31 %70 = OpVariable %71 UniformConstant -%74 = OpTypePointer Input %20 -%73 = OpVariable %74 Input -%77 = OpTypeFunction %2 -%86 = OpTypeVector %13 2 -%94 = OpTypeVector %13 4 -%105 = OpTypeVector %4 3 -%119 = OpVariable %74 Input -%141 = OpTypePointer Output %29 -%140 = OpVariable %141 Output -%151 = OpConstant %13 0 -%193 = OpVariable %141 Output -%222 = OpVariable %141 Output -%228 = OpTypeVector %8 2 -%231 = OpTypeSampledImage %22 -%234 = OpTypeSampledImage %23 -%251 = OpTypePointer Output %8 -%250 = OpVariable %251 Output -%258 = OpTypeSampledImage %31 -%263 = OpConstant %8 0.0 -%264 = OpTypeVector %8 3 -%266 = OpTypeSampledImage %32 -%271 = OpVariable %141 Output -%281 = OpConstant %13 1 -%284 = OpConstant %13 3 -%293 = OpVariable %141 Output -%76 = OpFunction %2 None %77 -%72 = OpLabel -%75 = OpLoad %20 %73 -%78 = OpLoad %12 %34 -%79 = OpLoad %14 %36 -%80 = OpLoad %16 %40 -%81 = OpLoad %17 %42 -%82 = OpLoad %19 %46 -%83 = OpLoad %18 %48 -OpBranch %84 -%84 = OpLabel -%85 = OpImageQuerySize %21 %80 -%87 = OpVectorShuffle %86 %75 %75 0 1 -%88 = OpBitcast %21 %87 -%89 = OpIMul %21 %85 %88 -%90 = OpCompositeConstruct %21 %3 %5 -%91 = OpSRem %21 %89 %90 -%92 = OpCompositeExtract %13 %75 2 -%93 = OpBitcast %4 %92 -%95 = OpImageFetch %94 %78 %91 Lod %93 -%96 = OpCompositeExtract %13 %75 2 +%73 = OpTypePointer UniformConstant %32 +%72 = OpVariable %73 UniformConstant +%75 = OpTypePointer UniformConstant %33 +%74 = OpVariable %75 UniformConstant +%78 = OpTypePointer Input %20 +%77 = OpVariable %78 Input +%81 = OpTypeFunction %2 +%90 = OpTypeVector %13 2 +%98 = OpTypeVector %13 4 +%109 = OpTypeVector %4 3 +%123 = OpVariable %78 Input +%145 = OpTypePointer Output %30 +%144 = OpVariable %145 Output +%155 = OpConstant %13 0 +%197 = OpVariable %145 Output +%226 = OpVariable %145 Output +%232 = OpTypeVector %8 2 +%235 = OpTypeSampledImage %22 +%238 = OpTypeSampledImage %23 +%255 = OpTypePointer Output %8 +%254 = OpVariable %255 Output +%262 = OpTypeSampledImage %32 +%267 = OpConstant %8 0.0 +%268 = OpTypeVector %8 3 +%270 = OpTypeSampledImage %33 +%275 = OpVariable %145 Output +%287 = OpConstant %13 1 +%290 = OpConstant %13 3 +%295 = OpTypeSampledImage %12 +%298 = OpTypeVector %4 4 +%299 = OpTypeSampledImage %24 +%310 = OpVariable %145 Output +%80 = OpFunction %2 None %81 +%76 = OpLabel +%79 = OpLoad %20 %77 +%82 = OpLoad %12 %35 +%83 = OpLoad %14 %37 +%84 = OpLoad %16 %41 +%85 = OpLoad %17 %43 +%86 = OpLoad %19 %47 +%87 = OpLoad %18 %49 +OpBranch %88 +%88 = OpLabel +%89 = OpImageQuerySize %21 %84 +%91 = OpVectorShuffle %90 %79 %79 0 1 +%92 = OpBitcast %21 %91 +%93 = OpIMul %21 %89 %92 +%94 = OpCompositeConstruct %21 %3 %5 +%95 = OpSRem %21 %93 %94 +%96 = OpCompositeExtract %13 %79 2 %97 = OpBitcast %4 %96 -%98 = OpImageFetch %94 %79 %91 Sample %97 -%99 = OpImageRead %94 %80 %91 -%100 = OpCompositeExtract %13 %75 2 +%99 = OpImageFetch %98 %82 %95 Lod %97 +%100 = OpCompositeExtract %13 %79 2 %101 = OpBitcast %4 %100 -%102 = OpCompositeExtract %13 %75 2 -%103 = OpBitcast %4 %102 -%104 = OpIAdd %4 %103 %6 -%106 = OpCompositeConstruct %105 %91 %101 -%107 = OpImageFetch %94 %81 %106 Lod %104 -%108 = OpCompositeExtract %13 %75 0 -%109 = OpBitcast %4 %108 -%110 = OpCompositeExtract %13 %75 2 -%111 = OpBitcast %4 %110 -%112 = OpImageFetch %94 %82 %109 Lod %111 -%113 = OpCompositeExtract %4 %91 0 -%114 = OpIAdd %94 %95 %98 -%115 = OpIAdd %94 %114 %99 -%116 = OpIAdd %94 %115 %107 -%117 = OpIAdd %94 %116 %112 -OpImageWrite %83 %113 %117 +%102 = OpImageFetch %98 %83 %95 Sample %101 +%103 = OpImageRead %98 %84 %95 +%104 = OpCompositeExtract %13 %79 2 +%105 = OpBitcast %4 %104 +%106 = OpCompositeExtract %13 %79 2 +%107 = OpBitcast %4 %106 +%108 = OpIAdd %4 %107 %6 +%110 = OpCompositeConstruct %109 %95 %105 +%111 = OpImageFetch %98 %85 %110 Lod %108 +%112 = OpCompositeExtract %13 %79 0 +%113 = OpBitcast %4 %112 +%114 = OpCompositeExtract %13 %79 2 +%115 = OpBitcast %4 %114 +%116 = OpImageFetch %98 %86 %113 Lod %115 +%117 = OpCompositeExtract %4 %95 0 +%118 = OpIAdd %98 %99 %102 +%119 = OpIAdd %98 %118 %103 +%120 = OpIAdd %98 %119 %111 +%121 = OpIAdd %98 %120 %116 +OpImageWrite %87 %117 %121 OpReturn OpFunctionEnd -%121 = OpFunction %2 None %77 -%118 = OpLabel -%120 = OpLoad %20 %119 -%122 = OpLoad %15 %38 -%123 = OpLoad %16 %40 -%124 = OpLoad %18 %48 -OpBranch %125 -%125 = OpLabel -%126 = OpImageQuerySize %21 %123 -%127 = OpVectorShuffle %86 %120 %120 0 1 -%128 = OpBitcast %21 %127 -%129 = OpIMul %21 %126 %128 -%130 = OpCompositeConstruct %21 %3 %5 -%131 = OpSRem %21 %129 %130 -%132 = OpCompositeExtract %13 %120 2 -%133 = OpBitcast %4 %132 -%134 = OpImageFetch %29 %122 %131 Sample %133 -%135 = OpCompositeExtract %8 %134 0 -%136 = OpCompositeExtract %4 %131 0 -%137 = OpConvertFToU %13 %135 -%138 = OpCompositeConstruct %94 %137 %137 %137 %137 -OpImageWrite %124 %136 %138 +%125 = OpFunction %2 None %81 +%122 = OpLabel +%124 = OpLoad %20 %123 +%126 = OpLoad %15 %39 +%127 = OpLoad %16 %41 +%128 = OpLoad %18 %49 +OpBranch %129 +%129 = OpLabel +%130 = OpImageQuerySize %21 %127 +%131 = OpVectorShuffle %90 %124 %124 0 1 +%132 = OpBitcast %21 %131 +%133 = OpIMul %21 %130 %132 +%134 = OpCompositeConstruct %21 %3 %5 +%135 = OpSRem %21 %133 %134 +%136 = OpCompositeExtract %13 %124 2 +%137 = OpBitcast %4 %136 +%138 = OpImageFetch %30 %126 %135 Sample %137 +%139 = OpCompositeExtract %8 %138 0 +%140 = OpCompositeExtract %4 %135 0 +%141 = OpConvertFToU %13 %139 +%142 = OpCompositeConstruct %98 %141 %141 %141 %141 +OpImageWrite %128 %140 %142 OpReturn OpFunctionEnd -%142 = OpFunction %2 None %77 -%139 = OpLabel -%143 = OpLoad %22 %50 -%144 = OpLoad %23 %52 -%145 = OpLoad %24 %54 -%146 = OpLoad %25 %56 -%147 = OpLoad %26 %58 -%148 = OpLoad %27 %60 -%149 = OpLoad %28 %62 -OpBranch %150 -%150 = OpLabel -%152 = OpImageQuerySizeLod %4 %143 %151 -%154 = OpImageQuerySizeLod %4 %143 %152 -%155 = OpImageQuerySizeLod %21 %144 %151 -%156 = OpImageQuerySizeLod %21 %144 %6 -%157 = OpImageQuerySizeLod %105 %145 %151 -%158 = OpVectorShuffle %21 %157 %157 0 1 -%159 = OpImageQuerySizeLod %105 %145 %6 -%160 = OpVectorShuffle %21 %159 %159 0 1 -%161 = OpImageQuerySizeLod %21 %146 %151 -%162 = OpImageQuerySizeLod %21 %146 %6 -%163 = OpImageQuerySizeLod %105 %147 %151 -%164 = OpVectorShuffle %21 %163 %163 0 0 -%165 = OpImageQuerySizeLod %105 %147 %6 -%166 = OpVectorShuffle %21 %165 %165 0 0 -%167 = OpImageQuerySizeLod %105 %148 %151 -%168 = OpImageQuerySizeLod %105 %148 %6 -%169 = OpImageQuerySize %21 %149 -%170 = OpCompositeExtract %4 %155 1 -%171 = OpIAdd %4 %152 %170 -%172 = OpCompositeExtract %4 %156 1 -%173 = OpIAdd %4 %171 %172 -%174 = OpCompositeExtract %4 %158 1 -%175 = OpIAdd %4 %173 %174 +%146 = OpFunction %2 None %81 +%143 = OpLabel +%147 = OpLoad %22 %51 +%148 = OpLoad %23 %53 +%149 = OpLoad %25 %58 +%150 = OpLoad %26 %60 +%151 = OpLoad %27 %62 +%152 = OpLoad %28 %64 +%153 = OpLoad %29 %66 +OpBranch %154 +%154 = OpLabel +%156 = OpImageQuerySizeLod %4 %147 %155 +%158 = OpImageQuerySizeLod %4 %147 %156 +%159 = OpImageQuerySizeLod %21 %148 %155 +%160 = OpImageQuerySizeLod %21 %148 %6 +%161 = OpImageQuerySizeLod %109 %149 %155 +%162 = OpVectorShuffle %21 %161 %161 0 1 +%163 = OpImageQuerySizeLod %109 %149 %6 +%164 = OpVectorShuffle %21 %163 %163 0 1 +%165 = OpImageQuerySizeLod %21 %150 %155 +%166 = OpImageQuerySizeLod %21 %150 %6 +%167 = OpImageQuerySizeLod %109 %151 %155 +%168 = OpVectorShuffle %21 %167 %167 0 0 +%169 = OpImageQuerySizeLod %109 %151 %6 +%170 = OpVectorShuffle %21 %169 %169 0 0 +%171 = OpImageQuerySizeLod %109 %152 %155 +%172 = OpImageQuerySizeLod %109 %152 %6 +%173 = OpImageQuerySize %21 %153 +%174 = OpCompositeExtract %4 %159 1 +%175 = OpIAdd %4 %156 %174 %176 = OpCompositeExtract %4 %160 1 %177 = OpIAdd %4 %175 %176 -%178 = OpCompositeExtract %4 %161 1 +%178 = OpCompositeExtract %4 %162 1 %179 = OpIAdd %4 %177 %178 -%180 = OpCompositeExtract %4 %162 1 +%180 = OpCompositeExtract %4 %164 1 %181 = OpIAdd %4 %179 %180 -%182 = OpCompositeExtract %4 %164 1 +%182 = OpCompositeExtract %4 %165 1 %183 = OpIAdd %4 %181 %182 %184 = OpCompositeExtract %4 %166 1 %185 = OpIAdd %4 %183 %184 -%186 = OpCompositeExtract %4 %167 2 +%186 = OpCompositeExtract %4 %168 1 %187 = OpIAdd %4 %185 %186 -%188 = OpCompositeExtract %4 %168 2 +%188 = OpCompositeExtract %4 %170 1 %189 = OpIAdd %4 %187 %188 -%190 = OpConvertSToF %8 %189 -%191 = OpCompositeConstruct %29 %190 %190 %190 %190 -OpStore %140 %191 +%190 = OpCompositeExtract %4 %171 2 +%191 = OpIAdd %4 %189 %190 +%192 = OpCompositeExtract %4 %172 2 +%193 = OpIAdd %4 %191 %192 +%194 = OpConvertSToF %8 %193 +%195 = OpCompositeConstruct %30 %194 %194 %194 %194 +OpStore %144 %195 OpReturn OpFunctionEnd -%194 = OpFunction %2 None %77 -%192 = OpLabel -%195 = OpLoad %23 %52 -%196 = OpLoad %24 %54 -%197 = OpLoad %25 %56 -%198 = OpLoad %26 %58 -%199 = OpLoad %27 %60 -%200 = OpLoad %28 %62 -OpBranch %201 -%201 = OpLabel -%202 = OpImageQueryLevels %4 %195 -%203 = OpImageQueryLevels %4 %196 -%204 = OpImageQuerySizeLod %105 %196 %151 -%205 = OpCompositeExtract %4 %204 2 -%206 = OpImageQueryLevels %4 %197 -%207 = OpImageQueryLevels %4 %198 -%208 = OpImageQuerySizeLod %105 %198 %151 +%198 = OpFunction %2 None %81 +%196 = OpLabel +%199 = OpLoad %23 %53 +%200 = OpLoad %25 %58 +%201 = OpLoad %26 %60 +%202 = OpLoad %27 %62 +%203 = OpLoad %28 %64 +%204 = OpLoad %29 %66 +OpBranch %205 +%205 = OpLabel +%206 = OpImageQueryLevels %4 %199 +%207 = OpImageQueryLevels %4 %200 +%208 = OpImageQuerySizeLod %109 %200 %155 %209 = OpCompositeExtract %4 %208 2 -%210 = OpImageQueryLevels %4 %199 -%211 = OpImageQuerySamples %4 %200 -%212 = OpIAdd %4 %205 %209 -%213 = OpIAdd %4 %212 %211 -%214 = OpIAdd %4 %213 %202 -%215 = OpIAdd %4 %214 %203 -%216 = OpIAdd %4 %215 %210 -%217 = OpIAdd %4 %216 %206 -%218 = OpIAdd %4 %217 %207 -%219 = OpConvertSToF %8 %218 -%220 = OpCompositeConstruct %29 %219 %219 %219 %219 -OpStore %193 %220 +%210 = OpImageQueryLevels %4 %201 +%211 = OpImageQueryLevels %4 %202 +%212 = OpImageQuerySizeLod %109 %202 %155 +%213 = OpCompositeExtract %4 %212 2 +%214 = OpImageQueryLevels %4 %203 +%215 = OpImageQuerySamples %4 %204 +%216 = OpIAdd %4 %209 %213 +%217 = OpIAdd %4 %216 %215 +%218 = OpIAdd %4 %217 %206 +%219 = OpIAdd %4 %218 %207 +%220 = OpIAdd %4 %219 %214 +%221 = OpIAdd %4 %220 %210 +%222 = OpIAdd %4 %221 %211 +%223 = OpConvertSToF %8 %222 +%224 = OpCompositeConstruct %30 %223 %223 %223 %223 +OpStore %197 %224 OpReturn OpFunctionEnd -%223 = OpFunction %2 None %77 -%221 = OpLabel -%224 = OpLoad %22 %50 -%225 = OpLoad %23 %52 -%226 = OpLoad %30 %64 -OpBranch %227 -%227 = OpLabel -%229 = OpCompositeConstruct %228 %7 %7 -%230 = OpCompositeExtract %8 %229 0 -%232 = OpSampledImage %231 %224 %226 -%233 = OpImageSampleImplicitLod %29 %232 %230 -%235 = OpSampledImage %234 %225 %226 -%236 = OpImageSampleImplicitLod %29 %235 %229 -%237 = OpSampledImage %234 %225 %226 -%238 = OpImageSampleImplicitLod %29 %237 %229 ConstOffset %33 -%239 = OpSampledImage %234 %225 %226 -%240 = OpImageSampleExplicitLod %29 %239 %229 Lod %9 -%241 = OpSampledImage %234 %225 %226 -%242 = OpImageSampleExplicitLod %29 %241 %229 Lod|ConstOffset %9 %33 -%243 = OpSampledImage %234 %225 %226 -%244 = OpImageSampleImplicitLod %29 %243 %229 Bias|ConstOffset %11 %33 -%245 = OpFAdd %29 %233 %236 -%246 = OpFAdd %29 %245 %238 -%247 = OpFAdd %29 %246 %240 -%248 = OpFAdd %29 %247 %242 -OpStore %222 %248 +%227 = OpFunction %2 None %81 +%225 = OpLabel +%228 = OpLoad %22 %51 +%229 = OpLoad %23 %53 +%230 = OpLoad %31 %68 +OpBranch %231 +%231 = OpLabel +%233 = OpCompositeConstruct %232 %7 %7 +%234 = OpCompositeExtract %8 %233 0 +%236 = OpSampledImage %235 %228 %230 +%237 = OpImageSampleImplicitLod %30 %236 %234 +%239 = OpSampledImage %238 %229 %230 +%240 = OpImageSampleImplicitLod %30 %239 %233 +%241 = OpSampledImage %238 %229 %230 +%242 = OpImageSampleImplicitLod %30 %241 %233 ConstOffset %34 +%243 = OpSampledImage %238 %229 %230 +%244 = OpImageSampleExplicitLod %30 %243 %233 Lod %9 +%245 = OpSampledImage %238 %229 %230 +%246 = OpImageSampleExplicitLod %30 %245 %233 Lod|ConstOffset %9 %34 +%247 = OpSampledImage %238 %229 %230 +%248 = OpImageSampleImplicitLod %30 %247 %233 Bias|ConstOffset %11 %34 +%249 = OpFAdd %30 %237 %240 +%250 = OpFAdd %30 %249 %242 +%251 = OpFAdd %30 %250 %244 +%252 = OpFAdd %30 %251 %246 +OpStore %226 %252 OpReturn OpFunctionEnd -%252 = OpFunction %2 None %77 -%249 = OpLabel -%253 = OpLoad %30 %66 -%254 = OpLoad %31 %68 -%255 = OpLoad %32 %70 -OpBranch %256 -%256 = OpLabel -%257 = OpCompositeConstruct %228 %7 %7 -%259 = OpSampledImage %258 %254 %253 -%260 = OpImageSampleDrefImplicitLod %8 %259 %257 %7 -%261 = OpSampledImage %258 %254 %253 -%262 = OpImageSampleDrefExplicitLod %8 %261 %257 %7 Lod %263 -%265 = OpCompositeConstruct %264 %7 %7 %7 -%267 = OpSampledImage %266 %255 %253 -%268 = OpImageSampleDrefExplicitLod %8 %267 %265 %7 Lod %263 -%269 = OpFAdd %8 %260 %262 -OpStore %250 %269 +%256 = OpFunction %2 None %81 +%253 = OpLabel +%257 = OpLoad %31 %70 +%258 = OpLoad %32 %72 +%259 = OpLoad %33 %74 +OpBranch %260 +%260 = OpLabel +%261 = OpCompositeConstruct %232 %7 %7 +%263 = OpSampledImage %262 %258 %257 +%264 = OpImageSampleDrefImplicitLod %8 %263 %261 %7 +%265 = OpSampledImage %262 %258 %257 +%266 = OpImageSampleDrefExplicitLod %8 %265 %261 %7 Lod %267 +%269 = OpCompositeConstruct %268 %7 %7 %7 +%271 = OpSampledImage %270 %259 %257 +%272 = OpImageSampleDrefExplicitLod %8 %271 %269 %7 Lod %267 +%273 = OpFAdd %8 %264 %266 +OpStore %254 %273 OpReturn OpFunctionEnd -%272 = OpFunction %2 None %77 -%270 = OpLabel -%273 = OpLoad %23 %52 -%274 = OpLoad %30 %64 -%275 = OpLoad %30 %66 -%276 = OpLoad %31 %68 -OpBranch %277 -%277 = OpLabel -%278 = OpCompositeConstruct %228 %7 %7 -%279 = OpSampledImage %234 %273 %274 -%280 = OpImageGather %29 %279 %278 %281 -%282 = OpSampledImage %234 %273 %274 -%283 = OpImageGather %29 %282 %278 %284 ConstOffset %33 -%285 = OpSampledImage %258 %276 %275 -%286 = OpImageDrefGather %29 %285 %278 %7 -%287 = OpSampledImage %258 %276 %275 -%288 = OpImageDrefGather %29 %287 %278 %7 ConstOffset %33 -%289 = OpFAdd %29 %280 %283 -%290 = OpFAdd %29 %289 %286 -%291 = OpFAdd %29 %290 %288 -OpStore %271 %291 +%276 = OpFunction %2 None %81 +%274 = OpLabel +%277 = OpLoad %23 %53 +%278 = OpLoad %12 %55 +%279 = OpLoad %24 %56 +%280 = OpLoad %31 %68 +%281 = OpLoad %31 %70 +%282 = OpLoad %32 %72 +OpBranch %283 +%283 = OpLabel +%284 = OpCompositeConstruct %232 %7 %7 +%285 = OpSampledImage %238 %277 %280 +%286 = OpImageGather %30 %285 %284 %287 +%288 = OpSampledImage %238 %277 %280 +%289 = OpImageGather %30 %288 %284 %290 ConstOffset %34 +%291 = OpSampledImage %262 %282 %281 +%292 = OpImageDrefGather %30 %291 %284 %7 +%293 = OpSampledImage %262 %282 %281 +%294 = OpImageDrefGather %30 %293 %284 %7 ConstOffset %34 +%296 = OpSampledImage %295 %278 %280 +%297 = OpImageGather %98 %296 %284 %155 +%300 = OpSampledImage %299 %279 %280 +%301 = OpImageGather %298 %300 %284 %155 +%302 = OpConvertUToF %30 %297 +%303 = OpConvertSToF %30 %301 +%304 = OpFAdd %30 %302 %303 +%305 = OpFAdd %30 %286 %289 +%306 = OpFAdd %30 %305 %292 +%307 = OpFAdd %30 %306 %294 +%308 = OpFAdd %30 %307 %304 +OpStore %275 %308 OpReturn OpFunctionEnd -%294 = OpFunction %2 None %77 -%292 = OpLabel -%295 = OpLoad %30 %64 -%296 = OpLoad %31 %68 -OpBranch %297 -%297 = OpLabel -%298 = OpCompositeConstruct %228 %7 %7 -%299 = OpSampledImage %258 %296 %295 -%300 = OpImageSampleImplicitLod %29 %299 %298 -%301 = OpCompositeExtract %8 %300 0 -%302 = OpSampledImage %258 %296 %295 -%303 = OpImageGather %29 %302 %298 %151 -%304 = OpCompositeConstruct %29 %301 %301 %301 %301 -%305 = OpFAdd %29 %304 %303 -OpStore %293 %305 +%311 = OpFunction %2 None %81 +%309 = OpLabel +%312 = OpLoad %31 %68 +%313 = OpLoad %32 %72 +OpBranch %314 +%314 = OpLabel +%315 = OpCompositeConstruct %232 %7 %7 +%316 = OpSampledImage %262 %313 %312 +%317 = OpImageSampleImplicitLod %30 %316 %315 +%318 = OpCompositeExtract %8 %317 0 +%319 = OpSampledImage %262 %313 %312 +%320 = OpImageGather %30 %319 %315 %155 +%321 = OpCompositeConstruct %30 %318 %318 %318 %318 +%322 = OpFAdd %30 %321 %320 +OpStore %310 %322 OpReturn OpFunctionEnd \ No newline at end of file diff --git a/tests/out/wgsl/image.wgsl b/tests/out/wgsl/image.wgsl index 062876f224..0c1dc48f88 100644 --- a/tests/out/wgsl/image.wgsl +++ b/tests/out/wgsl/image.wgsl @@ -19,14 +19,18 @@ var image_1d: texture_1d; @group(0) @binding(1) var image_2d: texture_2d; @group(0) @binding(2) -var image_2d_array: texture_2d_array; +var image_2d_u32_: texture_2d; @group(0) @binding(3) -var image_cube: texture_cube; +var image_2d_i32_: texture_2d; @group(0) @binding(4) -var image_cube_array: texture_cube_array; +var image_2d_array: texture_2d_array; @group(0) @binding(5) -var image_3d: texture_3d; +var image_cube: texture_cube; @group(0) @binding(6) +var image_cube_array: texture_cube_array; +@group(0) @binding(7) +var image_3d: texture_3d; +@group(0) @binding(8) var image_aa: texture_multisampled_2d; @group(1) @binding(0) var sampler_reg: sampler; @@ -120,7 +124,10 @@ fn gather() -> @location(0) vec4 { let s2d_offset_1 = textureGather(3, image_2d, sampler_reg, tc_2, vec2(3, 1)); let s2d_depth_1 = textureGatherCompare(image_2d_depth, sampler_cmp, tc_2, 0.5); let s2d_depth_offset = textureGatherCompare(image_2d_depth, sampler_cmp, tc_2, 0.5, vec2(3, 1)); - return (((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset); + let u = textureGather(0, image_2d_u32_, sampler_reg, tc_2); + let i = textureGather(0, image_2d_i32_, sampler_reg, tc_2); + let f = (vec4(u) + vec4(i)); + return ((((s2d_1 + s2d_offset_1) + s2d_depth_1) + s2d_depth_offset) + f); } @fragment