[spv-out] Add test for duplicate OpTypeImage instructions.

This commit is contained in:
Jim Blandy 2021-08-28 23:54:33 -07:00 committed by Dzmitry Malyshau
parent 81fbad182c
commit b226c5108a
5 changed files with 210 additions and 199 deletions

View File

@ -8,6 +8,8 @@ var image_depth_multisampled_src: texture_depth_multisampled_2d;
var image_storage_src: texture_storage_2d<rgba8uint>; var image_storage_src: texture_storage_2d<rgba8uint>;
[[group(0), binding(5)]] [[group(0), binding(5)]]
var image_array_src: texture_2d_array<u32>; var image_array_src: texture_2d_array<u32>;
[[group(0), binding(6)]]
var image_dup_src: texture_storage_1d<r32uint,read>; // for #1307
[[group(0), binding(2)]] [[group(0), binding(2)]]
var image_dst: texture_storage_1d<r32uint,write>; var image_dst: texture_storage_1d<r32uint,write>;

View File

@ -4,6 +4,7 @@ Texture2DMS<uint4> image_multisampled_src : register(t3);
Texture2DMS<float> image_depth_multisampled_src : register(t4); Texture2DMS<float> image_depth_multisampled_src : register(t4);
RWTexture2D<uint4> image_storage_src : register(u1); RWTexture2D<uint4> image_storage_src : register(u1);
Texture2DArray<uint4> image_array_src : register(t5); Texture2DArray<uint4> image_array_src : register(t5);
RWTexture1D<uint4> image_dup_src : register(u6);
RWTexture1D<uint4> image_dst : register(u2); RWTexture1D<uint4> image_dst : register(u2);
Texture1D<float4> image_1d : register(t0); Texture1D<float4> image_1d : register(t0);
Texture2D<float4> image_2d : register(t1); Texture2D<float4> image_2d : register(t1);

View File

@ -2,7 +2,7 @@
#include <metal_stdlib> #include <metal_stdlib>
#include <simd/simd.h> #include <simd/simd.h>
constant metal::int2 const_type7_ = {3, 1}; constant metal::int2 const_type8_ = {3, 1};
struct main1Input { struct main1Input {
}; };
@ -72,9 +72,9 @@ fragment sampleOutput sample(
) { ) {
metal::float2 tc = metal::float2(0.5); metal::float2 tc = metal::float2(0.5);
metal::float4 s2d = image_2d.sample(sampler_reg, tc); metal::float4 s2d = image_2d.sample(sampler_reg, tc);
metal::float4 s2d_offset = image_2d.sample(sampler_reg, tc, const_type7_); metal::float4 s2d_offset = image_2d.sample(sampler_reg, tc, const_type8_);
metal::float4 s2d_level = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284)); metal::float4 s2d_level = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284));
metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284), const_type7_); metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.299999952316284), const_type8_);
return sampleOutput { ((s2d + s2d_offset) + s2d_level) + s2d_level_offset }; return sampleOutput { ((s2d + s2d_offset) + s2d_level) + s2d_level_offset };
} }

View File

