Harden Metal ICB prepass edge cases

Fixes from an adversarial pre-push review of the prepass rework:

- Gate setSupportIndirectCommandBuffers on MTLMeshRenderPipelineDescriptor
  behind macOS 14 / iOS 17 / tvOS 18.1 / visionOS 2.1: the property is one
  OS generation newer there than mesh pipelines themselves, so macOS 13 /
  iOS 16 would hit an unrecognized selector. Mesh pipelines created without
  the flag fall back to the per-draw loop.
- Raise the ICB availability floor to macOS 10.15 / iOS 13 / tvOS 13:
  inheritPipelineState needs iOS/tvOS 13, and useResource:usage: only
  participates in hazard tracking from macOS 10.15 / iOS 13, which the
  deferred generation relies on.
- Fix the tvOS entry in the ICB feature-set table (tvOS_GPUFamily1_v2
  admitted the Apple2-class A8; the floor is tvOS_GPUFamily2_v2 / A10X).
- Drop the begin_render_pass debug_assert on the deferred queue: wgpu-core
  legally records discarded-surface fixup clears (render passes) in the
  internal pre-pass before draining the queue.
- Cache generation-kernel compile failures so a driver that rejects the
  library doesn't recompile it on every multi-draw call.
- Report the retry error (the pipeline's real problem) when creation fails
  both with and without the ICB flag, and log a breadcrumb when a pipeline
  is silently downgraded to non-ICB.
- Use the stage-scoped useResource for the ICB-referenced index buffer.
  The ICB itself keeps the stage-less variant: it is consumed by command
  ingestion, not a shader stage, and the stage-scoped call was observed on
  M4 to make the executed draws render nothing.
- Rename the stale MULTI_DRAW_INDIRECT_BIND_GROUP_FALLBACK test (bind
  groups no longer force a fallback; the test now verifies bind-group
  inheritance through the ICB path), document the test-count constant's
  coupling to ICB_MIN_DRAW_COUNT, document the hal contract in the
  changelog, and correct the ICB bind-count constant's rationale.

Author: matthargett

CHANGELOG.md

@@ -44,14 +44,15 @@ Bottom level categories:
 
 ### Added/New Features
 
-#### Hal
-
-- Add `BufferBinding::buffer`, a public read accessor for the bound buffer, which was previously inaccessible to out-of-tree `wgpu_hal::Api` implementations. By @danlehmann in [#9820](https://github.com/gfx-rs/wgpu/pull/9820).
-
 #### Metal
 
 - Metal now lowers fixed-count `multi_draw_indirect` / `multi_draw_indexed_indirect` (and their mesh-task counterpart) with 8 or more draws to GPU-generated indirect command buffers. The generation compute runs in the internal pre-pass command buffer that already carries indirect-draw validation, so the render pass is never interrupted. By @matthargett in [#9640](https://github.com/gfx-rs/wgpu/pull/9640).
 
+#### Hal
+
+- Add `BufferBinding::buffer`, a public read accessor for the bound buffer, which was previously inaccessible to out-of-tree `wgpu_hal::Api` implementations. By @danlehmann in [#9820](https://github.com/gfx-rs/wgpu/pull/9820).
+- Added `CommandEncoder::encode_deferred_multi_draws` (default no-op). Backends may defer part of the work for indirect multi-draws recorded in a render pass; after ending such a pass, callers must invoke this method while recording a command buffer that the queue executes before the pass's. wgpu-core does this in its internal pre-pass. By @matthargett in [#9640](https://github.com/gfx-rs/wgpu/pull/9640).
+
 ### Changes
 
 #### naga

tests/tests/wgpu-gpu/draw_indirect.rs

@@ -38,7 +38,7 @@ pub fn all_tests(vec: &mut Vec<GpuTestInitializer>) {
         MULTI_DRAW_INDIRECT_OVER_ICB_WORKGROUP,
         MULTI_DRAW_INDIRECT_FIRST_VERTEX_AND_INSTANCE,
         MULTI_DRAW_INDIRECT_MIXED_SEQUENCE,
-        MULTI_DRAW_INDIRECT_BIND_GROUP_FALLBACK,
+        MULTI_DRAW_INDIRECT_WITH_BIND_GROUPS,
         MULTI_DRAW_INDEXED_INDIRECT_U16,
         MULTI_DRAW_INDEXED_INDIRECT_POSITIVE_BASE_VERTEX,
         MULTI_DRAW_INDEXED_INDIRECT_NEGATIVE_BASE_VERTEX,
@@ -641,6 +641,9 @@ fn create_draw_indexed_indirect_buffer(
     })
 }
 
+/// Kept in sync with `ICB_MIN_DRAW_COUNT` in `wgpu-hal/src/metal/command.rs`
+/// so these tests exercise Metal's indirect-command-buffer lowering rather
+/// than the small-count per-draw loop.
 const ICB_MULTI_DRAW_TEST_COUNT: usize = 8;
 
 async fn run_multi_draw_indirect_over_icb_workgroup(ctx: TestingContext) {
@@ -876,7 +879,9 @@ async fn run_multi_draw_indirect_mixed_sequence(ctx: TestingContext) {
     assert_all_pixels_rgba8(&data, [u8::MAX; 4]);
 }
 
-async fn run_multi_draw_indirect_bind_group_fallback(ctx: TestingContext) {
+/// Multi-draw with an active bind group; on Metal's ICB path the bind-group
+/// bindings must be inherited correctly by the generated commands.
+async fn run_multi_draw_indirect_with_bind_groups(ctx: TestingContext) {
     let shader = ctx
         .device
         .create_shader_module(wgpu::ShaderModuleDescriptor {
@@ -1287,13 +1292,13 @@ static MULTI_DRAW_INDIRECT_MIXED_SEQUENCE: GpuTestConfiguration = GpuTestConfigu
     .run_async(run_multi_draw_indirect_mixed_sequence);
 
 #[gpu_test]
-static MULTI_DRAW_INDIRECT_BIND_GROUP_FALLBACK: GpuTestConfiguration = GpuTestConfiguration::new()
+static MULTI_DRAW_INDIRECT_WITH_BIND_GROUPS: GpuTestConfiguration = GpuTestConfiguration::new()
     .parameters(
         TestParameters::default()
             .downlevel_flags(wgpu::DownlevelFlags::INDIRECT_EXECUTION)
             .limits(wgpu::Limits::downlevel_defaults()),
     )
-    .run_async(run_multi_draw_indirect_bind_group_fallback);
+    .run_async(run_multi_draw_indirect_with_bind_groups);
 
 #[gpu_test]
 static MULTI_DRAW_INDEXED_INDIRECT_U16: GpuTestConfiguration = GpuTestConfiguration::new()

wgpu-hal/src/lib.rs

@@ -1684,12 +1684,20 @@ pub trait CommandEncoder: WasmNotSendSync + fmt::Debug {
         first_instance: u32,
         instance_count: u32,
     );
+    /// # Safety
+    ///
+    /// - If `draw_count > 1`, see the deferred-work obligation on
+    ///   [`encode_deferred_multi_draws`](CommandEncoder::encode_deferred_multi_draws).
     unsafe fn draw_indirect(
         &mut self,
         buffer: &<Self::A as Api>::Buffer,
         offset: wgt::BufferAddress,
         draw_count: u32,
     );
+    /// # Safety
+    ///
+    /// - If `draw_count > 1`, see the deferred-work obligation on
+    ///   [`encode_deferred_multi_draws`](CommandEncoder::encode_deferred_multi_draws).
     unsafe fn draw_indexed_indirect(
         &mut self,
         buffer: &<Self::A as Api>::Buffer,

wgpu-hal/src/metal/adapter.rs

@@ -583,7 +583,7 @@ const INDIRECT_DRAW_DISPATCH_SUPPORT: &[MTLFeatureSet] = &[
 /// kernels; widen this table as device validation lands.
 const INDIRECT_COMMAND_BUFFERS_RENDERING_SUPPORT: &[MTLFeatureSet] = &[
     MTLFeatureSet::iOS_GPUFamily5_v1,
-    MTLFeatureSet::tvOS_GPUFamily1_v2,
+    MTLFeatureSet::tvOS_GPUFamily2_v2,
     MTLFeatureSet::macOS_GPUFamily2_v1,
 ];
 
@@ -722,8 +722,12 @@ impl super::CapabilitiesQuery {
         // - First-generation Apple TV 4K (A10X, Apple3) is allowed because it
         //   was validated directly on hardware; other Apple3/Apple4 devices
         //   stay excluded until they're validated with this backend.
-        // `MTLIndirectCommandBuffer` and friends need macOS 10.14 / iOS 12.
-        let icb_api_check = available!(macos = 10.14, ios = 12.0, tvos = 12.0, visionos = 1.0);
+        // `MTLIndirectCommandBuffer` needs macOS 10.14 / iOS 12, but the
+        // pieces this backend depends on are newer: `inheritPipelineState`
+        // needs iOS/tvOS 13, and `useResource:usage:` only participates in
+        // hazard tracking from macOS 10.15 / iOS 13 — which the deferred
+        // generation relies on to order the ICB write against its execution.
+        let icb_api_check = available!(macos = 10.15, ios = 13.0, tvos = 13.0, visionos = 1.0);
         let icb_family_support = family_check
             && if os_type == super::OsType::Macos {
                 device.supportsFamily(MTLGPUFamily::Mac2)

wgpu-hal/src/metal/command.rs

@@ -11,10 +11,10 @@ use objc2_metal::{
     MTLComputeCommandEncoder, MTLComputePassDescriptor, MTLComputePipelineState,
     MTLCounterDontSample, MTLDevice, MTLFunction, MTLIndexType, MTLIndirectCommandBuffer,
     MTLIndirectCommandBufferDescriptor, MTLIndirectCommandType, MTLLibrary, MTLLoadAction,
-    MTLPrimitiveType, MTLRenderCommandEncoder, MTLRenderPassDescriptor, MTLResidencySet,
-    MTLResidencySetDescriptor, MTLResource, MTLResourceOptions, MTLResourceUsage, MTLSamplerState,
-    MTLScissorRect, MTLSize, MTLStoreAction, MTLTexture, MTLVertexAmplificationViewMapping,
-    MTLViewport, MTLVisibilityResultMode,
+    MTLPrimitiveType, MTLRenderCommandEncoder, MTLRenderPassDescriptor, MTLRenderStages,
+    MTLResidencySet, MTLResidencySetDescriptor, MTLResource, MTLResourceOptions, MTLResourceUsage,
+    MTLSamplerState, MTLScissorRect, MTLSize, MTLStoreAction, MTLTexture,
+    MTLVertexAmplificationViewMapping, MTLViewport, MTLVisibilityResultMode,
 };
 
 use super::{
@@ -40,9 +40,12 @@ const WORD_SIZE: usize = 4;
 /// on A12 through M4 hardware.
 const ICB_MIN_DRAW_COUNT: u32 = 8;
 
-/// Upper bound declared for `maxVertexBufferBindCount` when the ICB inherits
-/// buffer bindings: Metal's per-stage buffer argument table has 31 slots, so
-/// 31 covers every possible inherited binding.
+/// Value declared for `maxVertexBufferBindCount` on ICB descriptors. The
+/// generated commands never set buffers themselves (all bindings are
+/// inherited from the encoder), so per Metal's documentation these counts
+/// only size command-side binding storage; the full 31-slot argument-table
+/// size is declared anyway because driver validation of the interaction with
+/// `inheritBuffers` has proven underdocumented across OS generations.
 const ICB_MAX_INHERITED_BUFFER_BIND_COUNT: usize = 31;
 
 // Primitive-topology tags passed to the ICB generation kernels.
@@ -68,8 +71,8 @@ pub(super) struct IcbCommandPipelines {
     indexed_u16: IcbCommandPipeline,
     indexed_u32: IcbCommandPipeline,
     /// Compiled on first use: mesh commands in ICBs need a newer OS baseline
-    /// than plain draws.
-    mesh: Option<IcbCommandPipeline>,
+    /// than plain draws. Failures are cached.
+    mesh: Option<Result<IcbCommandPipeline, crate::DeviceError>>,
 }
 
 #[derive(Clone, Debug)]
@@ -189,15 +192,14 @@ impl IcbCommandPipelines {
         &mut self,
         shared: &super::AdapterShared,
     ) -> Result<IcbCommandPipeline, crate::DeviceError> {
-        if self.mesh.is_none() {
-            let library = Self::make_library(shared, ICB_MESH_GENERATION_SHADER)?;
-            self.mesh = Some(Self::make_pipeline_from_library(
-                shared,
-                &library,
-                "wgpu_generate_mesh_mdi_icb",
-            )?);
-        }
-        Ok(self.mesh.as_ref().unwrap().clone())
+        // Failures are cached too, so a driver that rejects the mesh kernel
+        // doesn't recompile it on every mesh multi-draw.
+        self.mesh
+            .get_or_insert_with(|| {
+                let library = Self::make_library(shared, ICB_MESH_GENERATION_SHADER)?;
+                Self::make_pipeline_from_library(shared, &library, "wgpu_generate_mesh_mdi_icb")
+            })
+            .clone()
     }
 }
 
@@ -378,18 +380,20 @@ impl super::CommandEncoder {
 
     fn get_icb_command_pipelines(&self) -> Result<IcbCommandPipelines, crate::DeviceError> {
         let mut pipelines = self.shared.icb_command_pipelines.lock();
-        if pipelines.is_none() {
-            *pipelines = Some(IcbCommandPipelines::new(&self.shared)?);
-        }
-        Ok(pipelines.as_ref().unwrap().clone())
+        // A compile failure is cached so a broken driver doesn't recompile
+        // the generation library on every multi-draw call.
+        pipelines
+            .get_or_insert_with(|| IcbCommandPipelines::new(&self.shared))
+            .clone()
     }
 
     fn get_icb_mesh_command_pipeline(&self) -> Result<IcbCommandPipeline, crate::DeviceError> {
         let mut pipelines = self.shared.icb_command_pipelines.lock();
-        if pipelines.is_none() {
-            *pipelines = Some(IcbCommandPipelines::new(&self.shared)?);
-        }
-        pipelines.as_mut().unwrap().mesh(&self.shared)
+        pipelines
+            .get_or_insert_with(|| IcbCommandPipelines::new(&self.shared))
+            .as_mut()
+            .map_err(|err| err.clone())?
+            .mesh(&self.shared)
     }
 
     fn icb_primitive_type_value(
@@ -568,10 +572,12 @@ impl super::CommandEncoder {
             {
                 // The generated commands reference the index buffer via a
                 // device pointer baked in at generation time, which residency
-                // tracking can't see.
-                encoder.useResource_usage(
+                // tracking can't see. Unlike the ICB, this is an ordinary
+                // vertex-stage read, so the stage-scoped variant is correct.
+                encoder.useResource_usage_stages(
                     ProtocolObject::from_ref(&**index_buffer),
                     MTLResourceUsage::Read,
+                    MTLRenderStages::Vertex,
                 );
             }
             encoder.executeCommandsInBuffer_withRange(
@@ -1342,9 +1348,6 @@ impl crate::CommandEncoder for super::CommandEncoder {
         assert!(self.state.blit.is_none());
         assert!(self.state.compute.is_none());
         assert!(self.state.render.is_none());
-        // Multi-draws deferred by the previous pass must have been encoded via
-        // `encode_deferred_multi_draws` before another pass begins.
-        debug_assert!(self.deferred_multi_draws.is_empty());
 
         autoreleasepool(|_| {
             let descriptor = MTLRenderPassDescriptor::new();

wgpu-hal/src/metal/device.rs

@@ -1754,7 +1754,17 @@ impl crate::Device for super::Device {
             // used with indirect command buffers"), so a failed creation is
             // retried without it and the multi-draw ICB lowering then skips
             // draws issued with that pipeline.
-            let request_icb_support = self.shared.private_caps.indirect_command_buffers_rendering;
+            //
+            // On `MTLMeshRenderPipelineDescriptor` the property needs a newer
+            // OS than mesh pipelines themselves (and than the property on the
+            // standard descriptor).
+            let request_icb_support = self.shared.private_caps.indirect_command_buffers_rendering
+                && match descriptor {
+                    MetalGenericRenderPipelineDescriptor::Standard(_) => true,
+                    MetalGenericRenderPipelineDescriptor::Mesh(_) => {
+                        available!(macos = 14.0, ios = 17.0, tvos = 18.1, visionos = 2.1)
+                    }
+                };
             if request_icb_support {
                 descriptor.setSupportIndirectCommandBuffers(true);
             }
@@ -1780,15 +1790,21 @@ impl crate::Device for super::Device {
                 Err(first_err) if request_icb_support => {
                     descriptor.setSupportIndirectCommandBuffers(false);
                     supports_indirect_command_buffers = false;
-                    create(&descriptor).map_err(|_| {
-                        // Report the original error: if the pipeline is
-                        // invalid regardless, the first message is the
-                        // relevant one.
+                    let raw = create(&descriptor).map_err(|retry_err| {
+                        // The retry error is the pipeline's real problem; the
+                        // first attempt may only have failed because of the
+                        // ICB flag.
                         crate::PipelineError::Linkage(
                             wgt::ShaderStages::VERTEX | wgt::ShaderStages::FRAGMENT,
-                            format!("new_render_pipeline_state: {first_err:?}"),
+                            format!("new_render_pipeline_state: {retry_err:?}"),
                         )
-                    })?
+                    })?;
+                    log::debug!(
+                        "created render pipeline {:?} without indirect command buffer support, \
+                         multi-draws recorded with it won't use ICBs: {first_err:?}",
+                        desc.label,
+                    );
+                    raw
                 }
                 Err(e) => {
                     return Err(crate::PipelineError::Linkage(

wgpu-hal/src/metal/mod.rs

@@ -407,7 +407,7 @@ struct AdapterShared {
     private_texture_format_caps: PrivateTextureFormatCapabilities,
     settings: Settings,
     presentation_timer: time::PresentationTimer,
-    icb_command_pipelines: Mutex<Option<command::IcbCommandPipelines>>,
+    icb_command_pipelines: Mutex<Option<Result<command::IcbCommandPipelines, crate::DeviceError>>>,
 }
 
 #[cfg(send_sync)]