From db455af04ba888e197d5831497f84830914248dd Mon Sep 17 00:00:00 2001
From: Dzmitry Malyshau <kvarkus@gmail.com>
Date: Thu, 1 Apr 2021 10:25:17 -0400
Subject: [PATCH] Add a snapshot for image load/store ops

---
 tests/in/image-copy.param.ron |  8 ++++++++
 tests/in/image-copy.wgsl      | 16 ++++++++++++++++
 tests/out/image-copy.msl.snap | 27 +++++++++++++++++++++++++++
 tests/snapshots.rs            |  7 +++++++
 4 files changed, 58 insertions(+)
 create mode 100644 tests/in/image-copy.param.ron
 create mode 100644 tests/in/image-copy.wgsl
 create mode 100644 tests/out/image-copy.msl.snap

diff --git a/tests/in/image-copy.param.ron b/tests/in/image-copy.param.ron
new file mode 100644
index 0000000000..f4cbc5f0bd
--- /dev/null
+++ b/tests/in/image-copy.param.ron
@@ -0,0 +1,8 @@
+(
+	spv_version: (1, 1),
+	spv_capabilities: [ Shader, Image1D, Sampled1D ],
+	mtl_bindings: {
+		(stage: Compute, group: 0, binding: 1): (texture: Some(0), mutable: false),
+		(stage: Compute, group: 0, binding: 2): (texture: Some(1), mutable: true),
+	}
+)
diff --git a/tests/in/image-copy.wgsl b/tests/in/image-copy.wgsl
new file mode 100644
index 0000000000..cb09083b9a
--- /dev/null
+++ b/tests/in/image-copy.wgsl
@@ -0,0 +1,16 @@
+[[group(0), binding(1)]]
+var image_src: [[access(read)]] texture_storage_2d<rgba8uint>;
+[[group(0), binding(2)]]
+var image_dst: [[access(write)]] texture_storage_1d<r32uint>;
+
+[[stage(compute), workgroup_size(16)]]
+fn main(
+    [[builtin(local_invocation_id)]] local_id: vec3<u32>,
+    //TODO: https://github.com/gpuweb/gpuweb/issues/1590
+    //[[builtin(workgroup_size)]] wg_size: vec3<u32>
+) {
+    const dim = textureDimensions(image_src);
+    const itc = dim * vec2<i32>(local_id.xy) % vec2<i32>(10, 20);
+    const value = textureLoad(image_src, itc);
+    textureStore(image_dst, itc.x, value);
+}
diff --git a/tests/out/image-copy.msl.snap b/tests/out/image-copy.msl.snap
new file mode 100644
index 0000000000..81329b6bd3
--- /dev/null
+++ b/tests/out/image-copy.msl.snap
@@ -0,0 +1,27 @@
+---
+source: tests/snapshots.rs
+expression: msl
+---
+#include <metal_stdlib>
+#include <simd/simd.h>
+
+constexpr constant int const_10i = 10;
+constexpr constant int const_20i = 20;
+typedef metal::texture2d<uint, metal::access::read> type;
+typedef metal::texture1d<uint, metal::access::write> type1;
+typedef metal::uint3 type2;
+typedef metal::uint2 type3;
+typedef metal::int2 type4;
+struct main1Input {
+};
+kernel void main1(
+  type2 local_id [[thread_position_in_threadgroup]]
+, type image_src [[texture(0)]]
+, type1 image_dst [[texture(1)]]
+) {
+    metal::int2 _expr12 = (int2(image_src.get_width(), image_src.get_height()) * static_cast<int2>(metal::uint2(local_id.x, local_id.y))) % metal::int2(const_10i, const_20i);
+    metal::uint4 _expr13 = image_src.read(metal::uint2(_expr12));
+    image_dst.write(_expr13, metal::uint(_expr12.x));
+    return;
+}
+
diff --git a/tests/snapshots.rs b/tests/snapshots.rs
index 30def8d2ce..9bddce1ca2 100644
--- a/tests/snapshots.rs
+++ b/tests/snapshots.rs
@@ -264,6 +264,13 @@ fn convert_wgsl_shadow() {
     convert_wgsl("shadow", Targets::SPIRV | Targets::METAL | Targets::GLSL);
 }
 
+#[cfg(feature = "wgsl-in")]
+#[test]
+fn convert_wgsl_image_copy() {
+    //SPIR-V is blocked by https://github.com/gfx-rs/naga/issues/646
+    convert_wgsl("image-copy", Targets::METAL);
+}
+
 #[cfg(feature = "wgsl-in")]
 #[test]
 fn convert_wgsl_texture_array() {