github/wgpu
1
1
Fork 0
mirror of https://github.com/gfx-rs/wgpu.git synced 2026-01-09 14:48:08 -05:00
Code Issues Packages Projects Releases Wiki Activity

Bindless Tests (#6732)

Browse Source 
* Move Partial Binding into Own File

* Texture Bindless Test

* Make It Work

* Tests

* Uniform Buffers

* BadCode

* Bugs!

* Exclude llvmpipe

* Combine Partial Binding Test

* MVK Issue

* Sampler Array Tests

* Make All Tests Partially Bound As Well
This commit is contained in:
Connor Fitzgerald
2024-12-16 04:08:05 -05:00
committed by GitHub
parent 2ec87d7c2f
commit 286082acef
9 changed files with 959 additions and 114 deletions
Download Patch File Download Diff File Expand all files Collapse all files

1
tests/Cargo.toml
View File

@@ -34,6 +34,7 @@ ctor.workspace = true
futures-lite.workspace = true
glam.workspace = true
itertools.workspace = true
image.workspace = true
libtest-mimic.workspace = true
log.workspace = true
parking_lot.workspace = true

265
tests/tests/binding_array/buffers.rs Normal file
View File

@@ -0,0 +1,265 @@
use std::num::{NonZeroU32, NonZeroU64};
use wgpu::*;
use wgpu_test::{gpu_test, FailureCase, GpuTestConfiguration, TestParameters, TestingContext};
#[gpu_test]
static BINDING_ARRAY_UNIFORM_BUFFERS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::BUFFER_BINDING_ARRAY
| Features::UNIFORM_BUFFER_AND_STORAGE_TEXTURE_ARRAY_NON_UNIFORM_INDEXING,
)
.limits(Limits {
max_uniform_buffers_per_shader_stage: 16,
..Limits::default()
})
// Naga bug on vulkan: https://github.com/gfx-rs/wgpu/issues/6733
//
// Causes varying errors on different devices, so we don't match more closely.
.expect_fail(FailureCase::backend(Backends::VULKAN))
// These issues cause a segfault on lavapipe
.skip(FailureCase::backend_adapter(Backends::VULKAN, "llvmpipe")),
)
.run_async(|ctx| async move { binding_array_buffers(ctx, BufferType::Uniform, false).await });
#[gpu_test]
static PARTIAL_BINDING_ARRAY_UNIFORM_BUFFERS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::BUFFER_BINDING_ARRAY
| Features::PARTIALLY_BOUND_BINDING_ARRAY
| Features::UNIFORM_BUFFER_AND_STORAGE_TEXTURE_ARRAY_NON_UNIFORM_INDEXING,
)
.limits(Limits {
max_uniform_buffers_per_shader_stage: 32,
..Limits::default()
})
// Naga bug on vulkan: https://github.com/gfx-rs/wgpu/issues/6733
//
// Causes varying errors on different devices, so we don't match more closely.
.expect_fail(FailureCase::backend(Backends::VULKAN))
// These issues cause a segfault on lavapipe
.skip(FailureCase::backend_adapter(Backends::VULKAN, "llvmpipe")),
)
.run_async(|ctx| async move { binding_array_buffers(ctx, BufferType::Uniform, true).await });
#[gpu_test]
static BINDING_ARRAY_STORAGE_BUFFERS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::BUFFER_BINDING_ARRAY
| Features::STORAGE_RESOURCE_BINDING_ARRAY
| Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
)
.limits(Limits {
max_storage_buffers_per_shader_stage: 17,
..Limits::default()
})
// See https://github.com/gfx-rs/wgpu/issues/6745.
.expect_fail(FailureCase::molten_vk()),
)
.run_async(|ctx| async move { binding_array_buffers(ctx, BufferType::Storage, false).await });
#[gpu_test]
static PARTIAL_BINDING_ARRAY_STORAGE_BUFFERS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::BUFFER_BINDING_ARRAY
| Features::PARTIALLY_BOUND_BINDING_ARRAY
| Features::STORAGE_RESOURCE_BINDING_ARRAY
| Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
)
.limits(Limits {
max_storage_buffers_per_shader_stage: 33,
..Limits::default()
})
// See https://github.com/gfx-rs/wgpu/issues/6745.
.expect_fail(FailureCase::molten_vk()),
)
.run_async(|ctx| async move { binding_array_buffers(ctx, BufferType::Storage, true).await });
enum BufferType {
Storage,
Uniform,
}
async fn binding_array_buffers(
ctx: TestingContext,
buffer_type: BufferType,
partial_binding: bool,
) {
let storage_mode = match buffer_type {
BufferType::Storage => "storage",
BufferType::Uniform => "uniform",
};
let shader = r#"
struct ImAU32 {
value: u32,
_padding: u32,
_padding2: u32,
_padding3: u32,
};
@group(0) @binding(0)
var<{storage_mode}> buffers: binding_array<ImAU32>;
@group(0) @binding(1)
var<storage, read_write> output_buffer: array<u32>;
@compute
@workgroup_size(16, 1, 1)
fn compMain(@builtin(global_invocation_id) id: vec3u) {
output_buffer[id.x] = buffers[id.x].value;
}
"#;
let shader = shader.replace("{storage_mode}", storage_mode);
let module = ctx
.device
.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Binding Array Buffer"),
source: wgpu::ShaderSource::Wgsl(shader.into()),
});
let image = image::load_from_memory(include_bytes!("../3x3_colors.png")).unwrap();
// Resize image to 4x4
let image = image
.resize_exact(4, 4, image::imageops::FilterType::Gaussian)
.into_rgba8();
// Create one buffer for each pixel
let mut buffers = Vec::with_capacity(64);
for data in image.pixels() {
let buffer = ctx.device.create_buffer(&BufferDescriptor {
label: None,
usage: match buffer_type {
BufferType::Storage => BufferUsages::STORAGE | BufferUsages::COPY_DST,
BufferType::Uniform => BufferUsages::UNIFORM | BufferUsages::COPY_DST,
},
// 16 to allow padding for uniform buffers
size: 16,
mapped_at_creation: true,
});
buffer.slice(..).get_mapped_range_mut()[0..4].copy_from_slice(&data.0);
buffer.unmap();
buffers.push(buffer);
}
let output_buffer = ctx.device.create_buffer(&BufferDescriptor {
label: None,
size: 4 * 4 * 4,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let multiplier = if partial_binding { 2 } else { 1 };
let bind_group_layout = ctx
.device
.create_bind_group_layout(&BindGroupLayoutDescriptor {
label: Some("Bind Group Layout"),
entries: &[
BindGroupLayoutEntry {
binding: 0,
visibility: ShaderStages::COMPUTE,
ty: BindingType::Buffer {
ty: match buffer_type {
BufferType::Storage => BufferBindingType::Storage { read_only: true },
BufferType::Uniform => BufferBindingType::Uniform,
},
has_dynamic_offset: false,
min_binding_size: Some(NonZeroU64::new(16).unwrap()),
},
count: Some(NonZeroU32::new(16 * multiplier).unwrap()),
},
BindGroupLayoutEntry {
binding: 1,
visibility: ShaderStages::COMPUTE,
ty: BindingType::Buffer {
ty: BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: Some(NonZeroU64::new(4).unwrap()),
},
count: None,
},
],
});
let buffer_references: Vec<_> = buffers
.iter()
.map(|b| b.as_entire_buffer_binding())
.collect();
let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
label: Some("Bind Group"),
layout: &bind_group_layout,
entries: &[
BindGroupEntry {
binding: 0,
resource: BindingResource::BufferArray(&buffer_references),
},
BindGroupEntry {
binding: 1,
resource: output_buffer.as_entire_binding(),
},
],
});
let pipeline_layout = ctx
.device
.create_pipeline_layout(&PipelineLayoutDescriptor {
label: Some("Pipeline Layout"),
bind_group_layouts: &[&bind_group_layout],
push_constant_ranges: &[],
});
let pipeline = ctx
.device
.create_compute_pipeline(&ComputePipelineDescriptor {
label: Some("Compute Pipeline"),
layout: Some(&pipeline_layout),
module: &module,
entry_point: Some("compMain"),
compilation_options: Default::default(),
cache: None,
});
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor { label: None });
{
let mut render_pass = encoder.begin_compute_pass(&ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
render_pass.set_pipeline(&pipeline);
render_pass.set_bind_group(0, &bind_group, &[]);
render_pass.dispatch_workgroups(1, 1, 1);
}
let readback_buffer = ctx.device.create_buffer(&BufferDescriptor {
label: None,
size: 4 * 4 * 4,
usage: BufferUsages::MAP_READ | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
encoder.copy_buffer_to_buffer(&output_buffer, 0, &readback_buffer, 0, 4 * 4 * 4);
ctx.queue.submit(Some(encoder.finish()));
let slice = readback_buffer.slice(..);
slice.map_async(MapMode::Read, |_| {});
ctx.device.poll(Maintain::Wait);
let data = slice.get_mapped_range();
assert_eq!(&data[..], &*image);
}

4
tests/tests/binding_array/mod.rs Normal file
View File

@@ -0,0 +1,4 @@
mod buffers;
mod sampled_textures;
mod samplers;
mod storage_textures;

234
tests/tests/binding_array/sampled_textures.rs Normal file
View File

@@ -0,0 +1,234 @@
use std::num::NonZeroU32;
use wgpu::*;
use wgpu_test::{
gpu_test, image::ReadbackBuffers, GpuTestConfiguration, TestParameters, TestingContext,
};
#[gpu_test]
static BINDING_ARRAY_SAMPLED_TEXTURES: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::TEXTURE_BINDING_ARRAY
| Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
)
.limits(Limits {
max_sampled_textures_per_shader_stage: 16,
..Limits::default()
}),
)
.run_async(|ctx| async move { binding_array_sampled_textures(ctx, false).await });
#[gpu_test]
static PARTIAL_BINDING_ARRAY_SAMPLED_TEXTURES: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::TEXTURE_BINDING_ARRAY
| Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING
| Features::PARTIALLY_BOUND_BINDING_ARRAY,
)
.limits(Limits {
max_sampled_textures_per_shader_stage: 32,
..Limits::default()
}),
)
.run_async(|ctx| async move { binding_array_sampled_textures(ctx, false).await });
/// Test to see how texture bindings array work and additionally making sure
/// that non-uniform indexing is working correctly.
///
/// If non-uniform indexing is not working correctly, AMD will produce the wrong
/// output due to non-native support for non-uniform indexing within a WARP.
async fn binding_array_sampled_textures(ctx: TestingContext, partially_bound: bool) {
let shader = r#"
@group(0) @binding(0)
var textures: binding_array<texture_2d<f32>>;
@vertex
fn vertMain(@builtin(vertex_index) id: u32) -> @builtin(position) vec4f {
var positions = array<vec2f, 3>(
vec2f(-1.0, -1.0),
vec2f(3.0, -1.0),
vec2f(-1.0, 3.0)
);
return vec4<f32>(positions[id], 0.0, 1.0);
}
@fragment
fn fragMain(@builtin(position) pos: vec4f) -> @location(0) vec4f {
let pixel = vec2u(floor(pos.xy));
let index = pixel.y * 4 + pixel.x;
return textureLoad(textures[index], vec2u(0), 0);
}
"#;
let module = ctx
.device
.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Binding Array Texture"),
source: wgpu::ShaderSource::Wgsl(shader.into()),
});
let image = image::load_from_memory(include_bytes!("../3x3_colors.png")).unwrap();
// Resize image to 4x4
let image = image
.resize_exact(4, 4, image::imageops::FilterType::Gaussian)
.into_rgba8();
// Create one texture for each pixel
let mut input_views = Vec::with_capacity(64);
for data in image.pixels() {
let texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
label: None,
size: Extent3d {
width: 1,
height: 1,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: TextureDimension::D2,
format: TextureFormat::Rgba8UnormSrgb,
usage: TextureUsages::TEXTURE_BINDING | TextureUsages::COPY_DST,
view_formats: &[],
});
ctx.queue.write_texture(
TexelCopyTextureInfo {
texture: &texture,
mip_level: 0,
origin: Origin3d::ZERO,
aspect: TextureAspect::All,
},
&data.0,
TexelCopyBufferLayout {
offset: 0,
bytes_per_row: Some(4),
rows_per_image: Some(1),
},
Extent3d {
width: 1,
height: 1,
depth_or_array_layers: 1,
},
);
input_views.push(texture.create_view(&TextureViewDescriptor::default()));
}
let output_texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
label: Some("Output Texture"),
size: Extent3d {
width: 4,
height: 4,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: TextureDimension::D2,
format: TextureFormat::Rgba8UnormSrgb,
usage: TextureUsages::RENDER_ATTACHMENT | TextureUsages::COPY_SRC,
view_formats: &[],
});
let output_view = output_texture.create_view(&TextureViewDescriptor::default());
let count = if partially_bound { 32 } else { 16 };
let bind_group_layout = ctx
.device
.create_bind_group_layout(&BindGroupLayoutDescriptor {
label: Some("Bind Group Layout"),
entries: &[BindGroupLayoutEntry {
binding: 0,
visibility: ShaderStages::FRAGMENT,
ty: BindingType::Texture {
sample_type: TextureSampleType::Float { filterable: false },
view_dimension: TextureViewDimension::D2,
multisampled: false,
},
count: Some(NonZeroU32::new(count).unwrap()),
}],
});
let input_view_references: Vec<_> = input_views.iter().collect();
let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
label: Some("Bind Group"),
layout: &bind_group_layout,
entries: &[BindGroupEntry {
binding: 0,
resource: BindingResource::TextureViewArray(&input_view_references),
}],
});
let pipeline_layout = ctx
.device
.create_pipeline_layout(&PipelineLayoutDescriptor {
label: Some("Pipeline Layout"),
bind_group_layouts: &[&bind_group_layout],
push_constant_ranges: &[],
});
let pipeline = ctx
.device
.create_render_pipeline(&RenderPipelineDescriptor {
label: Some("Render Pipeline"),
layout: Some(&pipeline_layout),
vertex: VertexState {
module: &module,
entry_point: Some("vertMain"),
buffers: &[],
compilation_options: PipelineCompilationOptions::default(),
},
fragment: Some(FragmentState {
module: &module,
entry_point: Some("fragMain"),
targets: &[Some(ColorTargetState {
format: TextureFormat::Rgba8UnormSrgb,
blend: None,
write_mask: ColorWrites::ALL,
})],
compilation_options: PipelineCompilationOptions::default(),
}),
primitive: PrimitiveState::default(),
depth_stencil: None,
multisample: MultisampleState::default(),
cache: None,
multiview: None,
});
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor { label: None });
{
let mut render_pass = encoder.begin_render_pass(&RenderPassDescriptor {
label: Some("Render Pass"),
color_attachments: &[Some(RenderPassColorAttachment {
view: &output_view,
resolve_target: None,
ops: Operations {
load: LoadOp::Clear(Color::BLACK),
store: StoreOp::Store,
},
})],
depth_stencil_attachment: None,
timestamp_writes: None,
occlusion_query_set: None,
});
render_pass.set_pipeline(&pipeline);
render_pass.set_bind_group(0, &bind_group, &[]);
render_pass.draw(0..3, 0..1);
}
let readback_buffers = ReadbackBuffers::new(&ctx.device, &output_texture);
readback_buffers.copy_from(&ctx.device, &mut encoder, &output_texture);
ctx.queue.submit(Some(encoder.finish()));
readback_buffers.assert_buffer_contents(&ctx, &image).await;
}

