From 0d6a06c6d99ef485c6826157662e3f42fa08ba2d Mon Sep 17 00:00:00 2001 From: Erich Gubler Date: Wed, 21 Feb 2024 15:20:50 -0500 Subject: [PATCH] Fix missing validation for `Device::clear_buffer` where `offset + size > buffer.size` was not checked when `size` was omitted. (#5282) MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit style: fix fmt. of `assert!(…)` in `clear_texture_via_buffer_copies` refactor: `command_encoder_clear_buffer`: s/end/end_offset fix: always check buffer clear `offset` for OOB Fuzz testing in Firefox encountered crashes for calls of `Global::command_encoder_clear_buffer` where: * `offset` is greater than `buffer.size`, but… * `size` is `None`. Oops! We should _always_ check this (i.e., even when `size` is `None`), because we have no guarantee that `offset` and the fallback value of `size` is in bounds. 😅 So, we change validation here to unconditionally compute `size` and run checks we previously gated behind `if let Some(size) = size { … }`. For convenience, the spec. link for this method: fix: `command_encoder_clear_buffer`: err. on `offset + size > u64::MAX` Rust would have made this operation either an overflow in release mode, or a panic in debug mode. Neither seem appropriate for this context, where I suspect an error should be returned instead. Web browsers, for instance, shouldn't crash simply because of an issue of this nature. Users may, quite reasonably, have bad arguments to this in early stages of development! --- CHANGELOG.md | 1 + tests/tests/buffer.rs | 221 +++++++++++++++++++++++++++++++++ wgpu-core/src/command/clear.rs | 45 ++++--- 3 files changed, 248 insertions(+), 19 deletions(-) diff --git a/CHANGELOG.md b/CHANGELOG.md index b7012833c2..7d08a0ea15 100644 --- a/CHANGELOG.md +++ b/CHANGELOG.md @@ -85,6 +85,7 @@ Bottom level categories: - 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) diff --git a/tests/tests/buffer.rs b/tests/tests/buffer.rs index c3b1dbea58..a5fcf3e595 100644 --- a/tests/tests/buffer.rs +++ b/tests/tests/buffer.rs @@ -164,3 +164,224 @@ static MAP_OFFSET: GpuTestConfiguration = GpuTestConfiguration::new().run_async( assert_eq!(*byte, 0); } }); + +/// 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 build a pipeline from a shader module with a +/// 32-byte variable and a bindgroup layout with a min_binding_size of +/// 16 for that variable's group/index. Pipeline creation should fail. +#[gpu_test] +static MINIMUM_BUFFER_BINDING_SIZE_LAYOUT: GpuTestConfiguration = GpuTestConfiguration::new() + .parameters(TestParameters::default().test_features_limits()) + .run_sync(|ctx| { + // 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 a: array; + @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 a: array; + @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, + size: 16, // too small for 32-byte var `a` in shader module + usage: wgpu::BufferUsages::STORAGE, + mapped_at_creation: false, + }); + + let bind_group = ctx.device.create_bind_group(&wgpu::BindGroupDescriptor { + label: None, + layout: &bind_group_layout, + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: buffer.as_entire_binding(), + }], + }); + + 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] +static CLEAR_OFFSET_OUTSIDE_RESOURCE_BOUNDS: GpuTestConfiguration = GpuTestConfiguration::new() + .parameters(TestParameters::default()) + .run_sync(|ctx| { + let size = 16; + + let buffer = ctx.device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size, + usage: wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + + 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" + )); + }); + +#[gpu_test] +static CLEAR_OFFSET_PLUS_SIZE_OUTSIDE_U64_BOUNDS: GpuTestConfiguration = + GpuTestConfiguration::new() + .parameters(TestParameters::default()) + .run_sync(|ctx| { + let buffer = ctx.device.create_buffer(&wgpu::BufferDescriptor { + label: None, + size: 16, // unimportant for this test + usage: wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + + let max_valid_offset = u64::MAX - (u64::MAX % wgpu::COPY_BUFFER_ALIGNMENT); + let smallest_aligned_invalid_size = wgpu::COPY_BUFFER_ALIGNMENT; + + ctx.device.push_error_scope(wgpu::ErrorFilter::Validation); + ctx.device + .create_command_encoder(&Default::default()) + .clear_buffer( + &buffer, + max_valid_offset, + Some(smallest_aligned_invalid_size), + ); + let err_msg = pollster::block_on(ctx.device.pop_error_scope()) + .unwrap() + .to_string(); + assert!(err_msg.contains(concat!( + "Clear starts at offset 18446744073709551612 with size of 4, ", + "but these added together exceed `u64::MAX`" + ))); + }); diff --git a/wgpu-core/src/command/clear.rs b/wgpu-core/src/command/clear.rs index 1a4b4cdeb1..f4c50dc853 100644 --- a/wgpu-core/src/command/clear.rs +++ b/wgpu-core/src/command/clear.rs @@ -40,6 +40,11 @@ pub enum ClearError { UnalignedFillSize(BufferAddress), #[error("Buffer offset {0:?} is not a multiple of `COPY_BUFFER_ALIGNMENT`")] UnalignedBufferOffset(BufferAddress), + #[error("Clear starts at offset {start_offset} with size of {requested_size}, but these added together exceed `u64::MAX`")] + OffsetPlusSizeExceeds64BitBounds { + start_offset: BufferAddress, + requested_size: BufferAddress, + }, #[error("Clear of {start_offset}..{end_offset} would end up overrunning the bounds of the buffer of size {buffer_size}")] BufferOverrun { start_offset: BufferAddress, @@ -118,25 +123,27 @@ impl Global { if offset % wgt::COPY_BUFFER_ALIGNMENT != 0 { return Err(ClearError::UnalignedBufferOffset(offset)); } - if let Some(size) = size { - if size % wgt::COPY_BUFFER_ALIGNMENT != 0 { - return Err(ClearError::UnalignedFillSize(size)); - } - let destination_end_offset = offset + size; - if destination_end_offset > dst_buffer.size { - return Err(ClearError::BufferOverrun { + + let size = size.unwrap_or(dst_buffer.size.saturating_sub(offset)); + if size % wgt::COPY_BUFFER_ALIGNMENT != 0 { + return Err(ClearError::UnalignedFillSize(size)); + } + let end_offset = + offset + .checked_add(size) + .ok_or(ClearError::OffsetPlusSizeExceeds64BitBounds { start_offset: offset, - end_offset: destination_end_offset, - buffer_size: dst_buffer.size, - }); - } + requested_size: size, + })?; + if end_offset > dst_buffer.size { + return Err(ClearError::BufferOverrun { + start_offset: offset, + end_offset, + buffer_size: dst_buffer.size, + }); } - let end = match size { - Some(size) => offset + size, - None => dst_buffer.size, - }; - if offset == end { + if offset == end_offset { log::trace!("Ignoring fill_buffer of size 0"); return Ok(()); } @@ -145,7 +152,7 @@ impl Global { cmd_buf_data.buffer_memory_init_actions.extend( dst_buffer.initialization_status.read().create_action( &dst_buffer, - offset..end, + offset..end_offset, MemoryInitKind::ImplicitlyInitialized, ), ); @@ -155,7 +162,7 @@ impl Global { let cmd_buf_raw = cmd_buf_data.encoder.open()?; unsafe { cmd_buf_raw.transition_buffers(dst_barrier.into_iter()); - cmd_buf_raw.clear_buffer(dst_raw, offset..end); + cmd_buf_raw.clear_buffer(dst_raw, offset..end_offset); } Ok(()) } @@ -367,7 +374,7 @@ fn clear_texture_via_buffer_copies( assert!( max_rows_per_copy > 0, "Zero buffer size is too small to fill a single row \ - of a texture with format {:?} and desc {:?}", + of a texture with format {:?} and desc {:?}", texture_desc.format, texture_desc.size );