nordic-dev.net/wgpu
nordic-dev.net/wgpu
2
0
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2024-11-25 16:24:24 +00:00
Code Issues Packages Projects Releases Wiki Activity

Add tests for sampling mipmapped and multisampled images.

Browse Source 
Prior to this change, nothing in the test suite exercised the SPIR-V backend
code for generating LOD or Sample arguments to OpImageFetch instructions.
This commit is contained in:
Jim Blandy 2021-07-19 12:47:22 -07:00 committed by Dzmitry Malyshau
parent feee1a2edb
commit f3c7537609
5 changed files with 264 additions and 222 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
tests/in/image.wgsl
View File

@ -1,5 +1,9 @@
[[group(0), binding(0)]]
var image_mipmapped_src: texture_2d<u32>;
[[group(0), binding(3)]]
var image_multisampled_src: texture_multisampled_2d<u32>;
[[group(0), binding(1)]] [[group(0), binding(1)]]
var image_src: [[access(read)]] texture_storage_2d<rgba8uint>; var image_storage_src: [[access(read)]] texture_storage_2d<rgba8uint>;
[[group(0), binding(2)]] [[group(0), binding(2)]]
var image_dst: [[access(write)]] texture_storage_1d<r32uint>; var image_dst: [[access(write)]] texture_storage_1d<r32uint>;
@ -9,10 +13,12 @@ fn main(
//TODO: https://github.com/gpuweb/gpuweb/issues/1590 //TODO: https://github.com/gpuweb/gpuweb/issues/1590
//[[builtin(workgroup_size)]] wg_size: vec3<u32> //[[builtin(workgroup_size)]] wg_size: vec3<u32>
) { ) {
let dim = textureDimensions(image_src); let dim = textureDimensions(image_storage_src);
let itc = dim * vec2<i32>(local_id.xy) % vec2<i32>(10, 20); let itc = dim * vec2<i32>(local_id.xy) % vec2<i32>(10, 20);
let value = textureLoad(image_src, itc); let value1 = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
textureStore(image_dst, itc.x, value); let value2 = textureLoad(image_multisampled_src, itc, i32(local_id.z));
let value3 = textureLoad(image_storage_src, itc);
textureStore(image_dst, itc.x, value1 + value2 + value3);
} }
[[group(0), binding(0)]] [[group(0), binding(0)]]

14
tests/out/hlsl/image.hlsl
View File

@ -1,4 +1,6 @@
Texture2D<uint4> image_src : register(t1); Texture2D<uint4> image_mipmapped_src : register(t0);
Texture2DMS<uint4> image_multisampled_src : register(t3);
Texture2D<uint4> image_storage_src : register(t1);
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);
@ -18,17 +20,19 @@ struct ComputeInput_main {
int2 NagaDimensions2D(Texture2D<uint4>) int2 NagaDimensions2D(Texture2D<uint4>)
{ {
uint4 ret; uint4 ret;
image_src.GetDimensions(0, ret.x, ret.y, ret.z); image_storage_src.GetDimensions(0, ret.x, ret.y, ret.z);
return ret.xy; return ret.xy;
} }
[numthreads(16, 1, 1)] [numthreads(16, 1, 1)]
void main(ComputeInput_main computeinput_main) void main(ComputeInput_main computeinput_main)
{ {
int2 dim = NagaDimensions2D(image_src); int2 dim = NagaDimensions2D(image_storage_src);
int2 itc = (mul(dim, int2(computeinput_main.local_id1.xy)) % int2(10, 20)); int2 itc = (mul(dim, int2(computeinput_main.local_id1.xy)) % int2(10, 20));
uint4 value = image_src.Load(int3(itc, 0)); uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(computeinput_main.local_id1.z)));
image_dst[itc.x] = value; uint4 value2_ = image_multisampled_src.Load(itc, int(computeinput_main.local_id1.z));
uint4 value3_ = image_storage_src.Load(int3(itc, 0));
image_dst[itc.x] = ((value1_ + value2_) + value3_);
return; return;
} }

18
tests/out/msl/image.msl
View File