251
tests/tests/binding_array/samplers.rs Normal file
View File

@@ -0,0 +1,251 @@
use std::num::{NonZeroU32, NonZeroU64};
use wgpu::*;
use wgpu_test::{gpu_test, GpuTestConfiguration, TestParameters, TestingContext};
#[gpu_test]
static BINDING_ARRAY_SAMPLERS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::TEXTURE_BINDING_ARRAY
| Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING,
)
.limits(Limits {
max_samplers_per_shader_stage: 2,
..Limits::default()
}),
)
.run_async(|ctx| async move { binding_array_samplers(ctx, false).await });
#[gpu_test]
static PARTIAL_BINDING_ARRAY_SAMPLERS: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::TEXTURE_BINDING_ARRAY
| Features::SAMPLED_TEXTURE_AND_STORAGE_BUFFER_ARRAY_NON_UNIFORM_INDEXING
| Features::PARTIALLY_BOUND_BINDING_ARRAY,
)
.limits(Limits {
max_samplers_per_shader_stage: 4,
..Limits::default()
}),
)
.run_async(|ctx| async move { binding_array_samplers(ctx, true).await });
async fn binding_array_samplers(ctx: TestingContext, partially_bound: bool) {
let shader = r#"
@group(0) @binding(0)
var samplers: binding_array<sampler>;
@group(0) @binding(1)
var texture: texture_2d<f32>;
@group(0) @binding(2)
var<storage, read_write> output_values: array<u32>;
@compute
@workgroup_size(2, 1, 1)
fn compMain(@builtin(global_invocation_id) id: vec3u) {
output_values[id.x] = pack4x8unorm(textureSampleLevel(texture, samplers[id.x], vec2f(0.25 + (0.5 * 0.25), 0.5), 0.0));
}
"#;
let module = ctx
.device
.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Binding Array Texture"),
source: wgpu::ShaderSource::Wgsl(shader.into()),
});
let input_image: [u8; 8] = [
255, 0, 0, 255, //
0, 255, 0, 255, //
];
let expected_output: [u8; 8] = [
191, 64, 0, 255, //
255, 0, 0, 255, //
];
let texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
label: None,
size: Extent3d {
width: 2,
height: 1,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: TextureDimension::D2,
format: TextureFormat::Rgba8Unorm,
usage: TextureUsages::TEXTURE_BINDING | TextureUsages::COPY_DST,
view_formats: &[],
});
ctx.queue.write_texture(
TexelCopyTextureInfo {
texture: &texture,
mip_level: 0,
origin: Origin3d::ZERO,
aspect: TextureAspect::All,
},
&input_image,
TexelCopyBufferLayout {
offset: 0,
bytes_per_row: Some(8),
rows_per_image: Some(1),
},
Extent3d {
width: 2,
height: 1,
depth_or_array_layers: 1,
},
);
let input_view = texture.create_view(&TextureViewDescriptor::default());
let samplers = [
ctx.device.create_sampler(&SamplerDescriptor {
label: None,
address_mode_u: AddressMode::ClampToEdge,
address_mode_v: AddressMode::ClampToEdge,
address_mode_w: AddressMode::ClampToEdge,
mag_filter: FilterMode::Linear,
min_filter: FilterMode::Linear,
mipmap_filter: FilterMode::Linear,
lod_min_clamp: 0.0,
lod_max_clamp: 1000.0,
compare: None,
anisotropy_clamp: 1,
border_color: None,
}),
ctx.device.create_sampler(&SamplerDescriptor {
label: None,
address_mode_u: AddressMode::ClampToEdge,
address_mode_v: AddressMode::ClampToEdge,
address_mode_w: AddressMode::ClampToEdge,
mag_filter: FilterMode::Nearest,
min_filter: FilterMode::Nearest,
mipmap_filter: FilterMode::Nearest,
lod_min_clamp: 0.0,
lod_max_clamp: 1000.0,
compare: None,
anisotropy_clamp: 1,
border_color: None,
}),
];
let output_buffer = ctx.device.create_buffer(&BufferDescriptor {
label: None,
size: 4 * 2,
usage: BufferUsages::STORAGE | BufferUsages::COPY_SRC,
mapped_at_creation: false,
});
let multiplier = if partially_bound { 2 } else { 1 };
let bind_group_layout = ctx
.device
.create_bind_group_layout(&BindGroupLayoutDescriptor {
label: Some("Bind Group Layout"),
entries: &[
BindGroupLayoutEntry {
binding: 0,
visibility: ShaderStages::COMPUTE,
ty: BindingType::Sampler(SamplerBindingType::Filtering),
count: Some(NonZeroU32::new(2 * multiplier).unwrap()),
},
BindGroupLayoutEntry {
binding: 1,
visibility: ShaderStages::COMPUTE,
ty: BindingType::Texture {
sample_type: wgpu::TextureSampleType::Float { filterable: true },
view_dimension: wgpu::TextureViewDimension::D2,
multisampled: false,
},
count: None,
},
BindGroupLayoutEntry {
binding: 2,
visibility: ShaderStages::COMPUTE,
ty: BindingType::Buffer {
ty: BufferBindingType::Storage { read_only: false },
has_dynamic_offset: false,
min_binding_size: Some(NonZeroU64::new(4).unwrap()),
},
count: None,
},
],
});
let sampler_references: Vec<_> = samplers.iter().collect();
let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
label: Some("Bind Group"),
layout: &bind_group_layout,
entries: &[
BindGroupEntry {
binding: 0,
resource: BindingResource::SamplerArray(&sampler_references),
},
BindGroupEntry {
binding: 1,
resource: BindingResource::TextureView(&input_view),
},
BindGroupEntry {
binding: 2,
resource: output_buffer.as_entire_binding(),
},
],
});
let pipeline_layout = ctx
.device
.create_pipeline_layout(&PipelineLayoutDescriptor {
label: Some("Pipeline Layout"),
bind_group_layouts: &[&bind_group_layout],
push_constant_ranges: &[],
});
let pipeline = ctx
.device
.create_compute_pipeline(&ComputePipelineDescriptor {
label: Some("Compute Pipeline"),
layout: Some(&pipeline_layout),
module: &module,
entry_point: Some("compMain"),
compilation_options: Default::default(),
cache: None,
});
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor { label: None });
{
let mut render_pass = encoder.begin_compute_pass(&ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
render_pass.set_pipeline(&pipeline);
render_pass.set_bind_group(0, &bind_group, &[]);
render_pass.dispatch_workgroups(1, 1, 1);
}
let readback_buffer = ctx.device.create_buffer(&BufferDescriptor {
label: None,
size: 4 * 2,
usage: BufferUsages::MAP_READ | BufferUsages::COPY_DST,
mapped_at_creation: false,
});
encoder.copy_buffer_to_buffer(&output_buffer, 0, &readback_buffer, 0, 4 * 2);
ctx.queue.submit(Some(encoder.finish()));
readback_buffer.slice(..).map_async(MapMode::Read, |_| {});
ctx.device.poll(Maintain::Wait);
let readback_buffer_slice = readback_buffer.slice(..).get_mapped_range();
assert_eq!(&readback_buffer_slice[0..8], &expected_output[..]);
}

