Add tests for sampling mipmapped and multisampled images.

Prior to this change, nothing in the test suite exercised the SPIR-V backend
code for generating LOD or Sample arguments to OpImageFetch instructions.

tests/in/image.wgsl

@@ -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)]]
-var image_src: [[access(read)]] texture_storage_2d<rgba8uint>;
+var image_storage_src: [[access(read)]] texture_storage_2d<rgba8uint>;
 [[group(0), binding(2)]]
 var image_dst: [[access(write)]] texture_storage_1d<r32uint>;
 
@@ -9,10 +13,12 @@ fn main(
     //TODO: https://github.com/gpuweb/gpuweb/issues/1590
     //[[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 value = textureLoad(image_src, itc);
-    textureStore(image_dst, itc.x, value);
+    let value1 = textureLoad(image_mipmapped_src, itc, i32(local_id.z));
+    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)]]

tests/out/hlsl/image.hlsl

@@ -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);
 Texture1D<float4> image_1d : register(t0);
 Texture2D<float4> image_2d : register(t1);
@@ -18,17 +20,19 @@ struct ComputeInput_main {
 int2 NagaDimensions2D(Texture2D<uint4>)
 {
     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;
 }
 
 [numthreads(16, 1, 1)]
 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));
-    uint4 value = image_src.Load(int3(itc, 0));
-    image_dst[itc.x] = value;
+    uint4 value1_ = image_mipmapped_src.Load(int3(itc, int(computeinput_main.local_id1.z)));
+    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;
 }
 

tests/out/msl/image.msl

@@ -2,19 +2,23 @@
 #include <metal_stdlib>
 #include <simd/simd.h>
 
-constant metal::int2 const_type3_ = {3, 1};
+constant metal::int2 const_type5_ = {3, 1};
 
 struct main1Input {
 };
 kernel void main1(
   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::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::uint4 value = image_src.read(metal::uint2(itc));
-    image_dst.write(value, metal::uint(itc.x));
+    metal::uint4 value1_ = image_mipmapped_src.read(metal::uint2(itc), static_cast<int>(local_id.z));
+    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;
 }
 
@@ -64,9 +68,9 @@ fragment sampleOutput sample(
 ) {
     metal::float2 tc = metal::float2(0.5);
     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_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 };
 }