@ -2,19 +2,23 @@
#include <metal_stdlib> #include <metal_stdlib>
#include <simd/simd.h> #include <simd/simd.h>
constant metal::int2 const_type3_ = {3, 1}; constant metal::int2 const_type5_ = {3, 1};
struct main1Input { struct main1Input {
}; };
kernel void main1( kernel void main1(
metal::uint3 local_id [[thread_position_in_threadgroup]] metal::uint3 local_id [[thread_position_in_threadgroup]]
, metal::texture2d<uint, metal::access::read> image_src [[user(fake0)]] , metal::texture2d<uint, metal::access::sample> image_mipmapped_src [[user(fake0)]]
, metal::texture2d_ms<uint, metal::access::read> image_multisampled_src [[user(fake0)]]
, metal::texture2d<uint, metal::access::read> image_storage_src [[user(fake0)]]
, metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]] , metal::texture1d<uint, metal::access::write> image_dst [[user(fake0)]]
) { ) {
metal::int2 dim = int2(image_src.get_width(), image_src.get_height()); metal::int2 dim = int2(image_storage_src.get_width(), image_storage_src.get_height());
metal::int2 itc = (dim * static_cast<int2>(local_id.xy)) % metal::int2(10, 20); metal::int2 itc = (dim * static_cast<int2>(local_id.xy)) % metal::int2(10, 20);
metal::uint4 value = image_src.read(metal::uint2(itc)); metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
image_dst.write(value, metal::uint(itc.x)); metal::uint4 value2_ = image_multisampled_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
metal::uint4 value3_ = image_storage_src.read(metal::uint2(itc));
image_dst.write((value1_ + value2_) + value3_, metal::uint(itc.x));
return; return;
} }
@ -64,9 +68,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_type3_); metal::float4 s2d_offset = image_2d.sample(sampler_reg, tc, const_type5_);
metal::float4 s2d_level = image_2d.sample(sampler_reg, tc, metal::level(2.3)); metal::float4 s2d_level = image_2d.sample(sampler_reg, tc, metal::level(2.3));
metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.3), const_type3_); metal::float4 s2d_level_offset = image_2d.sample(sampler_reg, tc, metal::level(2.3), const_type5_);
return sampleOutput { ((s2d + s2d_offset) + s2d_level) + s2d_level_offset }; return sampleOutput { ((s2d + s2d_offset) + s2d_level) + s2d_level_offset };
} }

426
tests/out/spv/image.spvasm
View File

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

14
tests/out/wgsl/image.wgsl
View File

@ -1,5 +1,9 @@
[[group(0), binding(0)]]
var image_mipmapped_src: texture_2d<u32>;
[[group(0), binding(3)]]
var image_multisampled_src: texture_multisampled_2d<u32>;
[[group(0), binding(1)]] [[group(0), binding(1)]]
var image_src: [[access(read)]] texture_storage_2d<rgba8uint>; var image_storage_src: [[access(read)]] texture_storage_2d<rgba8uint>;
[[group(0), binding(2)]] [[group(0), binding(2)]]
var image_dst: [[access(write)]] texture_storage_1d<r32uint>; var image_dst: [[access(write)]] texture_storage_1d<r32uint>;
[[group(0), binding(0)]] [[group(0), binding(0)]]
@ -25,10 +29,12 @@ var image_2d_depth: texture_depth_2d;
[[stage(compute), workgroup_size(16, 1, 1)]] [[stage(compute), workgroup_size(16, 1, 1)]]
fn main([[builtin(local_invocation_id)]] local_id: vec3<u32>) { fn main([[builtin(local_invocation_id)]] local_id: vec3<u32>) {
let dim: vec2<i32> = textureDimensions(image_src); let dim: vec2<i32> = textureDimensions(image_storage_src);
let itc: vec2<i32> = ((dim * vec2<i32>(local_id.xy)) % vec2<i32>(10, 20)); let itc: vec2<i32> = ((dim * vec2<i32>(local_id.xy)) % vec2<i32>(10, 20));
let value: vec4<u32> = textureLoad(image_src, itc); let value1_: vec4<u32> = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
textureStore(image_dst, itc.x, value); let value2_: vec4<u32> = textureLoad(image_multisampled_src, itc, i32(local_id.z));
let value3_: vec4<u32> = textureLoad(image_storage_src, itc);
textureStore(image_dst, itc.x, ((value1_ + value2_) + value3_));
return; return;
} }