github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-04-22 03:02:01 -04:00
Code Issues Packages Projects Releases Wiki Activity

Add a snapshot for image load/store ops

Browse Source
This commit is contained in:
Dzmitry Malyshau
2021-04-01 10:25:17 -04:00
committed by Dzmitry Malyshau
parent 6a5c7f10ad
commit db455af04b
4 changed files with 58 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
tests/in/image-copy.param.ron Normal file
View File

@@ -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),
}
)

16
tests/in/image-copy.wgsl Normal file
View File

@@ -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);
}

27
tests/out/image-copy.msl.snap Normal file
View File

@@ -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;
}

7
tests/snapshots.rs
View File

@@ -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() {