203
tests/tests/binding_array/storage_textures.rs Normal file
View File

@@ -0,0 +1,203 @@
use std::num::NonZeroU32;
use wgpu::*;
use wgpu_test::{
gpu_test, image::ReadbackBuffers, FailureCase, GpuTestConfiguration, TestParameters,
TestingContext,
};
#[gpu_test]
static BINDING_ARRAY_STORAGE_TEXTURES: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::TEXTURE_BINDING_ARRAY
| Features::STORAGE_RESOURCE_BINDING_ARRAY
| Features::UNIFORM_BUFFER_AND_STORAGE_TEXTURE_ARRAY_NON_UNIFORM_INDEXING
| Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES,
)
.limits(Limits {
max_storage_textures_per_shader_stage: 17,
..Limits::default()
})
.expect_fail(FailureCase::backend(Backends::METAL)),
)
.run_async(|ctx| async move { binding_array_storage_textures(ctx, false).await });
#[gpu_test]
static PARTIAL_BINDING_ARRAY_STORAGE_TEXTURES: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
Features::TEXTURE_BINDING_ARRAY
| Features::PARTIALLY_BOUND_BINDING_ARRAY
| Features::STORAGE_RESOURCE_BINDING_ARRAY
| Features::UNIFORM_BUFFER_AND_STORAGE_TEXTURE_ARRAY_NON_UNIFORM_INDEXING
| Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES,
)
.limits(Limits {
max_storage_textures_per_shader_stage: 33,
..Limits::default()
})
.expect_fail(FailureCase::backend(Backends::METAL)),
)
.run_async(|ctx| async move { binding_array_storage_textures(ctx, true).await });
async fn binding_array_storage_textures(ctx: TestingContext, partially_bound: bool) {
let shader = r#"
@group(0) @binding(0)
var textures: binding_array<texture_storage_2d<rgba8unorm, read_write> >;
@compute
@workgroup_size(4, 4, 1)
fn compMain(@builtin(global_invocation_id) id: vec3u) {
// Read from the 4x4 textures in 0-15, then write to the 4x4 texture in 16
let pixel = vec2u(id.xy);
let index = pixel.y * 4 + pixel.x;
let color = textureLoad(textures[index], vec2u(0));
textureStore(textures[16], pixel, color);
}
"#;
let module = ctx
.device
.create_shader_module(wgpu::ShaderModuleDescriptor {
label: Some("Binding Array Texture"),
source: wgpu::ShaderSource::Wgsl(shader.into()),
});
let image = image::load_from_memory(include_bytes!("../3x3_colors.png")).unwrap();
// Resize image to 4x4
let image = image
.resize_exact(4, 4, image::imageops::FilterType::Gaussian)
.into_rgba8();
// Create one texture for each pixel
let mut input_views = Vec::with_capacity(64);
for data in image.pixels() {
let texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
label: None,
size: Extent3d {
width: 1,
height: 1,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: TextureDimension::D2,
format: TextureFormat::Rgba8Unorm,
usage: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_DST,
view_formats: &[],
});
ctx.queue.write_texture(
TexelCopyTextureInfo {
texture: &texture,
mip_level: 0,
origin: Origin3d::ZERO,
aspect: TextureAspect::All,
},
&data.0,
TexelCopyBufferLayout {
offset: 0,
bytes_per_row: Some(4),
rows_per_image: Some(1),
},
Extent3d {
width: 1,
height: 1,
depth_or_array_layers: 1,
},
);
input_views.push(texture.create_view(&TextureViewDescriptor::default()));
}
let output_texture = ctx.device.create_texture(&wgpu::TextureDescriptor {
label: Some("Output Texture"),
size: Extent3d {
width: 4,
height: 4,
depth_or_array_layers: 1,
},
mip_level_count: 1,
sample_count: 1,
dimension: TextureDimension::D2,
format: TextureFormat::Rgba8Unorm,
usage: TextureUsages::STORAGE_BINDING | TextureUsages::COPY_SRC,
view_formats: &[],
});
let output_view = output_texture.create_view(&TextureViewDescriptor::default());
let multiplier = if partially_bound { 2 } else { 1 };
let bind_group_layout = ctx
.device
.create_bind_group_layout(&BindGroupLayoutDescriptor {
label: Some("Bind Group Layout"),
entries: &[BindGroupLayoutEntry {
binding: 0,
visibility: ShaderStages::COMPUTE,
ty: BindingType::StorageTexture {
access: StorageTextureAccess::ReadWrite,
format: TextureFormat::Rgba8Unorm,
view_dimension: TextureViewDimension::D2,
},
count: Some(NonZeroU32::new(4 * 4 * multiplier + 1).unwrap()),
}],
});
let mut input_view_references: Vec<_> = input_views.iter().collect();
input_view_references.push(&output_view);
let bind_group = ctx.device.create_bind_group(&BindGroupDescriptor {
label: Some("Bind Group"),
layout: &bind_group_layout,
entries: &[BindGroupEntry {
binding: 0,
resource: BindingResource::TextureViewArray(&input_view_references),
}],
});
let pipeline_layout = ctx
.device
.create_pipeline_layout(&PipelineLayoutDescriptor {
label: Some("Pipeline Layout"),
bind_group_layouts: &[&bind_group_layout],
push_constant_ranges: &[],
});
let pipeline = ctx
.device
.create_compute_pipeline(&ComputePipelineDescriptor {
label: Some("Compute Pipeline"),
layout: Some(&pipeline_layout),
module: &module,
entry_point: Some("compMain"),
compilation_options: Default::default(),
cache: None,
});
let mut encoder = ctx
.device
.create_command_encoder(&CommandEncoderDescriptor { label: None });
{
let mut render_pass = encoder.begin_compute_pass(&ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
render_pass.set_pipeline(&pipeline);
render_pass.set_bind_group(0, &bind_group, &[]);
render_pass.dispatch_workgroups(1, 1, 1);
}
let readback_buffers = ReadbackBuffers::new(&ctx.device, &output_texture);
readback_buffers.copy_from(&ctx.device, &mut encoder, &output_texture);
ctx.queue.submit(Some(encoder.finish()));
readback_buffers.assert_buffer_contents(&ctx, &image).await;
}

