diff --git a/CHANGELOG.md b/CHANGELOG.md index b7012833c26..7d08a0ea15c 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 c3b1dbea58f..a5fcf3e5953 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 1a4b4cdeb11..f4c50dc8532 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 );