@ -1,7 +1,7 @@
; SPIR-V ; SPIR-V
; Version: 1.1 ; Version: 1.1
; Generator: rspirv ; Generator: rspirv
; Bound: 213 ; Bound: 215
OpCapability SampledCubeArray OpCapability SampledCubeArray
OpCapability ImageQuery OpCapability ImageQuery
OpCapability Image1D OpCapability Image1D
@ -9,35 +9,36 @@ OpCapability Shader
OpCapability Sampled1D OpCapability Sampled1D
%1 = OpExtInstImport "GLSL.std.450" %1 = OpExtInstImport "GLSL.std.450"
OpMemoryModel Logical GLSL450 OpMemoryModel Logical GLSL450
OpEntryPoint GLCompute %67 "main" %64 OpEntryPoint GLCompute %69 "main" %66
OpEntryPoint Vertex %113 "queries" %111 OpEntryPoint Vertex %115 "queries" %113
OpEntryPoint Fragment %180 "sample" %179 OpEntryPoint Fragment %182 "sample" %181
OpEntryPoint Fragment %201 "sample_comparison" %199 OpEntryPoint Fragment %203 "sample_comparison" %201
OpExecutionMode %67 LocalSize 16 1 1 OpExecutionMode %69 LocalSize 16 1 1
OpExecutionMode %180 OriginUpperLeft OpExecutionMode %182 OriginUpperLeft
OpExecutionMode %201 OriginUpperLeft OpExecutionMode %203 OriginUpperLeft
OpSource GLSL 450 OpSource GLSL 450
OpName %31 "image_mipmapped_src" OpName %31 "image_mipmapped_src"
OpName %33 "image_multisampled_src" OpName %33 "image_multisampled_src"
OpName %35 "image_depth_multisampled_src" OpName %35 "image_depth_multisampled_src"
OpName %37 "image_storage_src" OpName %37 "image_storage_src"
OpName %39 "image_array_src" OpName %39 "image_array_src"
OpName %41 "image_dst" OpName %41 "image_dup_src"
OpName %43 "image_1d" OpName %43 "image_dst"
OpName %45 "image_2d" OpName %45 "image_1d"
OpName %47 "image_2d_array" OpName %47 "image_2d"
OpName %49 "image_cube" OpName %49 "image_2d_array"
OpName %51 "image_cube_array" OpName %51 "image_cube"
OpName %53 "image_3d" OpName %53 "image_cube_array"
OpName %55 "image_aa" OpName %55 "image_3d"
OpName %57 "sampler_reg" OpName %57 "image_aa"
OpName %59 "sampler_cmp" OpName %59 "sampler_reg"
OpName %61 "image_2d_depth" OpName %61 "sampler_cmp"
OpName %64 "local_id" OpName %63 "image_2d_depth"
OpName %67 "main" OpName %66 "local_id"
OpName %113 "queries" OpName %69 "main"
OpName %180 "sample" OpName %115 "queries"
OpName %201 "sample_comparison" OpName %182 "sample"
OpName %203 "sample_comparison"
OpDecorate %31 DescriptorSet 0 OpDecorate %31 DescriptorSet 0
OpDecorate %31 Binding 0 OpDecorate %31 Binding 0
OpDecorate %33 DescriptorSet 0 OpDecorate %33 DescriptorSet 0
@ -49,33 +50,36 @@ OpDecorate %37 DescriptorSet 0
OpDecorate %37 Binding 1 OpDecorate %37 Binding 1
OpDecorate %39 DescriptorSet 0 OpDecorate %39 DescriptorSet 0
OpDecorate %39 Binding 5 OpDecorate %39 Binding 5
OpDecorate %41 NonReadable OpDecorate %41 NonWritable
OpDecorate %41 DescriptorSet 0 OpDecorate %41 DescriptorSet 0
OpDecorate %41 Binding 2 OpDecorate %41 Binding 6
OpDecorate %43 NonReadable
OpDecorate %43 DescriptorSet 0 OpDecorate %43 DescriptorSet 0
OpDecorate %43 Binding 0 OpDecorate %43 Binding 2
OpDecorate %45 DescriptorSet 0 OpDecorate %45 DescriptorSet 0
OpDecorate %45 Binding 1 OpDecorate %45 Binding 0
OpDecorate %47 DescriptorSet 0 OpDecorate %47 DescriptorSet 0
OpDecorate %47 Binding 2 OpDecorate %47 Binding 1
OpDecorate %49 DescriptorSet 0 OpDecorate %49 DescriptorSet 0
OpDecorate %49 Binding 3 OpDecorate %49 Binding 2
OpDecorate %51 DescriptorSet 0 OpDecorate %51 DescriptorSet 0
OpDecorate %51 Binding 4 OpDecorate %51 Binding 3
OpDecorate %53 DescriptorSet 0 OpDecorate %53 DescriptorSet 0
OpDecorate %53 Binding 5 OpDecorate %53 Binding 4
OpDecorate %55 DescriptorSet 0 OpDecorate %55 DescriptorSet 0
OpDecorate %55 Binding 6 OpDecorate %55 Binding 5
OpDecorate %57 DescriptorSet 1 OpDecorate %57 DescriptorSet 0
OpDecorate %57 Binding 0 OpDecorate %57 Binding 6
OpDecorate %59 DescriptorSet 1 OpDecorate %59 DescriptorSet 1
OpDecorate %59 Binding 1 OpDecorate %59 Binding 0
OpDecorate %61 DescriptorSet 1 OpDecorate %61 DescriptorSet 1
OpDecorate %61 Binding 2 OpDecorate %61 Binding 1
OpDecorate %64 BuiltIn LocalInvocationId OpDecorate %63 DescriptorSet 1
OpDecorate %111 BuiltIn Position OpDecorate %63 Binding 2
OpDecorate %179 Location 0 OpDecorate %66 BuiltIn LocalInvocationId
OpDecorate %199 Location 0 OpDecorate %113 BuiltIn Position
OpDecorate %181 Location 0
OpDecorate %201 Location 0
%2 = OpTypeVoid %2 = OpTypeVoid
%4 = OpTypeInt 32 1 %4 = OpTypeInt 32 1
%3 = OpConstant %4 10 %3 = OpConstant %4 10
@ -117,189 +121,191 @@ OpDecorate %199 Location 0
%39 = OpVariable %40 UniformConstant %39 = OpVariable %40 UniformConstant
%42 = OpTypePointer UniformConstant %17 %42 = OpTypePointer UniformConstant %17
%41 = OpVariable %42 UniformConstant %41 = OpVariable %42 UniformConstant
%44 = OpTypePointer UniformConstant %20 %44 = OpTypePointer UniformConstant %17
%43 = OpVariable %44 UniformConstant %43 = OpVariable %44 UniformConstant
%46 = OpTypePointer UniformConstant %21 %46 = OpTypePointer UniformConstant %20
%45 = OpVariable %46 UniformConstant %45 = OpVariable %46 UniformConstant
%48 = OpTypePointer UniformConstant %22 %48 = OpTypePointer UniformConstant %21
%47 = OpVariable %48 UniformConstant %47 = OpVariable %48 UniformConstant
%50 = OpTypePointer UniformConstant %23 %50 = OpTypePointer UniformConstant %22
%49 = OpVariable %50 UniformConstant %49 = OpVariable %50 UniformConstant
%52 = OpTypePointer UniformConstant %24 %52 = OpTypePointer UniformConstant %23
%51 = OpVariable %52 UniformConstant %51 = OpVariable %52 UniformConstant
%54 = OpTypePointer UniformConstant %25 %54 = OpTypePointer UniformConstant %24
%53 = OpVariable %54 UniformConstant %53 = OpVariable %54 UniformConstant
%56 = OpTypePointer UniformConstant %26 %56 = OpTypePointer UniformConstant %25
%55 = OpVariable %56 UniformConstant %55 = OpVariable %56 UniformConstant
%58 = OpTypePointer UniformConstant %28 %58 = OpTypePointer UniformConstant %26
%57 = OpVariable %58 UniformConstant %57 = OpVariable %58 UniformConstant
%60 = OpTypePointer UniformConstant %28 %60 = OpTypePointer UniformConstant %28
%59 = OpVariable %60 UniformConstant %59 = OpVariable %60 UniformConstant
%62 = OpTypePointer UniformConstant %29 %62 = OpTypePointer UniformConstant %28
%61 = OpVariable %62 UniformConstant %61 = OpVariable %62 UniformConstant
%65 = OpTypePointer Input %18 %64 = OpTypePointer UniformConstant %29
%64 = OpVariable %65 Input %63 = OpVariable %64 UniformConstant
%68 = OpTypeFunction %2 %67 = OpTypePointer Input %18
%77 = OpTypeVector %12 2 %66 = OpVariable %67 Input
%85 = OpTypeVector %12 4 %70 = OpTypeFunction %2
%100 = OpTypeVector %4 3 %79 = OpTypeVector %12 2
%112 = OpTypePointer Output %27 %87 = OpTypeVector %12 4
%111 = OpVariable %112 Output %102 = OpTypeVector %4 3
%122 = OpConstant %12 0 %114 = OpTypePointer Output %27
%179 = OpVariable %112 Output %113 = OpVariable %114 Output
%184 = OpTypeVector %8 2 %124 = OpConstant %12 0
%186 = OpTypeSampledImage %21 %181 = OpVariable %114 Output
%200 = OpTypePointer Output %8 %186 = OpTypeVector %8 2
%199 = OpVariable %200 Output %188 = OpTypeSampledImage %21
%206 = OpTypeSampledImage %29 %202 = OpTypePointer Output %8
%211 = OpConstant %8 0.0 %201 = OpVariable %202 Output
%67 = OpFunction %2 None %68 %208 = OpTypeSampledImage %29
%63 = OpLabel %213 = OpConstant %8 0.0
%66 = OpLoad %18 %64 %69 = OpFunction %2 None %70
%69 = OpLoad %11 %31 %65 = OpLabel
%70 = OpLoad %13 %33 %68 = OpLoad %18 %66
%71 = OpLoad %14 %35 %71 = OpLoad %11 %31
%72 = OpLoad %15 %37 %72 = OpLoad %13 %33
%73 = OpLoad %16 %39 %73 = OpLoad %14 %35
%74 = OpLoad %17 %41 %74 = OpLoad %15 %37
OpBranch %75 %75 = OpLoad %16 %39
%75 = OpLabel %76 = OpLoad %17 %43
%76 = OpImageQuerySize %19 %72 OpBranch %77
%78 = OpVectorShuffle %77 %66 %66 0 1 %77 = OpLabel
%79 = OpBitcast %19 %78 %78 = OpImageQuerySize %19 %74
%80 = OpIMul %19 %76 %79 %80 = OpVectorShuffle %79 %68 %68 0 1
%81 = OpCompositeConstruct %19 %3 %5 %81 = OpBitcast %19 %80
%82 = OpSMod %19 %80 %81 %82 = OpIMul %19 %78 %81
%83 = OpCompositeExtract %12 %66 2 %83 = OpCompositeConstruct %19 %3 %5
%84 = OpBitcast %4 %83 %84 = OpSMod %19 %82 %83
%86 = OpImageFetch %85 %69 %82 Lod %84 %85 = OpCompositeExtract %12 %68 2
%87 = OpCompositeExtract %12 %66 2 %86 = OpBitcast %4 %85
%88 = OpBitcast %4 %87 %88 = OpImageFetch %87 %71 %84 Lod %86
%89 = OpImageFetch %85 %70 %82 Sample %88 %89 = OpCompositeExtract %12 %68 2
%90 = OpCompositeExtract %12 %66 2 %90 = OpBitcast %4 %89
%91 = OpBitcast %4 %90 %91 = OpImageFetch %87 %72 %84 Sample %90
%92 = OpImageFetch %27 %71 %82 Sample %91 %92 = OpCompositeExtract %12 %68 2
%93 = OpCompositeExtract %8 %92 0 %93 = OpBitcast %4 %92
%94 = OpImageRead %85 %72 %82 %94 = OpImageFetch %27 %73 %84 Sample %93
%95 = OpCompositeExtract %12 %66 2 %95 = OpCompositeExtract %8 %94 0
%96 = OpBitcast %4 %95 %96 = OpImageRead %87 %74 %84
%97 = OpCompositeExtract %12 %66 2 %97 = OpCompositeExtract %12 %68 2
%98 = OpBitcast %4 %97 %98 = OpBitcast %4 %97
%99 = OpIAdd %4 %98 %6 %99 = OpCompositeExtract %12 %68 2
%101 = OpCompositeConstruct %100 %82 %96 %100 = OpBitcast %4 %99
%102 = OpImageFetch %85 %73 %101 Lod %99 %101 = OpIAdd %4 %100 %6
%103 = OpCompositeExtract %4 %82 0 %103 = OpCompositeConstruct %102 %84 %98
%104 = OpIAdd %85 %86 %89 %104 = OpImageFetch %87 %75 %103 Lod %101
%105 = OpConvertFToU %12 %93 %105 = OpCompositeExtract %4 %84 0
%106 = OpCompositeConstruct %85 %105 %105 %105 %105 %106 = OpIAdd %87 %88 %91
%107 = OpIAdd %85 %104 %106 %107 = OpConvertFToU %12 %95
%108 = OpIAdd %85 %107 %94 %108 = OpCompositeConstruct %87 %107 %107 %107 %107
%109 = OpIAdd %85 %108 %102 %109 = OpIAdd %87 %106 %108
OpImageWrite %74 %103 %109 %110 = OpIAdd %87 %109 %96
%111 = OpIAdd %87 %110 %104
OpImageWrite %76 %105 %111
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%113 = OpFunction %2 None %68 %115 = OpFunction %2 None %70
%110 = OpLabel %112 = OpLabel
%114 = OpLoad %20 %43 %116 = OpLoad %20 %45
%115 = OpLoad %21 %45 %117 = OpLoad %21 %47
%116 = OpLoad %22 %47 %118 = OpLoad %22 %49
%117 = OpLoad %23 %49 %119 = OpLoad %23 %51
%118 = OpLoad %24 %51 %120 = OpLoad %24 %53
%119 = OpLoad %25 %53 %121 = OpLoad %25 %55
%120 = OpLoad %26 %55 %122 = OpLoad %26 %57
OpBranch %121 OpBranch %123
%121 = OpLabel %123 = OpLabel
%123 = OpImageQuerySizeLod %4 %114 %122 %125 = OpImageQuerySizeLod %4 %116 %124
%124 = OpImageQuerySizeLod %19 %115 %122 %126 = OpImageQuerySizeLod %19 %117 %124
%125 = OpImageQueryLevels %4 %115 %127 = OpImageQueryLevels %4 %117
%126 = OpImageQuerySizeLod %19 %115 %6 %128 = OpImageQuerySizeLod %19 %117 %6
%127 = OpImageQuerySizeLod %100 %116 %122 %129 = OpImageQuerySizeLod %102 %118 %124
%128 = OpVectorShuffle %19 %127 %127 0 1 %130 = OpVectorShuffle %19 %129 %129 0 1
%129 = OpImageQueryLevels %4 %116 %131 = OpImageQueryLevels %4 %118
%130 = OpImageQuerySizeLod %100 %116 %6 %132 = OpImageQuerySizeLod %102 %118 %6
%131 = OpVectorShuffle %19 %130 %130 0 1 %133 = OpVectorShuffle %19 %132 %132 0 1
%132 = OpImageQuerySizeLod %100 %116 %122 %134 = OpImageQuerySizeLod %102 %118 %124
%133 = OpCompositeExtract %4 %132 2 %135 = OpCompositeExtract %4 %134 2
%134 = OpImageQuerySizeLod %19 %117 %122 %136 = OpImageQuerySizeLod %19 %119 %124
%135 = OpImageQueryLevels %4 %117 %137 = OpImageQueryLevels %4 %119
%136 = OpImageQuerySizeLod %19 %117 %6 %138 = OpImageQuerySizeLod %19 %119 %6
%137 = OpImageQuerySizeLod %100 %118 %122 %139 = OpImageQuerySizeLod %102 %120 %124
%138 = OpVectorShuffle %19 %137 %137 0 0 %140 = OpVectorShuffle %19 %139 %139 0 0
%139 = OpImageQueryLevels %4 %118 %141 = OpImageQueryLevels %4 %120
%140 = OpImageQuerySizeLod %100 %118 %6 %142 = OpImageQuerySizeLod %102 %120 %6
%141 = OpVectorShuffle %19 %140 %140 0 0 %143 = OpVectorShuffle %19 %142 %142 0 0
%142 = OpImageQuerySizeLod %100 %118 %122 %144 = OpImageQuerySizeLod %102 %120 %124
%143 = OpCompositeExtract %4 %142 2 %145 = OpCompositeExtract %4 %144 2
%144 = OpImageQuerySizeLod %100 %119 %122 %146 = OpImageQuerySizeLod %102 %121 %124
%145 = OpImageQueryLevels %4 %119 %147 = OpImageQueryLevels %4 %121
%146 = OpImageQuerySizeLod %100 %119 %6 %148 = OpImageQuerySizeLod %102 %121 %6
%147 = OpImageQuerySamples %4 %120 %149 = OpImageQuerySamples %4 %122
%148 = OpCompositeExtract %4 %124 1
%149 = OpIAdd %4 %123 %148
%150 = OpCompositeExtract %4 %126 1 %150 = OpCompositeExtract %4 %126 1
%151 = OpIAdd %4 %149 %150 %151 = OpIAdd %4 %125 %150
%152 = OpCompositeExtract %4 %128 1 %152 = OpCompositeExtract %4 %128 1
%153 = OpIAdd %4 %151 %152 %153 = OpIAdd %4 %151 %152
%154 = OpCompositeExtract %4 %131 1 %154 = OpCompositeExtract %4 %130 1
%155 = OpIAdd %4 %153 %154 %155 = OpIAdd %4 %153 %154
%156 = OpIAdd %4 %155 %133 %156 = OpCompositeExtract %4 %133 1
%157 = OpCompositeExtract %4 %134 1 %157 = OpIAdd %4 %155 %156
%158 = OpIAdd %4 %156 %157 %158 = OpIAdd %4 %157 %135
%159 = OpCompositeExtract %4 %136 1 %159 = OpCompositeExtract %4 %136 1
%160 = OpIAdd %4 %158 %159 %160 = OpIAdd %4 %158 %159
%161 = OpCompositeExtract %4 %138 1 %161 = OpCompositeExtract %4 %138 1
%162 = OpIAdd %4 %160 %161 %162 = OpIAdd %4 %160 %161
%163 = OpCompositeExtract %4 %141 1 %163 = OpCompositeExtract %4 %140 1
%164 = OpIAdd %4 %162 %163 %164 = OpIAdd %4 %162 %163
%165 = OpIAdd %4 %164 %143 %165 = OpCompositeExtract %4 %143 1
%166 = OpCompositeExtract %4 %144 2 %166 = OpIAdd %4 %164 %165
%167 = OpIAdd %4 %165 %166 %167 = OpIAdd %4 %166 %145
%168 = OpCompositeExtract %4 %146 2 %168 = OpCompositeExtract %4 %146 2
%169 = OpIAdd %4 %167 %168 %169 = OpIAdd %4 %167 %168
%170 = OpIAdd %4 %169 %147 %170 = OpCompositeExtract %4 %148 2
%171 = OpIAdd %4 %170 %125 %171 = OpIAdd %4 %169 %170
%172 = OpIAdd %4 %171 %129 %172 = OpIAdd %4 %171 %149
%173 = OpIAdd %4 %172 %145 %173 = OpIAdd %4 %172 %127
%174 = OpIAdd %4 %173 %135 %174 = OpIAdd %4 %173 %131
%175 = OpIAdd %4 %174 %139 %175 = OpIAdd %4 %174 %147
%176 = OpConvertSToF %8 %175 %176 = OpIAdd %4 %175 %137
%177 = OpCompositeConstruct %27 %176 %176 %176 %176 %177 = OpIAdd %4 %176 %141
OpStore %111 %177 %178 = OpConvertSToF %8 %177
%179 = OpCompositeConstruct %27 %178 %178 %178 %178
OpStore %113 %179
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%180 = OpFunction %2 None %68 %182 = OpFunction %2 None %70
%178 = OpLabel %180 = OpLabel
%181 = OpLoad %21 %45 %183 = OpLoad %21 %47
%182 = OpLoad %28 %57 %184 = OpLoad %28 %59
OpBranch %183 OpBranch %185
%183 = OpLabel %185 = OpLabel
%185 = OpCompositeConstruct %184 %7 %7 %187 = OpCompositeConstruct %186 %7 %7
%187 = OpSampledImage %186 %181 %182 %189 = OpSampledImage %188 %183 %184
%188 = OpImageSampleImplicitLod %27 %187 %185 %190 = OpImageSampleImplicitLod %27 %189 %187
%189 = OpSampledImage %186 %181 %182 %191 = OpSampledImage %188 %183 %184
%190 = OpImageSampleImplicitLod %27 %189 %185 ConstOffset %30 %192 = OpImageSampleImplicitLod %27 %191 %187 ConstOffset %30
%191 = OpSampledImage %186 %181 %182 %193 = OpSampledImage %188 %183 %184
%192 = OpImageSampleExplicitLod %27 %191 %185 Lod %9 %194 = OpImageSampleExplicitLod %27 %193 %187 Lod %9
%193 = OpSampledImage %186 %181 %182 %195 = OpSampledImage %188 %183 %184
%194 = OpImageSampleExplicitLod %27 %193 %185 Lod|ConstOffset %9 %30 %196 = OpImageSampleExplicitLod %27 %195 %187 Lod|ConstOffset %9 %30
%195 = OpFAdd %27 %188 %190 %197 = OpFAdd %27 %190 %192
%196 = OpFAdd %27 %195 %192 %198 = OpFAdd %27 %197 %194
%197 = OpFAdd %27 %196 %194 %199 = OpFAdd %27 %198 %196
OpStore %179 %197 OpStore %181 %199
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd
%201 = OpFunction %2 None %68 %203 = OpFunction %2 None %70
%198 = OpLabel %200 = OpLabel
%202 = OpLoad %28 %59 %204 = OpLoad %28 %61
%203 = OpLoad %29 %61 %205 = OpLoad %29 %63
OpBranch %204 OpBranch %206
%204 = OpLabel %206 = OpLabel
%205 = OpCompositeConstruct %184 %7 %7 %207 = OpCompositeConstruct %186 %7 %7
%207 = OpSampledImage %206 %203 %202 %209 = OpSampledImage %208 %205 %204
%208 = OpImageSampleDrefImplicitLod %8 %207 %205 %7 %210 = OpImageSampleDrefImplicitLod %8 %209 %207 %7
%209 = OpSampledImage %206 %203 %202 %211 = OpSampledImage %208 %205 %204
%210 = OpImageSampleDrefExplicitLod %8 %209 %205 %7 Lod %211 %212 = OpImageSampleDrefExplicitLod %8 %211 %207 %7 Lod %213
%212 = OpFAdd %8 %208 %210 %214 = OpFAdd %8 %210 %212
OpStore %199 %212 OpStore %201 %214
OpReturn OpReturn
OpFunctionEnd OpFunctionEnd

View File

@ -8,6 +8,8 @@ var image_depth_multisampled_src: texture_depth_multisampled_2d;
var image_storage_src: texture_storage_2d<rgba8uint>; var image_storage_src: texture_storage_2d<rgba8uint>;
[[group(0), binding(5)]] [[group(0), binding(5)]]
var image_array_src: texture_2d_array<u32>; var image_array_src: texture_2d_array<u32>;
[[group(0), binding(6)]]
var image_dup_src: texture_storage_1d<r32uint>;
[[group(0), binding(2)]] [[group(0), binding(2)]]
var image_dst: texture_storage_1d<r32uint,write>; var image_dst: texture_storage_1d<r32uint,write>;
[[group(0), binding(0)]] [[group(0), binding(0)]]