102
tests/tests/partially_bounded_arrays/mod.rs
View File

@@ -1,102 +0,0 @@
use std::num::NonZeroU32;
use wgpu_test::{gpu_test, image::ReadbackBuffers, GpuTestConfiguration, TestParameters};
#[gpu_test]
static PARTIALLY_BOUNDED_ARRAY: GpuTestConfiguration = GpuTestConfiguration::new()
.parameters(
TestParameters::default()
.features(
wgpu::Features::TEXTURE_BINDING_ARRAY
| wgpu::Features::STORAGE_RESOURCE_BINDING_ARRAY
| wgpu::Features::PARTIALLY_BOUND_BINDING_ARRAY
| wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES,
)
.limits(wgpu::Limits::downlevel_defaults()),
)
.run_async(|ctx| async move {
let device = &ctx.device;
let texture_extent = wgpu::Extent3d {
width: 1,
height: 1,
depth_or_array_layers: 1,
};
let storage_texture = device.create_texture(&wgpu::TextureDescriptor {
label: None,
size: texture_extent,
mip_level_count: 1,
sample_count: 1,
dimension: wgpu::TextureDimension::D2,
format: wgpu::TextureFormat::Rgba32Float,
usage: wgpu::TextureUsages::TEXTURE_BINDING
| wgpu::TextureUsages::COPY_DST
| wgpu::TextureUsages::STORAGE_BINDING
| wgpu::TextureUsages::COPY_SRC,
view_formats: &[],
});
let texture_view = storage_texture.create_view(&wgpu::TextureViewDescriptor::default());
let bind_group_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
label: Some("bind group layout"),
entries: &[wgpu::BindGroupLayoutEntry {
binding: 0,
visibility: wgpu::ShaderStages::COMPUTE,
ty: wgpu::BindingType::StorageTexture {
access: wgpu::StorageTextureAccess::WriteOnly,
format: wgpu::TextureFormat::Rgba32Float,
view_dimension: wgpu::TextureViewDimension::D2,
},
count: NonZeroU32::new(4),
}],
});
let cs_module = device.create_shader_module(wgpu::include_wgsl!("shader.wgsl"));
let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
label: Some("main"),
bind_group_layouts: &[&bind_group_layout],
push_constant_ranges: &[],
});
let compute_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
label: None,
layout: Some(&pipeline_layout),
module: &cs_module,
entry_point: Some("main"),
compilation_options: Default::default(),
cache: None,
});
let bind_group = device.create_bind_group(&wgpu::BindGroupDescriptor {
entries: &[wgpu::BindGroupEntry {
binding: 0,
resource: wgpu::BindingResource::TextureViewArray(&[&texture_view]),
}],
layout: &bind_group_layout,
label: Some("bind group"),
});
let mut encoder =
device.create_command_encoder(&wgpu::CommandEncoderDescriptor { label: None });
{
let mut cpass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
label: None,
timestamp_writes: None,
});
cpass.set_pipeline(&compute_pipeline);
cpass.set_bind_group(0, &bind_group, &[]);
cpass.dispatch_workgroups(1, 1, 1);
}
let readback_buffers = ReadbackBuffers::new(&ctx.device, &storage_texture);
readback_buffers.copy_from(&ctx.device, &mut encoder, &storage_texture);
ctx.queue.submit(Some(encoder.finish()));
readback_buffers
.assert_buffer_contents(&ctx, bytemuck::bytes_of(&[4.0f32, 3.0, 2.0, 1.0]))
.await;
});

11
tests/tests/partially_bounded_arrays/shader.wgsl
View File

@@ -1,11 +0,0 @@
@group(0)
@binding(0)
var texture_array_storage: binding_array<texture_storage_2d<rgba32float,write>,1>;
@compute
@workgroup_size(1)
fn main(@builtin(global_invocation_id) global_id: vec3<u32>) {
textureStore(texture_array_storage[0],vec2<i32>(0,0), vec4<f32>(4.0,3.0,2.0,1.0));
}

2
tests/tests/root.rs
View File

@@ -13,6 +13,7 @@ mod regression {
mod bgra8unorm_storage;
mod bind_group_layout_dedup;
mod bind_groups;
mod binding_array;
mod buffer;
mod buffer_copy;
mod buffer_usages;
@@ -30,7 +31,6 @@ mod mem_leaks;
mod nv12_texture;
mod occlusion_query;
mod oob_indexing;
mod partially_bounded_arrays;
mod pipeline;
mod pipeline_cache;
mod poll;