Ensure the limit for storage binding size honors alignment (#8719)

Author: andyleiserson

cts_runner/test.lst

@@ -34,6 +34,7 @@ webgpu:api,validation,buffer,create:*
 webgpu:api,validation,buffer,destroy:*
 fails-if(dx12) webgpu:api,validation,capability_checks,limits,maxBindGroups:setBindGroup,*
 webgpu:api,validation,createBindGroup:buffer,effective_buffer_binding_size:*
+webgpu:api,validation,createBindGroup:buffer,resource_binding_size:*
 // Fails because we coerce a size of 0 in `GPUDevice.createBindGroup(…)` to `buffer.size - offset`.
 // FAIL webgpu:api,validation,createBindGroup:buffer_offset_and_size_for_bind_groups_match:*
 webgpu:api,validation,encoding,beginComputePass:*

wgpu-core/src/binding_model.rs

@@ -168,7 +168,7 @@ pub enum CreateBindGroupError {
     #[error("Binding declared as a single item, but bind group is using it as an array")]
     SingleBindingExpected,
     #[error("Effective buffer binding size {size} for storage buffers is expected to align to {alignment}, but size is {size}")]
-    UnalignedEffectiveBufferBindingSizeForStorage { alignment: u8, size: u64 },
+    UnalignedEffectiveBufferBindingSizeForStorage { alignment: u32, size: u64 },
     #[error("Buffer offset {0} does not respect device's requested `{1}` limit {2}")]
     UnalignedBufferOffset(wgt::BufferAddress, &'static str, u32),
     #[error(

wgpu-core/src/device/resource.rs

@@ -2796,16 +2796,13 @@ impl Device {
 
         let (bb, bind_size) = buffer.binding(bb.offset, bb.size, snatch_guard)?;
 
-        if matches!(binding_ty, wgt::BufferBindingType::Storage { .. }) {
-            let storage_buf_size_alignment = 4;
-
-            let aligned = bind_size % u64::from(storage_buf_size_alignment) == 0;
-            if !aligned {
-                return Err(Error::UnalignedEffectiveBufferBindingSizeForStorage {
-                    alignment: storage_buf_size_alignment,
-                    size: bind_size,
-                });
-            }
+        if matches!(binding_ty, wgt::BufferBindingType::Storage { .. })
+            && bind_size % u64::from(wgt::STORAGE_BINDING_SIZE_ALIGNMENT) != 0
+        {
+            return Err(Error::UnalignedEffectiveBufferBindingSizeForStorage {
+                alignment: wgt::STORAGE_BINDING_SIZE_ALIGNMENT,
+                size: bind_size,
+            });
         }
 
         let bind_end = bb.offset + bind_size;

wgpu-hal/build.rs

@@ -23,6 +23,7 @@ fn main() {
         ) },
         metal: { all(target_vendor = "apple", feature = "metal") },
         vulkan: { all(not(target_arch = "wasm32"), feature = "vulkan") },
+        any_backend: { any(dx12, metal, vulkan, gles) },
         // ⚠️ Keep in sync with target.cfg() definition in Cargo.toml and cfg_alias in `wgpu` crate ⚠️
         static_dxc: { all(target_os = "windows", feature = "static-dxc", not(target_arch = "aarch64"), target_env = "msvc") },
         supports_64bit_atomics: { target_has_atomic = "64" },

wgpu-hal/src/auxil/mod.rs

@@ -126,3 +126,21 @@ impl crate::TextureCopy {
         self.size = self.size.min(&max_src_size).min(&max_dst_size);
     }
 }
+
+/// Clamp the limits in `limits` to honor any HAL-imposed maximums.
+///
+/// Limits that do not have a HAL-defined maximum are left unchanged.
+#[cfg_attr(not(any_backend), allow(dead_code))]
+pub(crate) fn apply_hal_limits(mut limits: wgt::Limits) -> wgt::Limits {
+    // The Metal backend wants to have its own consistent view of the limits, so
+    // it may duplicate some of these limits.
+    limits.max_bind_groups = limits.max_bind_groups.min(crate::MAX_BIND_GROUPS as u32);
+    limits.max_storage_buffer_binding_size &= !(wgt::STORAGE_BINDING_SIZE_ALIGNMENT - 1);
+    limits.max_vertex_buffers = limits
+        .max_vertex_buffers
+        .min(crate::MAX_VERTEX_BUFFERS as u32);
+    limits.max_color_attachments = limits
+        .max_color_attachments
+        .min(crate::MAX_COLOR_ATTACHMENTS as u32);
+    limits
+}

wgpu-hal/src/dx12/adapter.rs

@@ -695,7 +695,7 @@ impl super::Adapter {
             info,
             features,
             capabilities: crate::Capabilities {
-                limits: wgt::Limits {
+                limits: auxil::apply_hal_limits(wgt::Limits {
                     max_texture_dimension_1d: Direct3D12::D3D12_REQ_TEXTURE1D_U_DIMENSION,
                     max_texture_dimension_2d: Direct3D12::D3D12_REQ_TEXTURE2D_U_OR_V_DIMENSION
                         .min(Direct3D12::D3D12_REQ_TEXTURECUBE_DIMENSION),
@@ -724,8 +724,7 @@ impl super::Adapter {
                     max_uniform_buffer_binding_size:
                         Direct3D12::D3D12_REQ_CONSTANT_BUFFER_ELEMENT_COUNT * 16,
                     max_storage_buffer_binding_size: auxil::MAX_I32_BINDING_SIZE,
-                    max_vertex_buffers: Direct3D12::D3D12_VS_INPUT_REGISTER_COUNT
-                        .min(crate::MAX_VERTEX_BUFFERS as u32),
+                    max_vertex_buffers: Direct3D12::D3D12_VS_INPUT_REGISTER_COUNT,
                     max_vertex_attributes: Direct3D12::D3D12_IA_VERTEX_INPUT_RESOURCE_SLOT_COUNT,
                     max_vertex_buffer_array_stride: Direct3D12::D3D12_SO_BUFFER_MAX_STRIDE_IN_BYTES,
                     // The immediates are part of the root signature which
@@ -806,7 +805,7 @@ impl super::Adapter {
                     },
 
                     max_multiview_view_count,
-                },
+                }),
                 alignments: crate::Alignments {
                     buffer_copy_offset: wgt::BufferSize::new(
                         Direct3D12::D3D12_TEXTURE_DATA_PLACEMENT_ALIGNMENT as u64,

wgpu-hal/src/gles/adapter.rs

@@ -680,15 +680,14 @@ impl super::Adapter {
 
         let max_color_attachments = unsafe {
             gl.get_parameter_i32(glow::MAX_COLOR_ATTACHMENTS)
-                .min(gl.get_parameter_i32(glow::MAX_DRAW_BUFFERS))
-                .min(crate::MAX_COLOR_ATTACHMENTS as i32) as u32
+                .min(gl.get_parameter_i32(glow::MAX_DRAW_BUFFERS)) as u32
         };
 
         // 16 bytes per sample is the maximum size of a color attachment.
         let max_color_attachment_bytes_per_sample =
             max_color_attachments * wgt::TextureFormat::MAX_TARGET_PIXEL_BYTE_COST;
 
-        let limits = wgt::Limits {
+        let limits = crate::auxil::apply_hal_limits(wgt::Limits {
             max_texture_dimension_1d: max_texture_size,
             max_texture_dimension_2d: max_texture_size,
             max_texture_dimension_3d: max_texture_3d_size,
@@ -720,8 +719,7 @@ impl super::Adapter {
                 (unsafe { gl.get_parameter_i32(glow::MAX_VERTEX_ATTRIB_BINDINGS) } as u32)
             } else {
                 16 // should this be different?
-            }
-            .min(crate::MAX_VERTEX_BUFFERS as u32),
+            },
             max_vertex_attributes: (unsafe { gl.get_parameter_i32(glow::MAX_VERTEX_ATTRIBS) }
                 as u32)
                 .min(super::MAX_VERTEX_ATTRIBUTES as u32),
@@ -815,7 +813,7 @@ impl super::Adapter {
             max_acceleration_structures_per_shader_stage: 0,
 
             max_multiview_view_count: 0,
-        };
+        });
 
         let mut workarounds = super::Workarounds::empty();
 

wgpu-hal/src/metal/adapter.rs

@@ -1136,72 +1136,92 @@ impl super::PrivateCapabilities {
         downlevel
             .flags
             .set(wgt::DownlevelFlags::ANISOTROPIC_FILTERING, true);
+
         let base = wgt::Limits::default();
-        crate::Capabilities {
-            limits: wgt::Limits {
-                max_texture_dimension_1d: self.max_texture_size as u32,
-                max_texture_dimension_2d: self.max_texture_size as u32,
-                max_texture_dimension_3d: self.max_texture_3d_size as u32,
-                max_texture_array_layers: self.max_texture_layers as u32,
-                max_bind_groups: 8,
-                max_bindings_per_bind_group: 65535,
-                max_dynamic_uniform_buffers_per_pipeline_layout: base
-                    .max_dynamic_uniform_buffers_per_pipeline_layout,
-                max_dynamic_storage_buffers_per_pipeline_layout: base
-                    .max_dynamic_storage_buffers_per_pipeline_layout,
-                max_sampled_textures_per_shader_stage: self.max_textures_per_stage,
-                max_samplers_per_shader_stage: self.max_samplers_per_stage,
-                max_storage_buffers_per_shader_stage: self.max_buffers_per_stage,
-                max_storage_textures_per_shader_stage: self.max_textures_per_stage,
-                max_uniform_buffers_per_shader_stage: self.max_buffers_per_stage,
-                max_binding_array_elements_per_shader_stage: self.max_binding_array_elements,
-                max_binding_array_sampler_elements_per_shader_stage: self
-                    .max_sampler_binding_array_elements,
-                max_uniform_buffer_binding_size: self.max_buffer_size.min(!0u32 as u64) as u32,
-                max_storage_buffer_binding_size: self.max_buffer_size.min(!0u32 as u64) as u32,
-                max_vertex_buffers: self.max_vertex_buffers,
-                max_vertex_attributes: 31,
-                max_vertex_buffer_array_stride: base.max_vertex_buffer_array_stride,
-                max_immediate_size: 0x1000,
-                min_uniform_buffer_offset_alignment: self.buffer_alignment as u32,
-                min_storage_buffer_offset_alignment: self.buffer_alignment as u32,
-                max_inter_stage_shader_components: self.max_varying_components,
-                max_color_attachments: (self.max_color_render_targets as u32)
-                    .min(crate::MAX_COLOR_ATTACHMENTS as u32),
-                max_color_attachment_bytes_per_sample: self.max_color_attachment_bytes_per_sample
-                    as u32,
-                max_compute_workgroup_storage_size: self.max_total_threadgroup_memory,
-                max_compute_invocations_per_workgroup: self.max_threads_per_group,
-                max_compute_workgroup_size_x: self.max_threads_per_group,
-                max_compute_workgroup_size_y: self.max_threads_per_group,
-                max_compute_workgroup_size_z: self.max_threads_per_group,
-                max_compute_workgroups_per_dimension: 0xFFFF,
-                max_buffer_size: self.max_buffer_size,
-                max_non_sampler_bindings: u32::MAX,
-
-                // See https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf, Maximum threadgroups per mesh shader grid
-                max_task_workgroup_total_count: 1024,
-                max_task_workgroups_per_dimension: 1024,
-                max_mesh_multiview_view_count: 0,
-                max_mesh_output_layers: self.max_texture_layers as u32,
-
-                max_blas_primitive_count: 0, // When added: 2^28 from https://developer.apple.com/documentation/metal/mtlaccelerationstructureusage/extendedlimits
-                max_blas_geometry_count: 0,  // When added: 2^24
-                max_tlas_instance_count: 0,  // When added: 2^24
-                // Unsure what this will be when added: acceleration structures count as a buffer so
-                // it may be worth using argument buffers for this all acceleration structures, then
-                // there will be no limit.
-                // From 2.17.7 in https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
-                // > [Acceleration structures] are opaque objects that can be bound directly using
-                // buffer binding points or via argument buffers
-                max_acceleration_structures_per_shader_stage: 0,
-
-                max_multiview_view_count: if self.supported_vertex_amplification_factor > 1 {
-                    self.supported_vertex_amplification_factor
-                } else {
-                    0
-                },
+
+        // Be careful adjusting limits here. The `AdapterShared` stores the
+        // original `PrivateCapabilities`, so code could accidentally use
+        // the wrong value.
+
+        let limits = wgt::Limits {
+            max_texture_dimension_1d: self.max_texture_size as u32,
+            max_texture_dimension_2d: self.max_texture_size as u32,
+            max_texture_dimension_3d: self.max_texture_3d_size as u32,
+            max_texture_array_layers: self.max_texture_layers as u32,
+            max_bind_groups: 8,
+            max_bindings_per_bind_group: 65535,
+            max_dynamic_uniform_buffers_per_pipeline_layout: base
+                .max_dynamic_uniform_buffers_per_pipeline_layout,
+            max_dynamic_storage_buffers_per_pipeline_layout: base
+                .max_dynamic_storage_buffers_per_pipeline_layout,
+            max_sampled_textures_per_shader_stage: self.max_textures_per_stage,
+            max_samplers_per_shader_stage: self.max_samplers_per_stage,
+            max_storage_buffers_per_shader_stage: self.max_buffers_per_stage,
+            max_storage_textures_per_shader_stage: self.max_textures_per_stage,
+            max_uniform_buffers_per_shader_stage: self.max_buffers_per_stage,
+            max_binding_array_elements_per_shader_stage: self.max_binding_array_elements,
+            max_binding_array_sampler_elements_per_shader_stage: self
+                .max_sampler_binding_array_elements,
+            // Note: any adjustment here will not be reflected in the stored `PrivateCapabilities`.
+            max_uniform_buffer_binding_size: self.max_buffer_size.min(!0u32 as u64) as u32,
+            max_storage_buffer_binding_size: self.max_buffer_size.min(!0u32 as u64) as u32
+                & !(wgt::STORAGE_BINDING_SIZE_ALIGNMENT - 1),
+            max_vertex_buffers: self.max_vertex_buffers,
+            max_vertex_attributes: 31,
+            max_vertex_buffer_array_stride: base.max_vertex_buffer_array_stride,
+            max_immediate_size: 0x1000,
+            min_uniform_buffer_offset_alignment: self.buffer_alignment as u32,
+            min_storage_buffer_offset_alignment: self.buffer_alignment as u32,
+            max_inter_stage_shader_components: self.max_varying_components,
+            max_color_attachments: self.max_color_render_targets as u32,
+            max_color_attachment_bytes_per_sample: self.max_color_attachment_bytes_per_sample
+                as u32,
+            max_compute_workgroup_storage_size: self.max_total_threadgroup_memory,
+            max_compute_invocations_per_workgroup: self.max_threads_per_group,
+            max_compute_workgroup_size_x: self.max_threads_per_group,
+            max_compute_workgroup_size_y: self.max_threads_per_group,
+            max_compute_workgroup_size_z: self.max_threads_per_group,
+            max_compute_workgroups_per_dimension: 0xFFFF,
+            max_buffer_size: self.max_buffer_size,
+            max_non_sampler_bindings: u32::MAX,
+
+            // See https://developer.apple.com/metal/Metal-Feature-Set-Tables.pdf, Maximum threadgroups per mesh shader grid
+            max_task_workgroup_total_count: 1024,
+            max_task_workgroups_per_dimension: 1024,
+            max_mesh_multiview_view_count: 0,
+            max_mesh_output_layers: self.max_texture_layers as u32,
+
+            max_blas_primitive_count: 0, // When added: 2^28 from https://developer.apple.com/documentation/metal/mtlaccelerationstructureusage/extendedlimits
+            max_blas_geometry_count: 0,  // When added: 2^24
+            max_tlas_instance_count: 0,  // When added: 2^24
+            // Unsure what this will be when added: acceleration structures count as a buffer so
+            // it may be worth using argument buffers for this all acceleration structures, then
+            // there will be no limit.
+            // From 2.17.7 in https://developer.apple.com/metal/Metal-Shading-Language-Specification.pdf
+            // > [Acceleration structures] are opaque objects that can be bound directly using
+            // buffer binding points or via argument buffers
+            max_acceleration_structures_per_shader_stage: 0,
+
+            max_multiview_view_count: if self.supported_vertex_amplification_factor > 1 {
+                self.supported_vertex_amplification_factor
+            } else {
+                0
             },
+        };
+
+        // Since a bunch of the limits are duplicated between `Limits` and
+        // `PrivateCapabilities`, reducing the limits at this point could make
+        // things inconsistent and lead to confusion. Make sure that doesn't
+        // happen.
+        debug_assert!(
+            crate::auxil::apply_hal_limits(limits.clone()) == limits,
+            "Limits were modified by apply_hal_limits\nOriginal:\n{:#?}\nModified:\n{:#?}",
+            limits,
+            crate::auxil::apply_hal_limits(limits.clone())
+        );
+
+        crate::Capabilities {
+            limits,
             alignments: crate::Alignments {
                 buffer_copy_offset: wgt::BufferSize::new(self.buffer_alignment).unwrap(),
                 buffer_copy_pitch: wgt::BufferSize::new(4).unwrap(),

wgpu-hal/src/metal/mod.rs

@@ -202,6 +202,8 @@ bitflags!(
     }
 );
 
+// TODO(https://github.com/gfx-rs/wgpu/issues/8715): Eliminate duplication with
+// `wgt::Limits`. Keeping multiple sets of limits creates a risk of confusion.
 #[allow(dead_code)]
 #[derive(Clone, Debug)]
 struct PrivateCapabilities {
@@ -277,6 +279,11 @@ struct PrivateCapabilities {
     max_binding_array_elements: ResourceIndex,
     max_sampler_binding_array_elements: ResourceIndex,
     buffer_alignment: u64,
+
+    /// Platform-reported maximum buffer size
+    ///
+    /// This value is clamped to `u32::MAX` for `wgt::Limits`, so you probably
+    /// shouldn't be looking at this copy.
     max_buffer_size: u64,
     max_texture_size: u64,
     max_texture_3d_size: u64,

wgpu-hal/src/vulkan/adapter.rs

@@ -1328,14 +1328,12 @@ impl PhysicalDeviceProperties {
             .map(|a| a.max_multiview_view_count.min(32))
             .unwrap_or(0);
 
-        wgt::Limits {
+        crate::auxil::apply_hal_limits(wgt::Limits {
             max_texture_dimension_1d: limits.max_image_dimension1_d,
             max_texture_dimension_2d: limits.max_image_dimension2_d,
             max_texture_dimension_3d: limits.max_image_dimension3_d,
             max_texture_array_layers: limits.max_image_array_layers,
-            max_bind_groups: limits
-                .max_bound_descriptor_sets
-                .min(crate::MAX_BIND_GROUPS as u32),
+            max_bind_groups: limits.max_bound_descriptor_sets,
             max_bindings_per_bind_group: wgt::Limits::default().max_bindings_per_bind_group,
             max_dynamic_uniform_buffers_per_pipeline_layout: limits
                 .max_descriptor_set_uniform_buffers_dynamic,
@@ -1354,9 +1352,7 @@ impl PhysicalDeviceProperties {
             max_storage_buffer_binding_size: limits
                 .max_storage_buffer_range
                 .min(crate::auxil::MAX_I32_BINDING_SIZE),
-            max_vertex_buffers: limits
-                .max_vertex_input_bindings
-                .min(crate::MAX_VERTEX_BUFFERS as u32),
+            max_vertex_buffers: limits.max_vertex_input_bindings,
             max_vertex_attributes: limits.max_vertex_input_attributes,
             max_vertex_buffer_array_stride: limits.max_vertex_input_binding_stride,
             max_immediate_size: limits.max_push_constants_size,
@@ -1365,9 +1361,7 @@ impl PhysicalDeviceProperties {
             max_inter_stage_shader_components: limits
                 .max_vertex_output_components
                 .min(limits.max_fragment_input_components),
-            max_color_attachments: limits
-                .max_color_attachments
-                .min(crate::MAX_COLOR_ATTACHMENTS as u32),
+            max_color_attachments: limits.max_color_attachments,
             max_color_attachment_bytes_per_sample,
             max_compute_workgroup_storage_size: limits.max_compute_shared_memory_size,
             max_compute_invocations_per_workgroup: limits.max_compute_work_group_invocations,
@@ -1389,7 +1383,7 @@ impl PhysicalDeviceProperties {
             max_acceleration_structures_per_shader_stage,
 
             max_multiview_view_count,
-        }
+        })
     }
 
     /// Return a `wgpu_hal::Alignments` structure describing this adapter.