[wgpu-core] Compute minimum binding size correctly for arrays. (gfx-r…

…s#5222) * [wgpu-core] Add tests for minimum binding size validation. * [wgpu-core] Compute minimum binding size correctly for arrays. In early versions of WGSL, `storage` or `uniform` global variables had to be either structs or runtime-sized arrays. This rule was relaxed, and now globals can have any type; Naga automatically wraps such variables in structs when required by the backend shading language. Under the old rules, whenever wgpu-core saw a `storage` or `uniform` global variable with an array type, it could assume it was a runtime-sized array, and take the stride as the minimum binding size. Under the new rules, wgpu-core must consider fixed-sized and runtime-sized arrays separately.
suspistew · Mar 23, 2024 · ad34bc6 · ad34bc6
1 parent 9f505e7
commit ad34bc6
Show file tree

Hide file tree

Showing 3 changed files with 145 additions and 35 deletions.
diff --git a/CHANGELOG.md b/CHANGELOG.md
@@ -106,25 +106,8 @@ By @cwfitzgerald in [#5325](https://github.com/gfx-rs/wgpu/pull/5325).
 
 #### General
 - Device lost callbacks are invoked when replaced and when global is dropped. By @bradwerth in [#5168](https://github.com/gfx-rs/wgpu/pull/5168)
-- Fix performance regression when allocating a large amount of resources of the same type. By @nical in [#5229](https://github.com/gfx-rs/wgpu/pull/5229)
-- Fix docs.rs wasm32 builds. By @cwfitzgerald in [#5310](https://github.com/gfx-rs/wgpu/pull/5310)
-- Improve error message when binding count limit hit. By @hackaugusto in [#5298](https://github.com/gfx-rs/wgpu/pull/5298)
-- Remove an unnecessary `clone` during GLSL shader injestion. By @a1phyr in [#5118](https://github.com/gfx-rs/wgpu/pull/5118).
-- Fix missing validation for `Device::clear_buffer` where `offset + size > buffer.size` was not checked when `size` was omitted. By @ErichDonGubler in [#5282](https://github.com/gfx-rs/wgpu/pull/5282).
-
-#### DX12
-- Fix `panic!` when dropping `Instance` without `InstanceFlags::VALIDATION`. By @hakolao in [#5134](https://github.com/gfx-rs/wgpu/pull/5134)
-
-#### OpenGL
-- Fix internal format for the `Etc2Rgba8Unorm` format. By @andristarr in [#5178](https://github.com/gfx-rs/wgpu/pull/5178)
-- Try to load `libX11.so.6` in addition to `libX11.so` on linux. [#5307](https://github.com/gfx-rs/wgpu/pull/5307)
-- Make use of `GL_EXT_texture_shadow_lod` to support sampling a cube depth texture with an explicit LOD. By @cmrschwarz in #[5171](https://github.com/gfx-rs/wgpu/pull/5171).
-
-#### `glsl-in`
-
-- Fix code generation from nested loops. By @cwfitzgerald and @teoxoy in [#5311](https://github.com/gfx-rs/wgpu/pull/5311)
-
-## v0.19.1 (2024-01-21)
+- Fix panic when creating a surface while no backend is available. By @wumpf [#5166](https://github.com/gfx-rs/wgpu/pull/5166)
+- Correctly compute minimum buffer size for array-typed `storage` and `uniform` vars. By @jimblandy [#5222](https://github.com/gfx-rs/wgpu/pull/5222)
 
 #### WGL
 

diff --git a/tests/tests/buffer.rs b/tests/tests/buffer.rs
@@ -166,10 +166,126 @@ static MAP_OFFSET: GpuTestConfiguration = GpuTestConfiguration::new().run_async(
 });
 
 #[gpu_test]
-static CLEAR_OFFSET_OUTSIDE_RESOURCE_BOUNDS: GpuTestConfiguration = GpuTestConfiguration::new()
-    .parameters(TestParameters::default())
+static MINIMUM_BUFFER_BINDING_SIZE_LAYOUT: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default().test_features_limits())
     .run_sync(|ctx| {
-        let size = 16;
+        // Create a shader module that statically uses a storage buffer.
+        let shader_module = ctx
+            .device
+            .create_shader_module(wgpu::ShaderModuleDescriptor {
+                label: None,
+                source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(
+                    r#"
+                        @group(0) @binding(0)
+                        var<storage, read_write> a: array<u32, 8>;
+                        @compute @workgroup_size(1)
+                        fn main() {
+                            a[0] = a[1];
+                        }
+            "#,
+                )),
+            });
+
+        let bind_group_layout =
+            ctx.device
+                .create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                    label: None,
+                    entries: &[wgpu::BindGroupLayoutEntry {
+                        binding: 0,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: std::num::NonZeroU64::new(16),
+                        },
+                        count: None,
+                    }],
+                });
+
+        let pipeline_layout = ctx
+            .device
+            .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: None,
+                bind_group_layouts: &[&bind_group_layout],
+                push_constant_ranges: &[],
+            });
+
+        wgpu_test::fail(&ctx.device, || {
+            ctx.device
+                .create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                    label: None,
+                    layout: Some(&pipeline_layout),
+                    module: &shader_module,
+                    entry_point: "main",
+                });
+        });
+    });
+
+/// The WebGPU algorithm [validating shader binding][vsb] requires
+/// implementations to check that buffer bindings are large enough to
+/// hold the WGSL `storage` or `uniform` variables they're bound to.
+///
+/// This test tries to dispatch a compute shader that uses a 32-byte
+/// variable with a bindgroup layout with a min_binding_size of zero
+/// (meaning, "validate at dispatch recording time") and a 16-byte
+/// binding. Command recording should fail.
+#[gpu_test]
+static MINIMUM_BUFFER_BINDING_SIZE_DISPATCH: GpuTestConfiguration = GpuTestConfiguration::new()
+    .parameters(TestParameters::default().test_features_limits())
+    .run_sync(|ctx| {
+        // This test tries to use a bindgroup layout with a
+        // min_binding_size of 16 to an index whose WGSL type requires 32
+        // bytes. Pipeline creation should fail.
+
+        // Create a shader module that statically uses a storage buffer.
+        let shader_module = ctx
+            .device
+            .create_shader_module(wgpu::ShaderModuleDescriptor {
+                label: None,
+                source: wgpu::ShaderSource::Wgsl(std::borrow::Cow::Borrowed(
+                    r#"
+                        @group(0) @binding(0)
+                        var<storage, read_write> a: array<u32, 8>;
+                        @compute @workgroup_size(1)
+                        fn main() {
+                            a[0] = a[1];
+                        }
+            "#,
+                )),
+            });
+
+        let bind_group_layout =
+            ctx.device
+                .create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                    label: None,
+                    entries: &[wgpu::BindGroupLayoutEntry {
+                        binding: 0,
+                        visibility: wgpu::ShaderStages::COMPUTE,
+                        ty: wgpu::BindingType::Buffer {
+                            ty: wgpu::BufferBindingType::Storage { read_only: false },
+                            has_dynamic_offset: false,
+                            min_binding_size: None,
+                        },
+                        count: None,
+                    }],
+                });
+
+        let pipeline_layout = ctx
+            .device
+            .create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                label: None,
+                bind_group_layouts: &[&bind_group_layout],
+                push_constant_ranges: &[],
+            });
+
+        let pipeline = ctx
+            .device
+            .create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: None,
+                layout: Some(&pipeline_layout),
+                module: &shader_module,
+                entry_point: "main",
+            });
 
         let buffer = ctx.device.create_buffer(&wgpu::BufferDescriptor {
             label: None,
@@ -180,16 +296,21 @@ static CLEAR_OFFSET_OUTSIDE_RESOURCE_BOUNDS: GpuTestConfiguration = GpuTestConfi
 
         let out_of_bounds = size.checked_add(wgpu::COPY_BUFFER_ALIGNMENT).unwrap();
 
-        ctx.device.push_error_scope(wgpu::ErrorFilter::Validation);
-        ctx.device
-            .create_command_encoder(&Default::default())
-            .clear_buffer(&buffer, out_of_bounds, None);
-        let err_msg = pollster::block_on(ctx.device.pop_error_scope())
-            .unwrap()
-            .to_string();
-        assert!(err_msg.contains(
-            "Clear of 20..20 would end up overrunning the bounds of the buffer of size 16"
-        ));
+        wgpu_test::fail(&ctx.device, || {
+            let mut encoder = ctx.device.create_command_encoder(&Default::default());
+
+            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: None,
+                timestamp_writes: None,
+            });
+
+            pass.set_bind_group(0, &bind_group, &[]);
+            pass.set_pipeline(&pipeline);
+            pass.dispatch_workgroups(1, 1, 1);
+
+            drop(pass);
+            let _ = encoder.finish();
+        });
     });
 
 #[gpu_test]

diff --git a/wgpu-core/src/validation.rs b/wgpu-core/src/validation.rs
@@ -892,9 +892,15 @@ impl Interface {
                     class,
                 },
                 naga::TypeInner::Sampler { comparison } => ResourceType::Sampler { comparison },
-                naga::TypeInner::Array { stride, .. } => ResourceType::Buffer {
-                    size: wgt::BufferSize::new(stride as u64).unwrap(),
-                },
+                naga::TypeInner::Array { stride, size, .. } => {
+                    let size = match size {
+                        naga::ArraySize::Constant(size) => size.get() * stride,
+                        naga::ArraySize::Dynamic => stride,
+                    };
+                    ResourceType::Buffer {
+                        size: wgt::BufferSize::new(size as u64).unwrap(),
+                    }
+                }
                 ref other => ResourceType::Buffer {
                     size: wgt::BufferSize::new(other.size(module.to_ctx()) as u64).unwrap(),
                 },