From 286082acefb304d69dd47d13cb0ea07e69bfc0dd Mon Sep 17 00:00:00 2001 From: Connor Fitzgerald Date: Mon, 16 Dec 2024 04:08:05 -0500 Subject: [PATCH] Bindless Tests (#6732) * 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 --- tests/Cargo.toml | 1 + tests/tests/binding_array/buffers.rs | 265 ++++++++++++++++++ tests/tests/binding_array/mod.rs | 4 + tests/tests/binding_array/sampled_textures.rs | 234 ++++++++++++++++ tests/tests/binding_array/samplers.rs | 251 +++++++++++++++++ tests/tests/binding_array/storage_textures.rs | 203 ++++++++++++++ tests/tests/partially_bounded_arrays/mod.rs | 102 ------- .../partially_bounded_arrays/shader.wgsl | 11 - tests/tests/root.rs | 2 +- 9 files changed, 959 insertions(+), 114 deletions(-) create mode 100644 tests/tests/binding_array/buffers.rs create mode 100644 tests/tests/binding_array/mod.rs create mode 100644 tests/tests/binding_array/sampled_textures.rs create mode 100644 tests/tests/binding_array/samplers.rs create mode 100644 tests/tests/binding_array/storage_textures.rs delete mode 100644 tests/tests/partially_bounded_arrays/mod.rs delete mode 100644 tests/tests/partially_bounded_arrays/shader.wgsl diff --git a/tests/Cargo.toml b/tests/Cargo.toml index db91fb8665..127226b0bd 100644 --- a/tests/Cargo.toml +++ b/tests/Cargo.toml @@ -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 diff --git a/tests/tests/binding_array/buffers.rs b/tests/tests/binding_array/buffers.rs new file mode 100644 index 0000000000..1ef9818302 --- /dev/null +++ b/tests/tests/binding_array/buffers.rs @@ -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; + + @group(0) @binding(1) + var output_buffer: array; + + @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); +} diff --git a/tests/tests/binding_array/mod.rs b/tests/tests/binding_array/mod.rs new file mode 100644 index 0000000000..4b8972fcdb --- /dev/null +++ b/tests/tests/binding_array/mod.rs @@ -0,0 +1,4 @@ +mod buffers; +mod sampled_textures; +mod samplers; +mod storage_textures; diff --git a/tests/tests/binding_array/sampled_textures.rs b/tests/tests/binding_array/sampled_textures.rs new file mode 100644 index 0000000000..e95476e876 --- /dev/null +++ b/tests/tests/binding_array/sampled_textures.rs @@ -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>; + + @vertex + fn vertMain(@builtin(vertex_index) id: u32) -> @builtin(position) vec4f { + var positions = array( + vec2f(-1.0, -1.0), + vec2f(3.0, -1.0), + vec2f(-1.0, 3.0) + ); + + return vec4(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; +} diff --git a/tests/tests/binding_array/samplers.rs b/tests/tests/binding_array/samplers.rs new file mode 100644 index 0000000000..d4ff2a24b5 --- /dev/null +++ b/tests/tests/binding_array/samplers.rs @@ -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; + @group(0) @binding(1) + var texture: texture_2d; + @group(0) @binding(2) + var output_values: array; + + @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[..]); +} diff --git a/tests/tests/binding_array/storage_textures.rs b/tests/tests/binding_array/storage_textures.rs new file mode 100644 index 0000000000..ed8e6b4edb --- /dev/null +++ b/tests/tests/binding_array/storage_textures.rs @@ -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 >; + + @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; +} diff --git a/tests/tests/partially_bounded_arrays/mod.rs b/tests/tests/partially_bounded_arrays/mod.rs deleted file mode 100644 index 669c13c511..0000000000 --- a/tests/tests/partially_bounded_arrays/mod.rs +++ /dev/null @@ -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; - }); diff --git a/tests/tests/partially_bounded_arrays/shader.wgsl b/tests/tests/partially_bounded_arrays/shader.wgsl deleted file mode 100644 index 7d475800fa..0000000000 --- a/tests/tests/partially_bounded_arrays/shader.wgsl +++ /dev/null @@ -1,11 +0,0 @@ -@group(0) -@binding(0) -var texture_array_storage: binding_array,1>; - -@compute -@workgroup_size(1) -fn main(@builtin(global_invocation_id) global_id: vec3) { - - textureStore(texture_array_storage[0],vec2(0,0), vec4(4.0,3.0,2.0,1.0)); - -} diff --git a/tests/tests/root.rs b/tests/tests/root.rs index dac56a9db0..9df2b12248 100644 --- a/tests/tests/root.rs +++ b/tests/tests/root.rs @@ -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;