Metal: terrain draws its sections through indirect command buffers

Where the vendored wgpu-hal runs multi-draws from indirect command buffers
(Apple GPUs, M1 and later), terrain.wgsl takes its textures and samplers
through binding arrays of one in group 2 (wgpu keeps binding arrays out of
groups with a uniform buffer), the only way Metal accepts its pipelines in
an ICB. The shader samples the same textures the same way: the source is
rewritten at load (every use becomes element 0), and a test checks the
rewrite validates and changes nothing else.

The section multi-draws then cost the CPU one command each, so the flat
draws (one index list per frame built by a compute pass, which Metal used
from 4,096 draws) are off there, and draw call counts are reported per
multi-draw as on Vulkan. WGPU_METAL_ICB=0 gives the old path back.

Co-Authored-By: Claude Opus 5.5 <[email protected]>

Cargo.toml

@@ -19,8 +19,9 @@ large_enum_variant = "allow"
 # wgpu-hal 30.0.1 with gfx-rs/wgpu#10124: wgpu 30 creates a Vulkan descriptor
 # pool with no pool sizes while it creates the device, which Vulkan allows only
 # since 1.3.215, and Intel's Windows driver for 6th-10th gen iGPUs (HD 520/530,
-# UHD 620/630) crashes on it with an access violation (#746). Drop this and the
-# folder once a wgpu release has the fix.
+# UHD 620/630) crashes on it with an access violation (#746). Also: Metal runs
+# indirect multi-draws from indirect command buffers (third_party/wgpu-hal/
+# README.md). Drop this and the folder once wgpu releases have both.
 [patch.crates-io]
 wgpu-hal = { path = "third_party/wgpu-hal" }
 

THIRD_PARTY_NOTICES.md

@@ -33,6 +33,13 @@ SOFTWARE.
 ## wgpu-hal
 
 `third_party/wgpu-hal` is [wgpu-hal](https://github.com/gfx-rs/wgpu) 30.0.1 by the gfx-rs
-developers, with one upstream fix applied
-([gfx-rs/wgpu#10124](https://github.com/gfx-rs/wgpu/pull/10124)). It is licensed under MIT or
-Apache-2.0; both licenses are in that folder (`LICENSE.MIT`, `LICENSE.APACHE`).
+developers, with two changes:
+
+- one upstream fix applied ([gfx-rs/wgpu#10124](https://github.com/gfx-rs/wgpu/pull/10124));
+- Horizon's own change to the Metal backend: indirect multi-draws run from Metal indirect
+  command buffers generated on the GPU (`src/metal/icb.rs`), the approach upstream has proposed
+  in [gfx-rs/wgpu#9640](https://github.com/gfx-rs/wgpu/pull/9640). The vendored README lists
+  the files it touches.
+
+It is licensed under MIT or Apache-2.0; both licenses are in that folder (`LICENSE.MIT`,
+`LICENSE.APACHE`), and Horizon's changes to it are under the same terms.

crates/horizon-client/src/gfx.rs

@@ -321,7 +321,7 @@ pub struct Renderer {
     debug_line_capacity: u64,
     /// The terrain resources with the debug crosshair's projection
     /// (`crosshair_globals`).
-    crosshair_group: wgpu::BindGroup,
+    crosshair_group: TerrainGroups,
     crosshair_globals: wgpu::Buffer,
     entity_pipeline: wgpu::RenderPipeline,
     crumbling_pipeline: wgpu::RenderPipeline,
@@ -355,7 +355,7 @@ pub struct Renderer {
     entity_buffer: wgpu::Buffer,
     entity_capacity: u64,
     entity_batch: EntityBatch,
-    group: wgpu::BindGroup,
+    group: TerrainGroups,
     globals: wgpu::Buffer,
     lightmap: wgpu::Texture,
     /// The block sprites, [`SPRITE_SPLIT`] layers per array (the first is
@@ -369,6 +369,10 @@ pub struct Renderer {
     /// The mipmap levels sampled (the Mipmap Levels option), at most `mip_levels - 1`.
     mipmaps: u32,
     terrain_layout: wgpu::BindGroupLayout,
+    /// Terrain's textures as binding arrays, where multi-draws run from Metal
+    /// indirect command buffers ([`icb_multi_draws`]).
+    terrain_texture_layout: Option<wgpu::BindGroupLayout>,
+    icb_draws: bool,
     sprite_sampler: wgpu::Sampler,
     lightmap_view: wgpu::TextureView,
     lightmap_sampler: wgpu::Sampler,
@@ -497,49 +501,164 @@ fn sprite_views(
         })
         .collect()
 }
-/// terrain.wgsl's group 0.
+/// The group of terrain.wgsl's textures and samplers where multi-draws run
+/// from indirect command buffers ([`terrain_with_binding_arrays`]): wgpu
+/// keeps binding arrays out of groups that hold a uniform buffer.
+const TERRAIN_TEXTURE_GROUP: u32 = 2;
+/// terrain.wgsl's bind groups: group 0, and its textures and samplers in
+/// [`TERRAIN_TEXTURE_GROUP`] as binding arrays of one where multi-draws run
+/// from indirect command buffers.
+struct TerrainGroups {
+    globals: wgpu::BindGroup,
+    textures: Option<wgpu::BindGroup>,
+}
+impl TerrainGroups {
+    fn bind(&self, pass: &mut wgpu::RenderPass<'_>) {
+        pass.set_bind_group(0, &self.globals, &[]);
+        if let Some(textures) = &self.textures {
+            pass.set_bind_group(TERRAIN_TEXTURE_GROUP, textures, &[]);
+        }
+    }
+}
+/// terrain.wgsl's groups; `texture_layout` (the binding-array layout) puts
+/// the textures and samplers in a group of their own.
 #[allow(clippy::too_many_arguments)]
 fn terrain_group(
     device: &wgpu::Device,
     label: &str,
     layout: &wgpu::BindGroupLayout,
+    texture_layout: Option<&wgpu::BindGroupLayout>,
     globals: &wgpu::Buffer,
     sprites: &[wgpu::TextureView],
     sprite_sampler: &wgpu::Sampler,
     lightmap: &wgpu::TextureView,
     lightmap_sampler: &wgpu::Sampler,
-) -> wgpu::BindGroup {
-    let view = |binding, view| wgpu::BindGroupEntry {
-        binding,
-        resource: wgpu::BindingResource::TextureView(view),
+) -> TerrainGroups {
+    fn texture<'a>(
+        arrays: bool,
+        binding: u32,
+        view: &'a [&'a wgpu::TextureView; 1],
+    ) -> wgpu::BindGroupEntry<'a> {
+        let resource = if arrays {
+            wgpu::BindingResource::TextureViewArray(view)
+        } else {
+            wgpu::BindingResource::TextureView(view[0])
+        };
+        wgpu::BindGroupEntry { binding, resource }
+    }
+    fn sampler<'a>(
+        arrays: bool,
+        binding: u32,
+        sampler: &'a [&'a wgpu::Sampler; 1],
+    ) -> wgpu::BindGroupEntry<'a> {
+        let resource = if arrays {
+            wgpu::BindingResource::SamplerArray(sampler)
+        } else {
+            wgpu::BindingResource::Sampler(sampler[0])
+        };
+        wgpu::BindGroupEntry { binding, resource }
+    }
+    let arrays = texture_layout.is_some();
+    // One-element slices for binding arrays; their first element otherwise.
+    let sprites: Vec<[&wgpu::TextureView; 1]> = sprites.iter().map(|v| [v]).collect();
+    let (lightmap, sprite_sampler, lightmap_sampler) =
+        ([lightmap], [sprite_sampler], [lightmap_sampler]);
+    let globals = wgpu::BindGroupEntry {
+        binding: 0,
+        resource: globals.as_entire_binding(),
     };
-    device.create_bind_group(&wgpu::BindGroupDescriptor {
-        label: Some(label),
-        layout,
-        entries: &[
-            wgpu::BindGroupEntry {
-                binding: 0,
-                resource: globals.as_entire_binding(),
-            },
-            view(1, &sprites[0]),
-            wgpu::BindGroupEntry {
-                binding: 2,
-                resource: wgpu::BindingResource::Sampler(sprite_sampler),
-            },
-            view(3, lightmap),
-            wgpu::BindGroupEntry {
-                binding: 4,
-                resource: wgpu::BindingResource::Sampler(lightmap_sampler),
-            },
-            view(5, &sprites[1]),
-            view(6, &sprites[2]),
-            view(7, &sprites[3]),
-            view(8, &sprites[4]),
-            view(9, &sprites[5]),
-            view(10, &sprites[6]),
-            view(11, &sprites[7]),
-        ],
-    })
+    let textures = [
+        texture(arrays, 1, &sprites[0]),
+        sampler(arrays, 2, &sprite_sampler),
+        texture(arrays, 3, &lightmap),
+        sampler(arrays, 4, &lightmap_sampler),
+        texture(arrays, 5, &sprites[1]),
+        texture(arrays, 6, &sprites[2]),
+        texture(arrays, 7, &sprites[3]),
+        texture(arrays, 8, &sprites[4]),
+        texture(arrays, 9, &sprites[5]),
+        texture(arrays, 10, &sprites[6]),
+        texture(arrays, 11, &sprites[7]),
+    ];
+    let group = |layout, entries: &[wgpu::BindGroupEntry<'_>]| {
+        device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some(label),
+            layout,
+            entries,
+        })
+    };
+    match texture_layout {
+        Some(texture_layout) => TerrainGroups {
+            globals: group(layout, &[globals]),
+            textures: Some(group(texture_layout, &textures)),
+        },
+        None => {
+            let mut entries = vec![globals];
+            entries.extend(textures);
+            TerrainGroups {
+                globals: group(layout, &entries),
+                textures: None,
+            }
+        }
+    }
+}
+/// Whether the vendored wgpu-hal draws indirect multi-draws from Metal
+/// indirect command buffers (`metal/icb.rs`): one command for a whole
+/// multi-draw instead of one draw call per entry. It does on Apple GPUs from
+/// the M1 on (`Icbs::new` there: keep the two in step) unless
+/// `WGPU_METAL_ICB=0`. Pipelines only qualify with no texture or sampler bound
+/// directly, so terrain then takes its own through binding arrays
+/// ([`terrain_with_binding_arrays`]).
+fn icb_multi_draws(info: &wgpu::AdapterInfo, features: wgpu::Features) -> bool {
+    info.backend == wgpu::Backend::Metal
+        && info.name.starts_with("Apple M")
+        && features.contains(
+            wgpu::Features::TEXTURE_BINDING_ARRAY | wgpu::Features::INDIRECT_FIRST_INSTANCE,
+        )
+        && std::env::var_os("WGPU_METAL_ICB").is_none_or(|v| v != "0")
+}
+/// terrain.wgsl with every texture and sampler of group 0 declared as a
+/// binding array of one in [`TERRAIN_TEXTURE_GROUP`] (same bindings) and used
+/// through its element 0. The shader samples
+/// the same textures the same way; on Metal they are then reached through an
+/// argument buffer, which pipelines drawn from indirect command buffers need
+/// ([`icb_multi_draws`]).
+fn terrain_with_binding_arrays(source: &str) -> String {
+    let mut names: Vec<&str> = Vec::new();
+    let mut out = String::with_capacity(source.len() + 1024);
+    out.push_str("enable wgpu_binding_array;\n");
+    for line in source.lines() {
+        let declaration = line
+            .strip_prefix("@group(0) @binding(")
+            .and_then(|_| line.split_once(" var "))
+            .and_then(|(head, rest)| {
+                let (name, ty) = rest.split_once(": ")?;
+                let ty = ty.strip_suffix(';')?;
+                (ty.starts_with("texture_") || ty == "sampler").then_some((head, name, ty))
+            });
+        if let Some((head, name, ty)) = declaration {
+            names.push(name);
+            let head = head.replacen("@group(0)", &format!("@group({TERRAIN_TEXTURE_GROUP})"), 1);
+            out.push_str(&format!("{head} var {name}: binding_array<{ty}, 1>;\n"));
+            continue;
+        }
+        // Every other use of those names (whole identifiers) is element 0.
+        let mut rest = line;
+        while let Some(start) = rest.find(|c: char| c.is_ascii_alphanumeric() || c == '_') {
+            let end = rest[start..]
+                .find(|c: char| !(c.is_ascii_alphanumeric() || c == '_'))
+                .map_or(rest.len(), |n| start + n);
+            let word = &rest[start..end];
+            out.push_str(&rest[..end]);
+            if names.contains(&word) && !rest[..start].ends_with('.') {
+                out.push_str("[0]");
+            }
+            rest = &rest[end..];
+        }
+        out.push_str(rest);
+        out.push('\n');
+    }
+    out
 }
 /// CPU texture uploads prepared by the resource worker. Consumed by renderer
 /// construction so the temporary mip chains are released after GPU upload.
@@ -780,11 +899,19 @@ async fn open_device(
         }
         // Mappable primary buffers: section pages that no longer fit in
         // video memory go to system memory.
+        // Binding arrays on Metal: terrain takes its textures through them
+        // where multi-draws use indirect command buffers (`icb_multi_draws`).
+        let binding_arrays = if info.backend == wgpu::Backend::Metal {
+            wgpu::Features::TEXTURE_BINDING_ARRAY
+        } else {
+            wgpu::Features::empty()
+        };
         let features = timing::GpuTimer::features(&adapter)
             | cull::Cull::features(&adapter)
             | (adapter.features()
                 & (wgpu::Features::INDIRECT_FIRST_INSTANCE
-                    | wgpu::Features::MAPPABLE_PRIMARY_BUFFERS));
+                    | wgpu::Features::MAPPABLE_PRIMARY_BUFFERS
+                    | binding_arrays));
         let driver = format!("{} {}", info.driver, info.driver_info);
         startup_step(format_args!(
             "requesting the device (driver {}) with features [{features}] and the adapter's \
@@ -973,9 +1100,14 @@ impl Renderer {
             startup_step(format_args!("creating the pipelines"));
         }
         let gpu_timer = timing::GpuTimer::new(&device, &queue);
+        let icb_draws = icb_multi_draws(&adapter.get_info(), device.features());
         let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
             label: Some("Horizon terrain"),
-            source: wgpu::ShaderSource::Wgsl(include_str!("terrain.wgsl").into()),
+            source: wgpu::ShaderSource::Wgsl(if icb_draws {
+                terrain_with_binding_arrays(include_str!("terrain.wgsl")).into()
+            } else {
+                include_str!("terrain.wgsl").into()
+            }),
         });
         let entry = |binding, visibility, ty| wgpu::BindGroupLayoutEntry {
             binding,
@@ -997,58 +1129,84 @@ impl Renderer {
             ("SPRITE_ARRAYS", f64::from(sprite_arrays)),
             ("SPRITE_SPLIT", f64::from(split)),
         ];
-        let bindings = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
-            label: Some("Horizon terrain bindings"),
-            entries: &[
-                entry(
-                    0,
-                    vf,
-                    wgpu::BindingType::Buffer {
-                        ty: wgpu::BufferBindingType::Uniform,
-                        has_dynamic_offset: false,
-                        min_binding_size: None,
-                    },
-                ),
-                entry(
-                    1,
-                    vf,
-                    wgpu::BindingType::Texture {
-                        sample_type: wgpu::TextureSampleType::Float { filterable: true },
-                        view_dimension: wgpu::TextureViewDimension::D2Array,
-                        multisampled: false,
-                    },
-                ),
-                entry(
-                    1 + 1,
-                    vf,
-                    wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
-                ),
-                entry(
-                    3,
-                    vf,
-                    wgpu::BindingType::Texture {
-                        sample_type: wgpu::TextureSampleType::Float { filterable: true },
-                        view_dimension: wgpu::TextureViewDimension::D2,
-                        multisampled: false,
-                    },
-                ),
-                entry(
-                    4,
-                    vf,
-                    wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
-                ),
-                entry(5, vf, sprite_array),
-                entry(6, vf, sprite_array),
-                entry(7, vf, sprite_array),
-                entry(8, vf, sprite_array),
-                entry(9, vf, sprite_array),
-                entry(10, vf, sprite_array),
-                entry(11, vf, sprite_array),
-            ],
-        });
+        let globals_entry = entry(
+            0,
+            vf,
+            wgpu::BindingType::Buffer {
+                ty: wgpu::BufferBindingType::Uniform,
+                has_dynamic_offset: false,
+                min_binding_size: None,
+            },
+        );
+        let texture_entries = [
+            entry(
+                1,
+                vf,
+                wgpu::BindingType::Texture {
+                    sample_type: wgpu::TextureSampleType::Float { filterable: true },
+                    view_dimension: wgpu::TextureViewDimension::D2Array,
+                    multisampled: false,
+                },
+            ),
+            entry(
+                1 + 1,
+                vf,
+                wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
+            ),
+            entry(
+                3,
+                vf,
+                wgpu::BindingType::Texture {
+                    sample_type: wgpu::TextureSampleType::Float { filterable: true },
+                    view_dimension: wgpu::TextureViewDimension::D2,
+                    multisampled: false,
+                },
+            ),
+            entry(
+                4,
+                vf,
+                wgpu::BindingType::Sampler(wgpu::SamplerBindingType::Filtering),
+            ),
+            entry(5, vf, sprite_array),
+            entry(6, vf, sprite_array),
+            entry(7, vf, sprite_array),
+            entry(8, vf, sprite_array),
+            entry(9, vf, sprite_array),
+            entry(10, vf, sprite_array),
+            entry(11, vf, sprite_array),
+        ];
+        let bind_group_layout = |label, entries: &[wgpu::BindGroupLayoutEntry]| {
+            device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+                label: Some(label),
+                entries,
+            })
+        };
+        let (bindings, texture_bindings) = if icb_draws {
+            (
+                bind_group_layout("Horizon terrain bindings", &[globals_entry]),
+                Some(bind_group_layout(
+                    "Horizon terrain textures",
+                    // Binding arrays of one.
+                    &texture_entries.map(|entry| wgpu::BindGroupLayoutEntry {
+                        count: Some(std::num::NonZeroU32::MIN),
+                        ..entry
+                    }),
+                )),
+            )
+        } else {
+            let mut entries = vec![globals_entry];
+            entries.extend(texture_entries);
+            (
+                bind_group_layout("Horizon terrain bindings", &entries),
+                None,
+            )
+        };
         let layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
             label: Some("Horizon terrain layout"),
-            bind_group_layouts: &[Some(&bindings)],
+            bind_group_layouts: &match &texture_bindings {
+                Some(textures) => vec![Some(&bindings), None, Some(textures)],
+                None => vec![Some(&bindings)],
+            },
             ..Default::default()
         });
         let instance_attributes = wgpu::vertex_attr_array![5=>Float32x4];
@@ -1336,6 +1494,7 @@ impl Renderer {
             &device,
             "Horizon terrain resources",
             &bindings,
+            texture_bindings.as_ref(),
             &globals,
             &sprite_views,
             &sprite_sampler,
@@ -1346,6 +1505,7 @@ impl Renderer {
             &device,
             "Debug crosshair resources",
             &bindings,
+            texture_bindings.as_ref(),
             &crosshair_globals,
             &sprite_views,
             &sprite_sampler,
@@ -2002,7 +2162,7 @@ impl Renderer {
             1 << 20,
         );
         // Sections as one index list per frame where multi-draw is emulated.
-        let flat = flat::wanted(&adapter, &device).then(|| {
+        let flat = flat::wanted(&adapter, &device, icb_draws).then(|| {
             let buffers = [Some(wgpu::VertexBufferLayout {
                 array_stride: std::mem::size_of::<TerrainVertex>() as u64,
                 step_mode: wgpu::VertexStepMode::Vertex,
@@ -2137,6 +2297,8 @@ impl Renderer {
             mip_levels,
             mipmaps: mip_levels - 1,
             terrain_layout: bindings,
+            terrain_texture_layout: texture_bindings,
+            icb_draws,
             sprite_sampler,
             lightmap_view,
             lightmap_sampler,
@@ -2636,6 +2798,7 @@ impl Renderer {
             &self.device,
             "Horizon terrain resources",
             &self.terrain_layout,
+            self.terrain_texture_layout.as_ref(),
             &self.globals,
             &views,
             &self.sprite_sampler,
@@ -3188,19 +3351,21 @@ impl Renderer {
         self.translucency_cursor += budget;
         let translucent_draws = self.draws.len() as u32 - solid_draws;
         self.section_draws = self.draws.len();
-        // Metal replays a multi-draw as one draw call per entry.
-        self.section_draw_calls = if self.multi_draw && self.info.backend != wgpu::Backend::Metal {
-            (0..self.draws.len())
-                .filter(|&i| {
-                    i == 0
-                        || i == phase1_draws as usize
-                        || i == solid_draws as usize
-                        || self.draw_pages[i] != self.draw_pages[i - 1]
-                })
-                .count()
-        } else {
-            self.draws.len()
-        };
+        // Metal replays a multi-draw as one draw call per entry, unless it
+        // runs from an indirect command buffer.
+        self.section_draw_calls =
+            if self.multi_draw && (self.info.backend != wgpu::Backend::Metal || self.icb_draws) {
+                (0..self.draws.len())
+                    .filter(|&i| {
+                        i == 0
+                            || i == phase1_draws as usize
+                            || i == solid_draws as usize
+                            || self.draw_pages[i] != self.draw_pages[i - 1]
+                    })
+                    .count()
+            } else {
+                self.draws.len()
+            };
         let debug_depth = frame.debug_lines.len() as u32;
         let debug_top = frame.debug_lines_on_top.len() as u32;
         let crosshair = frame.debug_crosshair.len() as u32;
@@ -3496,7 +3661,7 @@ impl Renderer {
                 }),
                 ..Default::default()
             });
-            pass.set_bind_group(0, &self.group, &[]);
+            self.group.bind(&mut pass);
             // LevelRenderer: no sky inside lava or powder snow (sky fog ends within 2 blocks).
             let sky_visible = match frame.fog {
                 Some(f) => f.sky_end > 2.0,
@@ -3509,18 +3674,18 @@ impl Renderer {
                 }
                 if frame.title {
                     self.panorama.draw(&mut pass);
-                    pass.set_bind_group(0, &self.group, &[]);
+                    self.group.bind(&mut pass);
                 } else if celestial_vertices > 0 {
                     pass.set_pipeline(&self.celestial_pipeline);
                     pass.set_bind_group(0, &self.entity_group, &[]);
                     pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint32);
                     pass.set_vertex_buffer(0, self.celestial_buffer.slice(..));
                     pass.draw_indexed(0..(celestial_vertices as u32 / 4) * 6, 0, 0..1);
-                    pass.set_bind_group(0, &self.group, &[]);
+                    self.group.bind(&mut pass);
                 }
                 if !frame.title && frame.stars.is_some() {
                     self.stars.draw(&mut pass);
-                    pass.set_bind_group(0, &self.group, &[]);
+                    self.group.bind(&mut pass);
                 }
                 if !frame.title && frame.dark_disc {
                     pass.set_pipeline(&self.dark_disc);
@@ -3533,7 +3698,7 @@ impl Renderer {
                     flat.draw(&mut pass, &flat.calls[..flat.solid_calls], |p| {
                         self.arena.buffer(p)
                     });
-                    pass.set_bind_group(0, &self.group, &[]);
+                    self.group.bind(&mut pass);
                     pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint32);
                 }
                 None => {
@@ -3576,7 +3741,7 @@ impl Renderer {
                                 ),
                                 ..Default::default()
                             });
-                            pass.set_bind_group(0, &self.group, &[]);
+                            self.group.bind(&mut pass);
                             pass.set_pipeline(&self.solid);
                             pass.set_index_buffer(
                                 self.indices.slice(..),
@@ -3600,14 +3765,14 @@ impl Renderer {
                 pass.set_bind_group(0, &self.entity_group, &[]);
                 pass.set_vertex_buffer(0, self.entity_buffer.slice(..));
                 pass.draw_indexed(0..entity_quads * 6, 0, 0..1);
-                pass.set_bind_group(0, &self.group, &[]);
+                self.group.bind(&mut pass);
             }
             if crumbling_quads > 0 {
                 pass.set_pipeline(&self.crumbling_pipeline);
                 pass.set_bind_group(0, &self.entity_group, &[]);
                 pass.set_vertex_buffer(0, self.entity_buffer.slice(..));
                 pass.draw_indexed(0..crumbling_quads * 6, (entity_quads * 4) as i32, 0..1);
-                pass.set_bind_group(0, &self.group, &[]);
+                self.group.bind(&mut pass);
             }
             // BoatRenderer's water patch (RenderTypes.waterMask): depth only,
             // before translucent terrain, so no water shows inside the boat.
@@ -3620,14 +3785,14 @@ impl Renderer {
                     ((entity_quads + crumbling_quads + translucent_crumbling_quads) * 4) as i32,
                     0..1,
                 );
-                pass.set_bind_group(0, &self.group, &[]);
+                self.group.bind(&mut pass);
             }
             if additive_quads > 0 {
                 pass.set_pipeline(&self.additive_pipeline);
                 pass.set_bind_group(0, &self.entity_group, &[]);
                 pass.set_vertex_buffer(0, self.entity_buffer.slice(..));
                 pass.draw_indexed(0..additive_quads * 6, (world_quads * 4) as i32, 0..1);
-                pass.set_bind_group(0, &self.group, &[]);
+                self.group.bind(&mut pass);
             }
             if outline_edges > 0 && !frame.outline_after_terrain {
                 pass.set_pipeline(&self.lines);
@@ -3650,7 +3815,7 @@ impl Renderer {
                     flat.draw(&mut pass, &flat.calls[flat.solid_calls..], |p| {
                         self.arena.buffer(p)
                     });
-                    pass.set_bind_group(0, &self.group, &[]);
+                    self.group.bind(&mut pass);
                     pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint32);
                 }
                 None => {
@@ -3684,7 +3849,7 @@ impl Renderer {
                     let first = side as u32 * 6;
                     pass.draw_indexed(first..first + 6, 0, 0..1);
                 }
-                pass.set_bind_group(0, &self.group, &[]);
+                self.group.bind(&mut pass);
             }
             if translucent_crumbling_quads > 0 {
                 pass.set_pipeline(&self.crumbling_pipeline);
@@ -3836,7 +4001,7 @@ impl Renderer {
                 ..Default::default()
             });
             pass.set_pipeline(&self.debug_lines_on_top);
-            pass.set_bind_group(0, &self.crosshair_group, &[]);
+            self.crosshair_group.bind(&mut pass);
             pass.set_vertex_buffer(0, self.debug_line_buffer.slice(..));
             let first = debug_depth + debug_top;
             pass.draw(0..4, first..first + crosshair);
@@ -4313,6 +4478,66 @@ mod tests {
         .validate(&module)
         .unwrap();
     }
+    /// The binding-array terrain shader (Metal with indirect command buffers)
+    /// validates, declares every group 0 texture and sampler as an array of
+    /// one and reaches them only through element 0, and keeps everything else.
+    #[test]
+    fn terrain_with_binding_arrays_uses_element_zero_of_each() {
+        let source = include_str!("terrain.wgsl");
+        let arrays = super::terrain_with_binding_arrays(source);
+        let module = naga::front::wgsl::parse_str(&arrays).unwrap();
+        naga::valid::Validator::new(
+            naga::valid::ValidationFlags::all(),
+            naga::valid::Capabilities::all(),
+        )
+        .validate(&module)
+        .unwrap();
+        let arrays_of_one = module
+            .global_variables
+            .iter()
+            .filter(|(_, v)| {
+                matches!(
+                    module.types[v.ty].inner,
+                    naga::TypeInner::BindingArray {
+                        size: naga::ArraySize::Constant(n),
+                        ..
+                    } if n.get() == 1
+                )
+            })
+            .count();
+        assert_eq!(
+            arrays_of_one, 11,
+            "8 sprite arrays, the lightmap, 2 samplers"
+        );
+        for (before, after) in [
+            (
+                "textureSample(sprites, sprite_sampler, uv",
+                "textureSample(sprites[0], sprite_sampler[0], uv",
+            ),
+            (
+                "textureSampleGrad(sprites_7, sprite_sampler,",
+                "textureSampleGrad(sprites_7[0], sprite_sampler[0],",
+            ),
+            (
+                "textureSampleLevel(lightmap, lightmap_sampler,",
+                "textureSampleLevel(lightmap[0], lightmap_sampler[0],",
+            ),
+            ("fn sample_lightmap(", "fn sample_lightmap("),
+        ] {
+            assert!(source.contains(before), "{before}");
+            assert!(arrays.contains(after), "{after}");
+        }
+        // Nothing else changes: removing the array syntax gives the source back.
+        let mut plain = arrays
+            .strip_prefix("enable wgpu_binding_array;\n")
+            .unwrap()
+            .replace("[0]", "")
+            .replace("@group(2)", "@group(0)");
+        for ty in ["texture_2d_array<f32>", "texture_2d<f32>", "sampler"] {
+            plain = plain.replace(&format!("binding_array<{ty}, 1>"), ty);
+        }
+        assert_eq!(plain.trim_end(), source.trim_end());
+    }
     #[test]
     fn terrain_attributes_follow_the_terrain_vertex() {
         use super::TerrainVertex;

crates/horizon-client/src/gfx/flat.rs

@@ -71,10 +71,15 @@ pub struct Flat {
 /// saves 1.3 ms.
 const MIN_DRAWS: usize = 4096;
 
-/// Whether flat draws apply: the backend emulates multi-draw, and vertex
+/// Whether flat draws apply: the backend emulates multi-draw (Metal, unless
+/// multi-draws run from indirect command buffers: `icb_draws`), and vertex
 /// shaders can read storage buffers large enough for an index page.
 /// `HORIZON_FLAT_DRAWS=1` draws every frame flat, `0` none.
-pub fn wanted(adapter: &wgpu::Adapter, device: &wgpu::Device) -> bool {
+pub fn wanted(adapter: &wgpu::Adapter, device: &wgpu::Device, icb_draws: bool) -> bool {
+    // terrain.wgsl's textures are in group 2 then, and flat draws use group 1.
+    if icb_draws {
+        return false;
+    }
     match std::env::var("HORIZON_FLAT_DRAWS").as_deref() {
         Ok("0") => return false,
         Ok("1") => return supported(adapter, device),

crates/horizon-client/tests/multi_draw.rs

@@ -0,0 +1,300 @@
+//! Indirect multi-draws draw exactly what their entries say, on whatever
+//! path the backend takes: on Metal, the vendored wgpu-hal lowers multi-draws
+//! of 16 or more entries under pipelines without direct textures to indirect
+//! command buffers generated on the GPU (third_party/wgpu-hal, `metal/icb.rs`).
+//!
+//! Each entry draws one quad covering one pixel of a 1-row target, chosen by
+//! its instance (first_instance), so the target shows which entries drew, with
+//! which instance and which vertices: entries with no instances or no indices
+//! must leave their pixel clear, offsets into the index buffer must be kept.
+
+use wgpu::util::DeviceExt;
+
+const WIDTH: u32 = 96;
+
+const SHADER: &str = r"
+struct Out {
+    @builtin(position) clip: vec4<f32>,
+    @location(0) @interpolate(flat) value: u32,
+}
+// Vertex v is corner v % 4 of quad v / 4 (indexed draws), at pixel
+// `instance`; the value written tells the pixel which instance and which
+// quad drew it.
+@vertex fn vs(@builtin(vertex_index) v: u32, @builtin(instance_index) instance: u32) -> Out {
+    return corner(v % 4u, v / 4u, instance);
+}
+// Non-indexed draws: vertex v is corner 0 1 2 0 2 3 [v % 6] of quad v / 6.
+@vertex fn vs_list(@builtin(vertex_index) v: u32, @builtin(instance_index) instance: u32) -> Out {
+    let corners = array(0u, 1u, 2u, 0u, 2u, 3u);
+    return corner(corners[v % 6u], v / 6u, instance);
+}
+fn corner(c: u32, quad: u32, instance: u32) -> Out {
+    let x0 = (f32(instance) / f32(WIDTH)) * 2.0 - 1.0;
+    let x1 = (f32(instance + 1u) / f32(WIDTH)) * 2.0 - 1.0;
+    let x = select(x0, x1, c == 1u || c == 2u);
+    let y = select(-1.0, 1.0, c >= 2u);
+    var out: Out;
+    out.clip = vec4(x, y, 0.0, 1.0);
+    out.value = (instance + 1u) | (quad << 16u);
+    return out;
+}
+@fragment fn fs(in: Out) -> @location(0) vec4<u32> {
+    return vec4(in.value, 0u, 0u, 1u);
+}
+";
+
+fn device() -> Option<(wgpu::Device, wgpu::Queue)> {
+    let instance = wgpu::Instance::new(wgpu::InstanceDescriptor::new_without_display_handle());
+    let adapter = pollster::block_on(instance.request_adapter(&Default::default())).ok()?;
+    if !adapter
+        .features()
+        .contains(wgpu::Features::INDIRECT_FIRST_INSTANCE)
+    {
+        return None;
+    }
+    pollster::block_on(adapter.request_device(&wgpu::DeviceDescriptor {
+        required_features: wgpu::Features::INDIRECT_FIRST_INSTANCE,
+        ..Default::default()
+    }))
+    .ok()
+}
+
+/// Draws `entries` (indexed with `format` when given, else non-indexed) as
+/// one multi-draw per run of `runs` (first entry, entries) and returns the
+/// target's row.
+fn draw(
+    device: &wgpu::Device,
+    queue: &wgpu::Queue,
+    entries: &[[u32; 5]],
+    indexed: Option<wgpu::IndexFormat>,
+    runs: &[(u32, u32)],
+) -> Vec<u32> {
+    let shader = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+        label: None,
+        source: wgpu::ShaderSource::Wgsl(SHADER.replace("WIDTH", &format!("{WIDTH}.0")).into()),
+    });
+    let format = wgpu::TextureFormat::R32Uint;
+    let pipeline = device.create_render_pipeline(&wgpu::RenderPipelineDescriptor {
+        label: Some("multi-draw test"),
+        layout: None,
+        vertex: wgpu::VertexState {
+            module: &shader,
+            entry_point: Some(if indexed.is_some() { "vs" } else { "vs_list" }),
+            compilation_options: Default::default(),
+            buffers: &[],
+        },
+        primitive: wgpu::PrimitiveState::default(),
+        depth_stencil: None,
+        multisample: Default::default(),
+        fragment: Some(wgpu::FragmentState {
+            module: &shader,
+            entry_point: Some("fs"),
+            compilation_options: Default::default(),
+            targets: &[Some(format.into())],
+        }),
+        multiview_mask: None,
+        cache: None,
+    });
+    let target = device.create_texture(&wgpu::TextureDescriptor {
+        label: None,
+        size: wgpu::Extent3d {
+            width: WIDTH,
+            height: 1,
+            depth_or_array_layers: 1,
+        },
+        mip_level_count: 1,
+        sample_count: 1,
+        dimension: wgpu::TextureDimension::D2,
+        format,
+        usage: wgpu::TextureUsages::RENDER_ATTACHMENT | wgpu::TextureUsages::COPY_SRC,
+        view_formats: &[],
+    });
+    // Two quads per 4 vertices: corners 0 1 2, 0 2 3. Index i of quad q is
+    // 4q + corner, so first_index and base_vertex both show in the result.
+    let quads: Vec<u32> = (0..64u32)
+        .flat_map(|q| [0, 1, 2, 0, 2, 3].map(|c| q * 4 + c))
+        .collect();
+    // Two padding indices in front: the index buffer is bound at an offset.
+    let mut index_bytes = Vec::new();
+    match indexed {
+        Some(wgpu::IndexFormat::Uint16) => {
+            for i in [0u32, 0].iter().chain(&quads) {
+                index_bytes.extend_from_slice(&(*i as u16).to_le_bytes());
+            }
+        }
+        _ => {
+            for i in [0u32, 0].iter().chain(&quads) {
+                index_bytes.extend_from_slice(&i.to_le_bytes());
+            }
+        }
+    }
+    let index_offset = match indexed {
+        Some(wgpu::IndexFormat::Uint16) => 4,
+        _ => 8,
+    };
+    let indices = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+        label: None,
+        contents: &index_bytes,
+        usage: wgpu::BufferUsages::INDEX,
+    });
+    let args: Vec<u32> = match indexed {
+        Some(_) => entries.iter().flatten().copied().collect(),
+        // Non-indexed: vertex count, instances, first vertex, first instance.
+        None => entries
+            .iter()
+            .flat_map(|e| [e[0], e[1], e[2], e[4]])
+            .collect(),
+    };
+    // One unused entry in front: the multi-draw starts at an offset.
+    let stride = if indexed.is_some() { 5 } else { 4 };
+    let mut padded = vec![7u32; stride];
+    padded.extend(args);
+    let indirect = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+        label: None,
+        contents: bytemuck::cast_slice(&padded),
+        usage: wgpu::BufferUsages::INDIRECT,
+    });
+    let readback = device.create_buffer(&wgpu::BufferDescriptor {
+        label: None,
+        size: u64::from(WIDTH) * 4,
+        usage: wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::MAP_READ,
+        mapped_at_creation: false,
+    });
+    let view = target.create_view(&Default::default());
+    let mut encoder = device.create_command_encoder(&Default::default());
+    {
+        let mut pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+            label: None,
+            color_attachments: &[Some(wgpu::RenderPassColorAttachment {
+                view: &view,
+                depth_slice: None,
+                resolve_target: None,
+                ops: wgpu::Operations {
+                    load: wgpu::LoadOp::Clear(wgpu::Color::TRANSPARENT),
+                    store: wgpu::StoreOp::Store,
+                },
+            })],
+            ..Default::default()
+        });
+        pass.set_pipeline(&pipeline);
+        if let Some(format) = indexed {
+            pass.set_index_buffer(indices.slice(index_offset..), format);
+        }
+        for &(first, count) in runs {
+            let offset = (u64::from(first) + 1) * stride as u64 * 4;
+            match indexed {
+                Some(_) => pass.multi_draw_indexed_indirect(&indirect, offset, count),
+                None => pass.multi_draw_indirect(&indirect, offset, count),
+            }
+        }
+    }
+    encoder.copy_texture_to_buffer(
+        target.as_image_copy(),
+        wgpu::TexelCopyBufferInfo {
+            buffer: &readback,
+            layout: wgpu::TexelCopyBufferLayout {
+                offset: 0,
+                bytes_per_row: None,
+                rows_per_image: None,
+            },
+        },
+        wgpu::Extent3d {
+            width: WIDTH,
+            height: 1,
+            depth_or_array_layers: 1,
+        },
+    );
+    queue.submit([encoder.finish()]);
+    readback.slice(..).map_async(wgpu::MapMode::Read, |_| {});
+    device
+        .poll(wgpu::PollType::wait_indefinitely())
+        .expect("poll");
+    let mapped = readback.slice(..).get_mapped_range().expect("mapped");
+    bytemuck::cast_slice(&mapped).to_vec()
+}
+
+/// `count` entries [index count, instances, first index, base vertex,
+/// first instance]: entry i draws quad (i * 5) % 61 at pixel i + 3, every
+/// seventh draws no instances and every eleventh no indices.
+fn entries(count: u32) -> Vec<[u32; 5]> {
+    (0..count)
+        .map(|i| {
+            let quad = (i * 5) % 61;
+            let indices = if i % 11 == 10 { 0 } else { 6 };
+            let instances = if i % 7 == 6 { 0 } else { 1 };
+            // Half the entries reach their quad through base_vertex.
+            let (first, base) = if i % 2 == 0 {
+                (quad * 6, 0)
+            } else {
+                (0, quad * 4)
+            };
+            [indices, instances, first, base, i + 3]
+        })
+        .collect()
+}
+
+/// What `entries` must leave in the row: entry i's pixel holds its
+/// instance + 1 and its quad, the others nothing.
+fn expected(entries: &[[u32; 5]], non_indexed: bool) -> Vec<u32> {
+    let mut row = vec![0; WIDTH as usize];
+    for (i, e) in entries.iter().enumerate() {
+        if e[0] == 0 || e[1] == 0 {
+            continue;
+        }
+        let quad = if non_indexed {
+            e[2] / 6
+        } else {
+            (i as u32 * 5) % 61
+        };
+        row[e[4] as usize] = (e[4] + 1) | (quad << 16);
+    }
+    row
+}
+
+#[test]
+fn multi_draws_draw_every_entry_as_given() {
+    let Some((device, queue)) = device() else {
+        eprintln!("no GPU with indirect first instance: skipped");
+        return;
+    };
+    for count in [3, 16, 90] {
+        let entries = entries(count);
+        for format in [wgpu::IndexFormat::Uint16, wgpu::IndexFormat::Uint32] {
+            let row = draw(&device, &queue, &entries, Some(format), &[(0, count)]);
+            assert_eq!(row, expected(&entries, false), "{count} {format:?} draws");
+        }
+    }
+    // Consecutive multi-draws of one list, with short ones between (which
+    // Metal draws one by one): the long ones share a generation dispatch.
+    let entries = entries(90);
+    let runs = [(0, 20), (20, 3), (23, 30), (53, 1), (54, 36)];
+    let row = draw(
+        &device,
+        &queue,
+        &entries,
+        Some(wgpu::IndexFormat::Uint32),
+        &runs,
+    );
+    assert_eq!(row, expected(&entries, false), "runs {runs:?}");
+}
+
+#[test]
+fn non_indexed_multi_draws_draw_every_entry_as_given() {
+    let Some((device, queue)) = device() else {
+        eprintln!("no GPU with indirect first instance: skipped");
+        return;
+    };
+    for count in [3, 16, 90] {
+        // Non-indexed: vertex count, instances, first vertex (a quad's
+        // first), first instance.
+        let entries: Vec<[u32; 5]> = entries(count)
+            .into_iter()
+            .map(|[n, inst, _, _, first_instance]| {
+                let quad = (first_instance * 3) % 61;
+                [n, inst, quad * 6, 0, first_instance]
+            })
+            .collect();
+        let row = draw(&device, &queue, &entries, None, &[(0, count)]);
+        assert_eq!(row, expected(&entries, true), "{count} non-indexed draws");
+    }
+}

third_party/wgpu-hal/Cargo.toml

@@ -449,6 +449,8 @@ features = [
     "MTLAccelerationStructureTypes",
     "MTLAccelerationStructureCommandEncoder",
     "MTLResidencySet",
+    "MTLArgumentEncoder",
+    "MTLIndirectCommandBuffer",
 ]
 optional = true
 default-features = false

third_party/wgpu-hal/README.md

@@ -1,3 +1,24 @@
+# Horizon's copy of wgpu-hal 30.0.1
+
+This is wgpu-hal 30.0.1 from crates.io, patched into the build with `[patch.crates-io]` in the
+workspace `Cargo.toml`. Changes from the release:
+
+1. **Vulkan descriptor pool** (`src/vulkan/descriptor.rs`): the upstream fix
+   [gfx-rs/wgpu#10124](https://github.com/gfx-rs/wgpu/pull/10124) for Intel's Windows driver
+   crashing at device creation (#746).
+2. **Metal multi-draw through indirect command buffers** (`src/metal/icb.rs`, plus hooks in
+   `src/metal/{adapter,command,device,mod}.rs` marked "Horizon patch", and the
+   `MTLArgumentEncoder` and `MTLIndirectCommandBuffer` features of `objc2-metal` in
+   `Cargo.toml`). The release encodes one draw call per entry of an indirect multi-draw; on Apple
+   GPUs (M1 and later) a multi-draw of 16 or more entries under a pipeline whose layout binds no
+   texture or sampler directly is now one `executeCommandsInBuffer`, with the commands written on
+   the GPU from the indirect arguments before the pass. `WGPU_METAL_ICB=0` turns it off. Upstream
+   proposes the same in [gfx-rs/wgpu#9640](https://github.com/gfx-rs/wgpu/pull/9640), with a
+   wgpu-core hook where this copy relies on wgpu-core 30 encoding each render pass's pre-pass
+   right after it (see the comment at the top of `icb.rs`).
+
+Drop each change once a wgpu release has it, and the folder once none is left.
+
 # `wgpu_hal`: a cross-platform unsafe graphics abstraction
 
 This crate defines a set of traits abstracting over modern graphics APIs,

third_party/wgpu-hal/src/metal/adapter.rs

@@ -110,6 +110,11 @@ impl crate::Adapter for super::Adapter {
                     features,
                     counters: Default::default(),
                     limits: limits.clone(),
+                    // Horizon patch: multi-draws through indirect command buffers.
+                    icb: super::icb::Icbs::new(
+                        &self.shared.device,
+                        self.shared.private_caps.msl_version,
+                    ),
                 },
                 queue: super::Queue {
                     shared: Arc::new(QueueShared {

third_party/wgpu-hal/src/metal/command.rs

@@ -43,6 +43,8 @@ impl Default for super::CommandState {
             vertex_buffer_size_map: Default::default(),
             immediates: Vec::new(),
             pending_timer_queries: Vec::new(),
+            icb_pipeline: false,
+            icb_index_used: None,
         }
     }
 }
@@ -157,6 +159,63 @@ impl super::CommandEncoder {
         self.raw_cmd_buf.as_deref()
     }
 
+    /// Horizon patch: encodes a multi-draw as one execution of an indirect
+    /// command buffer range, whose commands are generated before the pass
+    /// (see `icb`). Returns false, encoding nothing, where it doesn't apply.
+    unsafe fn draw_through_icb(
+        &mut self,
+        buffer: &super::Buffer,
+        offset: wgt::BufferAddress,
+        draw_count: u32,
+        indexed: bool,
+    ) -> bool {
+        if draw_count < super::icb::MIN_DRAWS || !self.state.icb_pipeline {
+            return false;
+        }
+        let Some(icbs) = &self.icb else {
+            return false;
+        };
+        let index = match (indexed, &self.state.index) {
+            (false, _) => None,
+            (true, Some(index)) => Some((
+                // SAFETY: the index buffer outlives the command buffer.
+                unsafe { Retained::retain(index.buffer_ptr.as_ptr()) }.unwrap(),
+                index.offset,
+                index.raw_type,
+            )),
+            (true, None) => return false,
+        };
+        let index_buffer = index.as_ref().map(|(buffer, _, _)| Retained::as_ptr(buffer));
+        let Some(slot) = super::icb::lower(
+            icbs,
+            &mut self.icb_leases,
+            &mut self.icb_pending,
+            buffer,
+            offset,
+            draw_count,
+            index,
+            self.state.raw_primitive_type,
+        ) else {
+            return false;
+        };
+        // The index buffer, reached only through the commands, is declared
+        // once per pass.
+        let declare = index_buffer.filter(|&ptr| self.state.icb_index_used != Some(ptr));
+        if declare.is_some() {
+            self.state.icb_index_used = declare;
+        }
+        let encoder = self.state.render.as_ref().unwrap();
+        unsafe {
+            super::icb::execute(
+                encoder,
+                &slot,
+                draw_count,
+                declare.map(|ptr| &*ptr),
+            )
+        };
+        true
+    }
+
     fn enter_blit(&mut self) -> Retained<ProtocolObject<dyn MTLBlitCommandEncoder>> {
         if self.state.blit.is_none() {
             self.leave_acceleration_structure_builder();
@@ -520,6 +579,13 @@ impl crate::CommandEncoder for super::CommandEncoder {
             cmd_buf_ref.to_owned()
         });
 
+        // Horizon patch: a command buffer that follows one with ICB
+        // executions is its render pass's pre-pass (see `icb`).
+        debug_assert!(
+            self.icb_carry.is_empty() || label.is_none_or(|l| l.contains("Pre Pass")),
+            "ICB generation would land in {label:?}, not in the pass's pre-pass"
+        );
+
         // Queries should either be closed out, or cleared in `discard_encoding`.
         // This assertion is here mainly to facilitate comparison with the other
         // backends, which clear in `begin_encoding` rather than `discard_encoding`.
@@ -542,6 +608,9 @@ impl crate::CommandEncoder for super::CommandEncoder {
             encoder.endEncoding();
         }
         self.state.pending_timer_queries.clear();
+        self.icb_leases.clear();
+        self.icb_pending.clear();
+        self.icb_carry.clear();
         let had_command_buffer = self.raw_cmd_buf.is_some();
         // Clear the Option first so the underlying `metal::CommandBuffer` is
         // dropped before we update the counter.
@@ -566,9 +635,20 @@ impl crate::CommandEncoder for super::CommandEncoder {
         debug_assert!(self.state.compute.is_none());
         debug_assert!(self.state.pending_timer_queries.is_empty());
 
+        let raw = self.raw_cmd_buf.take().unwrap();
+        // Horizon patch: the previous command buffer held render passes with
+        // multi-draws lowered to ICBs: this one is their pre-pass (see `icb`),
+        // so their commands are generated here, after everything else in it.
+        if let (Some(icbs), false) = (&self.icb, self.icb_carry.is_empty()) {
+            super::icb::encode(icbs, &raw, &self.icb_carry);
+        }
+        self.icb_carry.clear();
+        core::mem::swap(&mut self.icb_carry, &mut self.icb_pending);
+
         Ok(super::CommandBuffer {
-            raw: self.raw_cmd_buf.take().unwrap(),
+            raw,
             queue_shared: Arc::clone(&self.queue_shared),
+            _icb_leases: core::mem::take(&mut self.icb_leases),
         })
     }
 
@@ -576,7 +656,7 @@ impl crate::CommandEncoder for super::CommandEncoder {
     where
         I: Iterator<Item = super::CommandBuffer>,
     {
-        //do nothing
+        self.icb_carry.clear();
     }
 
     unsafe fn transition_buffers<'a, T>(&mut self, _barriers: T)
@@ -896,6 +976,7 @@ impl crate::CommandEncoder for super::CommandEncoder {
     ) -> Result<(), crate::DeviceError> {
         self.begin_pass();
         self.state.index = None;
+        self.state.icb_index_used = None;
 
         assert!(self.state.blit.is_none());
         assert!(self.state.compute.is_none());
@@ -1274,6 +1355,7 @@ impl crate::CommandEncoder for super::CommandEncoder {
 
     unsafe fn set_render_pipeline(&mut self, pipeline: &super::RenderPipeline) {
         self.state.raw_primitive_type = pipeline.raw_primitive_type;
+        self.state.icb_pipeline = pipeline.icb;
         match pipeline.vs_info {
             Some(ref info) => self.state.stage_infos.vs.assign_from(info),
             None => self.state.stage_infos.vs.clear(),
@@ -1592,6 +1674,9 @@ impl crate::CommandEncoder for super::CommandEncoder {
         mut offset: wgt::BufferAddress,
         draw_count: u32,
     ) {
+        if unsafe { self.draw_through_icb(buffer, offset, draw_count, false) } {
+            return;
+        }
         let encoder = self.state.render.as_ref().unwrap();
         for _ in 0..draw_count {
             unsafe {
@@ -1611,6 +1696,9 @@ impl crate::CommandEncoder for super::CommandEncoder {
         mut offset: wgt::BufferAddress,
         draw_count: u32,
     ) {
+        if unsafe { self.draw_through_icb(buffer, offset, draw_count, true) } {
+            return;
+        }
         let encoder = self.state.render.as_ref().unwrap();
         let index = self.state.index.as_ref().unwrap();
         for _ in 0..draw_count {

third_party/wgpu-hal/src/metal/device.rs

@@ -439,14 +439,23 @@ impl super::Device {
     ) -> super::Device {
         let capabilities_query = super::CapabilitiesQuery::new(&raw);
         let shared = super::AdapterShared::new(raw, &capabilities_query);
+        let icb = super::icb::Icbs::new(&shared.device, shared.private_caps.msl_version);
         super::Device {
             shared: Arc::new(shared),
             features,
             counters: Default::default(),
             limits: limits.clone(),
+            icb,
         }
     }
 
+    /// Horizon patch: whether multi-draws of pipelines whose layout binds no
+    /// texture or sampler directly to the vertex or fragment stage go
+    /// through indirect command buffers (see `icb`).
+    pub fn multi_draws_use_icbs(&self) -> bool {
+        self.icb.as_ref().is_some_and(|icbs| icbs.ready())
+    }
+
     pub unsafe fn buffer_from_raw(
         raw: Retained<ProtocolObject<dyn MTLBuffer>>,
         size: wgt::BufferAddress,
@@ -776,6 +785,10 @@ impl crate::Device for super::Device {
             state: super::CommandState::default(),
             temp: super::Temp::default(),
             counters: Arc::clone(&self.counters),
+            icb: self.icb.clone(),
+            icb_leases: Vec::new(),
+            icb_pending: Vec::new(),
+            icb_carry: Vec::new(),
         })
     }
 
@@ -949,6 +962,10 @@ impl crate::Device for super::Device {
             }
         }
 
+        let icb_compatible = [&stage_data.vs, &stage_data.fs]
+            .iter()
+            .all(|info| info.counters.textures == 0 && info.counters.samplers == 0);
+
         let immediates_infos = stage_data.map_ref(|info| {
             info.pc_buffer.map(|buffer_index| super::ImmediateDataInfo {
                 count: info.pc_limit,
@@ -974,6 +991,7 @@ impl crate::Device for super::Device {
             total_immediates: desc.immediate_size,
             per_stage_map,
             binding_array_length_map,
+            icb_compatible,
         })
     }
 
@@ -1743,13 +1761,40 @@ impl crate::Device for super::Device {
                 };
             }
 
+            // Horizon patch: pipelines that can run from an indirect command
+            // buffer get its multi-draws lowered to one (see `icb`). Metal
+            // rejects some (fragment functions with direct textures, which
+            // the layout check rules out); those are created again without.
+            let mut icb = desc.multiview_mask.is_none()
+                && desc.layout.icb_compatible
+                && matches!(descriptor, MetalGenericRenderPipelineDescriptor::Standard(_))
+                && self.icb.as_ref().is_some_and(|icbs| icbs.ready());
+            let icb_raw = match descriptor {
+                MetalGenericRenderPipelineDescriptor::Standard(ref d) if icb => {
+                    d.setSupportIndirectCommandBuffers(true);
+                    let raw = self.shared.device.newRenderPipelineStateWithDescriptor_error(d);
+                    if let Err(ref e) = raw {
+                        log::debug!(
+                            "Metal rejected an ICB-capable {:?}, so its multi-draws use the \
+                             per-draw loop: {e:?}",
+                            desc.label
+                        );
+                        d.setSupportIndirectCommandBuffers(false);
+                        icb = false;
+                    }
+                    raw.ok()
+                }
+                _ => None,
+            };
+
             // Create the pipeline from descriptor
-            let raw = match descriptor {
-                MetalGenericRenderPipelineDescriptor::Standard(d) => self
+            let raw = match (icb_raw, descriptor) {
+                (Some(raw), _) => Ok(raw),
+                (None, MetalGenericRenderPipelineDescriptor::Standard(d)) => self
                     .shared
                     .device
                     .newRenderPipelineStateWithDescriptor_error(&d),
-                MetalGenericRenderPipelineDescriptor::Mesh(d) => self
+                (None, MetalGenericRenderPipelineDescriptor::Mesh(d)) => self
                     .shared
                     .device
                     .newRenderPipelineStateWithMeshDescriptor_options_reflection_error(
@@ -1787,6 +1832,7 @@ impl crate::Device for super::Device {
                     None
                 },
                 depth_stencil,
+                icb,
             })
         })
     }

third_party/wgpu-hal/src/metal/icb.rs

@@ -0,0 +1,508 @@
+//! Multi-draw indirect through Metal indirect command buffers (ICBs).
+//!
+//! Horizon patch (not in wgpu-hal 30.0.1). Metal has no multi-draw, so the
+//! backend used to encode one `drawIndexedPrimitives:...indirectBuffer:` per
+//! entry of a multi-draw: thousands of driver calls a frame for a terrain
+//! pass. Here a multi-draw of [`MIN_DRAWS`] or more under a pipeline that can
+//! run in an ICB is encoded as one `executeCommandsInBuffer:withRange:`, and
+//! a compute kernel writes the range's commands from the indirect arguments
+//! on the GPU, so the draw count costs the CPU nothing.
+//!
+//! The kernels must run after whatever wrote the arguments and before the
+//! render pass. wgpu-core 30 records every render pass in its own hal command
+//! buffer and then, on the same hal encoder, the pass's "Pre Pass" buffer
+//! (barriers, query resets, indirect validation), which it submits before the
+//! pass. The generation for a pass is therefore encoded at the end of the
+//! buffer that follows it (see `CommandEncoder::end_encoding`): after the
+//! indirect validation, before the pass. This is the same place as upstream's
+//! proposal (gfx-rs/wgpu#9640), which gets an explicit wgpu-core hook for it;
+//! drop this file once a wgpu release lowers multi-draws to ICBs.
+//!
+//! Commands inherit the pipeline and every buffer binding from the render
+//! encoder; only the index buffer travels in the command, so it is declared
+//! with `useResource`. Metal accepts such pipelines only when no stage binds a
+//! texture or sampler directly (argument buffers are fine), so only pipeline
+//! layouts without direct textures and samplers in the vertex and fragment
+//! stages get ICB-capable pipelines; everything else keeps the loop.
+
+use alloc::{string::ToString as _, sync::Arc, vec::Vec};
+use std::sync::OnceLock;
+
+use objc2::{rc::Retained, runtime::ProtocolObject};
+use objc2_foundation::{NSRange, NSString};
+use objc2_metal::{
+    MTLArgumentEncoder, MTLBuffer, MTLCommandBuffer, MTLCommandEncoder, MTLCompileOptions,
+    MTLComputeCommandEncoder, MTLComputePipelineState, MTLDevice, MTLDispatchType, MTLFunction,
+    MTLGPUFamily, MTLIndexType, MTLIndirectCommandBuffer, MTLIndirectCommandBufferDescriptor,
+    MTLIndirectCommandType, MTLLanguageVersion, MTLLibrary, MTLPrimitiveType,
+    MTLRenderCommandEncoder, MTLRenderStages, MTLResource, MTLResourceOptions, MTLResourceUsage,
+    MTLSize,
+};
+use parking_lot::Mutex;
+
+/// Multi-draws with fewer draws keep the per-draw loop: below this the
+/// kernel dispatch costs about what the draws do.
+pub(super) const MIN_DRAWS: u32 = 16;
+
+/// Commands per ICB; a multi-draw larger than this gets an ICB of its own.
+const CHUNK: u32 = 1 << 14;
+
+/// ICBs kept for reuse once the GPU is done with them.
+const POOL: usize = 16;
+
+const SOURCE: &str = r#"
+#include <metal_stdlib>
+using namespace metal;
+
+struct DrawArgs { uint vertex_count; uint instance_count; uint first_vertex; uint first_instance; };
+struct DrawIndexedArgs { uint index_count; uint instance_count; uint first_index; int base_vertex; uint first_instance; };
+struct Icb { command_buffer commands [[id(0)]]; };
+// first: first command; index_base: index buffer elements before the bound offset.
+struct Params { uint count; uint first; uint primitive; uint index_base; };
+
+static primitive_type primitive(uint p) {
+    switch (p) {
+        case 0: return primitive_type::point;
+        case 1: return primitive_type::line;
+        case 2: return primitive_type::line_strip;
+        case 4: return primitive_type::triangle_strip;
+        default: return primitive_type::triangle;
+    }
+}
+
+kernel void draws(const device DrawArgs* args [[buffer(0)]], device Icb& icb [[buffer(1)]],
+                  constant Params& p [[buffer(2)]], uint i [[thread_position_in_grid]]) {
+    if (i >= p.count) return;
+    render_command cmd(icb.commands, p.first + i);
+    DrawArgs a = args[i];
+    if (a.vertex_count == 0 || a.instance_count == 0) { cmd.reset(); return; }
+    cmd.draw_primitives(primitive(p.primitive), a.first_vertex, a.vertex_count, a.instance_count,
+                        a.first_instance);
+}
+
+template <typename I>
+static void indexed(const device DrawIndexedArgs* args, device Icb& icb, constant Params& p,
+                    const device I* indices, uint i) {
+    if (i >= p.count) return;
+    render_command cmd(icb.commands, p.first + i);
+    DrawIndexedArgs a = args[i];
+    if (a.index_count == 0 || a.instance_count == 0) { cmd.reset(); return; }
+    cmd.draw_indexed_primitives(primitive(p.primitive), a.index_count,
+                                indices + p.index_base + a.first_index, a.instance_count,
+                                uint(a.base_vertex), a.first_instance);
+}
+
+kernel void draws_u16(const device DrawIndexedArgs* args [[buffer(0)]], device Icb& icb [[buffer(1)]],
+                      constant Params& p [[buffer(2)]], const device ushort* indices [[buffer(3)]],
+                      uint i [[thread_position_in_grid]]) {
+    indexed(args, icb, p, indices, i);
+}
+
+kernel void draws_u32(const device DrawIndexedArgs* args [[buffer(0)]], device Icb& icb [[buffer(1)]],
+                      constant Params& p [[buffer(2)]], const device uint* indices [[buffer(3)]],
+                      uint i [[thread_position_in_grid]]) {
+    indexed(args, icb, p, indices, i);
+}
+"#;
+
+#[repr(C)]
+#[derive(Clone, Copy)]
+struct Params {
+    count: u32,
+    first: u32,
+    primitive: u32,
+    index_base: u32,
+}
+
+type Raw<T> = Retained<ProtocolObject<T>>;
+
+struct Kernels {
+    draws: Raw<dyn MTLComputePipelineState>,
+    draws_u16: Raw<dyn MTLComputePipelineState>,
+    draws_u32: Raw<dyn MTLComputePipelineState>,
+    /// Encodes an ICB into the kernels' `Icb` argument (buffer 1); locked,
+    /// as encoders on several threads may lease ICBs.
+    arguments: Mutex<Raw<dyn MTLArgumentEncoder>>,
+}
+
+/// The device's ICB support: the generation kernels and a pool of ICBs.
+pub(super) struct Icbs {
+    device: Raw<dyn MTLDevice>,
+    msl_version: MTLLanguageVersion,
+    kernels: OnceLock<Option<Kernels>>,
+    pool: Mutex<Vec<Icb>>,
+}
+
+impl core::fmt::Debug for Icbs {
+    fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result {
+        f.debug_struct("Icbs").finish_non_exhaustive()
+    }
+}
+
+unsafe impl Send for Icbs {}
+unsafe impl Sync for Icbs {}
+
+impl Icbs {
+    /// ICB support for `device`, where it is known to work: Apple GPUs from
+    /// the M1 and A14 on (not the paravirtual device of macOS VMs).
+    /// `WGPU_METAL_ICB=0` turns it off.
+    pub(super) fn new(
+        device: &Raw<dyn MTLDevice>,
+        msl_version: MTLLanguageVersion,
+    ) -> Option<Arc<Self>> {
+        if std::env::var_os("WGPU_METAL_ICB").is_some_and(|v| v == "0") {
+            return None;
+        }
+        let supported = device.supportsFamily(MTLGPUFamily::Apple7)
+            && !device.name().to_string().contains("Paravirtual");
+        supported.then(|| {
+            Arc::new(Self {
+                device: device.clone(),
+                msl_version,
+                kernels: OnceLock::new(),
+                pool: Mutex::new(Vec::new()),
+            })
+        })
+    }
+
+    /// The generation kernels, compiled on first use; `None` if they don't
+    /// compile (multi-draws then keep the loop).
+    fn kernels(&self) -> Option<&Kernels> {
+        self.kernels
+            .get_or_init(|| {
+                let kernels = self.compile();
+                if let Err(ref e) = kernels {
+                    log::warn!("Metal ICB kernels: {e}; multi-draws use the per-draw loop");
+                }
+                kernels.ok()
+            })
+            .as_ref()
+    }
+
+    /// Whether multi-draws can be lowered (compiling the kernels now if not
+    /// yet: called when an ICB-capable pipeline is created).
+    pub(super) fn ready(&self) -> bool {
+        self.kernels().is_some()
+    }
+
+    fn compile(&self) -> Result<Kernels, alloc::string::String> {
+        use alloc::format;
+        let options = MTLCompileOptions::new();
+        options.setLanguageVersion(self.msl_version.max(MTLLanguageVersion::Version2_1));
+        let library = self
+            .device
+            .newLibraryWithSource_options_error(&NSString::from_str(SOURCE), Some(&options))
+            .map_err(|e| format!("{e:?}"))?;
+        let pipeline = |name: &str| {
+            let function = library
+                .newFunctionWithName(&NSString::from_str(name))
+                .ok_or_else(|| format!("no function {name}"))?;
+            let state = self
+                .device
+                .newComputePipelineStateWithFunction_error(&function)
+                .map_err(|e| format!("{e:?}"))?;
+            Ok::<_, alloc::string::String>((function, state))
+        };
+        let (draws_fn, draws) = pipeline("draws")?;
+        let (_, draws_u16) = pipeline("draws_u16")?;
+        let (_, draws_u32) = pipeline("draws_u32")?;
+        let arguments = unsafe { draws_fn.newArgumentEncoderWithBufferIndex(1) };
+        Ok(Kernels {
+            draws,
+            draws_u16,
+            draws_u32,
+            arguments: Mutex::new(arguments),
+        })
+    }
+
+    /// An ICB for at least `count` commands, from the pool if one fits.
+    fn lease(self: &Arc<Self>, count: u32) -> Option<Lease> {
+        let kernels = self.kernels()?;
+        let reused = {
+            let mut pool = self.pool.lock();
+            pool.iter()
+                .position(|icb| icb.capacity >= count)
+                .map(|i| pool.swap_remove(i))
+        };
+        let icb = match reused {
+            Some(icb) => icb,
+            None => {
+                let capacity = count.max(CHUNK).next_power_of_two();
+                let descriptor = MTLIndirectCommandBufferDescriptor::new();
+                descriptor.setCommandTypes(
+                    MTLIndirectCommandType::Draw | MTLIndirectCommandType::DrawIndexed,
+                );
+                descriptor.setInheritPipelineState(true);
+                descriptor.setInheritBuffers(true);
+                // Inheriting buffers, these are unused; some Apple GPUs
+                // (A12, A13) fault unless they are at the maximum.
+                descriptor.setMaxVertexBufferBindCount(31);
+                descriptor.setMaxFragmentBufferBindCount(31);
+                let raw = unsafe {
+                    self.device
+                        .newIndirectCommandBufferWithDescriptor_maxCommandCount_options(
+                            &descriptor,
+                            capacity as usize,
+                            MTLResourceOptions::StorageModePrivate,
+                        )
+                }?;
+                raw.setLabel(Some(&NSString::from_str("wgpu multi-draw ICB")));
+                let arguments = kernels.arguments.lock();
+                let argument = self.device.newBufferWithLength_options(
+                    arguments.encodedLength(),
+                    MTLResourceOptions::StorageModeShared,
+                )?;
+                unsafe {
+                    arguments.setArgumentBuffer_offset(Some(&argument), 0);
+                    arguments.setIndirectCommandBuffer_atIndex(Some(&raw), 0);
+                }
+                drop(arguments);
+                Icb {
+                    raw,
+                    argument,
+                    capacity,
+                }
+            }
+        };
+        Some(Lease {
+            icb: Some(icb),
+            used: 0,
+            pool: Arc::clone(self),
+        })
+    }
+}
+
+struct Icb {
+    raw: Raw<dyn MTLIndirectCommandBuffer>,
+    /// The kernels' `Icb` argument, pointing at `raw`.
+    argument: Raw<dyn MTLBuffer>,
+    capacity: u32,
+}
+
+/// An ICB in use by one command buffer; back to the pool when that command
+/// buffer is dropped, which wgpu-core does only once the GPU is done with it.
+pub(super) struct Lease {
+    icb: Option<Icb>,
+    used: u32,
+    pool: Arc<Icbs>,
+}
+
+impl core::fmt::Debug for Lease {
+    fn fmt(&self, f: &mut core::fmt::Formatter<'_>) -> core::fmt::Result {
+        f.debug_struct("Lease").field("used", &self.used).finish()
+    }
+}
+
+impl Drop for Lease {
+    fn drop(&mut self) {
+        if let Some(icb) = self.icb.take() {
+            let mut pool = self.pool.pool.lock();
+            if pool.len() < POOL {
+                pool.push(icb);
+            }
+        }
+    }
+}
+
+unsafe impl Send for Lease {}
+unsafe impl Sync for Lease {}
+
+/// One multi-draw's commands to write before its pass runs.
+pub(super) struct Generation {
+    icb: Raw<dyn MTLIndirectCommandBuffer>,
+    argument: Raw<dyn MTLBuffer>,
+    first: u32,
+    count: u32,
+    args: Raw<dyn MTLBuffer>,
+    args_offset: u64,
+    /// Index buffer, its offset and type, for indexed draws.
+    index: Option<(Raw<dyn MTLBuffer>, u64, MTLIndexType)>,
+    primitive: MTLPrimitiveType,
+}
+
+unsafe impl Send for Generation {}
+unsafe impl Sync for Generation {}
+
+/// Where a multi-draw's commands go.
+pub(super) struct Slot {
+    icb: Raw<dyn MTLIndirectCommandBuffer>,
+    first: u32,
+}
+
+/// Most entries skipped between two multi-draws whose commands one dispatch
+/// still writes together (it writes the skipped entries' commands too, which
+/// never run).
+const MAX_GAP: u64 = 64;
+
+/// Reserves the commands of a multi-draw of `count` entries at `args_offset`
+/// in `args`, and queues their generation: within the previous multi-draw's
+/// dispatch when that one read the same buffers up to a few entries before
+/// (consecutive runs of one indirect list, as a frame's terrain draws are),
+/// else as a dispatch of its own.
+#[allow(clippy::too_many_arguments)]
+pub(super) fn lower(
+    icbs: &Arc<Icbs>,
+    leases: &mut Vec<Lease>,
+    pending: &mut Vec<Generation>,
+    args: &super::Buffer,
+    args_offset: u64,
+    count: u32,
+    index: Option<(Raw<dyn MTLBuffer>, u64, MTLIndexType)>,
+    primitive: MTLPrimitiveType,
+) -> Option<Slot> {
+    let stride = if index.is_some() { 20 } else { 16 };
+    if let (Some(last), Some(lease)) = (pending.last_mut(), leases.last_mut()) {
+        let end = last.args_offset + u64::from(last.count) * stride;
+        let gap = args_offset
+            .checked_sub(end)
+            .map(|bytes| (bytes / stride, bytes % stride));
+        let same_index = match (&last.index, &index) {
+            (None, None) => true,
+            (Some((a, a_offset, a_type)), Some((b, b_offset, b_type))) => {
+                core::ptr::eq(&**a, &**b) && a_offset == b_offset && a_type == b_type
+            }
+            _ => false,
+        };
+        if let (Some((gap, 0)), Some(icb), true) = (gap, &lease.icb, same_index) {
+            let first = lease.used + gap as u32;
+            if gap <= MAX_GAP
+                && core::ptr::eq(&*last.icb, &*icb.raw)
+                && core::ptr::eq(&*last.args, &*args.raw)
+                && last.primitive == primitive
+                && last.first + last.count == lease.used
+                && u64::from(first) + u64::from(count) <= u64::from(icb.capacity)
+            {
+                lease.used = first + count;
+                last.count = lease.used - last.first;
+                return Some(Slot {
+                    icb: icb.raw.clone(),
+                    first,
+                });
+            }
+        }
+    }
+    let fits = leases.last().is_some_and(|lease| {
+        lease
+            .icb
+            .as_ref()
+            .is_some_and(|icb| icb.capacity - lease.used >= count)
+    });
+    if !fits {
+        leases.push(icbs.lease(count)?);
+    }
+    let lease = leases.last_mut()?;
+    let icb = lease.icb.as_ref()?;
+    let first = lease.used;
+    lease.used += count;
+    pending.push(Generation {
+        icb: icb.raw.clone(),
+        argument: icb.argument.clone(),
+        first,
+        count,
+        args: args.raw.clone(),
+        args_offset,
+        index,
+        primitive,
+    });
+    Some(Slot {
+        icb: icb.raw.clone(),
+        first,
+    })
+}
+
+/// Encodes `slot`'s execution into the render pass: `count` commands.
+///
+/// # Safety
+///
+/// `index`, if any, must stay alive until the command buffer completes.
+pub(super) unsafe fn execute(
+    encoder: &ProtocolObject<dyn MTLRenderCommandEncoder>,
+    slot: &Slot,
+    count: u32,
+    index: Option<&ProtocolObject<dyn MTLBuffer>>,
+) {
+    if let Some(index) = index {
+        encoder.useResource_usage_stages(
+            ProtocolObject::<dyn MTLResource>::from_ref(index),
+            MTLResourceUsage::Read,
+            MTLRenderStages::Vertex,
+        );
+    }
+    unsafe {
+        encoder.executeCommandsInBuffer_withRange(
+            &slot.icb,
+            NSRange::new(slot.first as usize, count as usize),
+        )
+    };
+}
+
+/// Encodes the kernels that write `generations`' commands into `cmd_buf`.
+pub(super) fn encode(
+    icbs: &Icbs,
+    cmd_buf: &ProtocolObject<dyn MTLCommandBuffer>,
+    generations: &[Generation],
+) {
+    let Some(kernels) = icbs.kernels() else {
+        return;
+    };
+    let Some(encoder) = cmd_buf.computeCommandEncoderWithDispatchType(MTLDispatchType::Concurrent)
+    else {
+        return;
+    };
+    encoder.setLabel(Some(&NSString::from_str("wgpu multi-draw ICB generation")));
+    for g in generations {
+        let (pipeline, index_base) = match &g.index {
+            None => (&kernels.draws, 0),
+            Some((buffer, offset, index_type)) => {
+                // Bind at a 4-byte aligned offset; the rest in elements.
+                let aligned = offset & !3;
+                let stride = if *index_type == MTLIndexType::UInt16 {
+                    2
+                } else {
+                    4
+                };
+                unsafe { encoder.setBuffer_offset_atIndex(Some(buffer), aligned as usize, 3) };
+                let pipeline = if *index_type == MTLIndexType::UInt16 {
+                    &kernels.draws_u16
+                } else {
+                    &kernels.draws_u32
+                };
+                (pipeline, ((offset - aligned) / stride) as u32)
+            }
+        };
+        let params = Params {
+            count: g.count,
+            first: g.first,
+            primitive: g.primitive.0 as u32,
+            index_base,
+        };
+        encoder.setComputePipelineState(pipeline);
+        unsafe {
+            encoder.setBuffer_offset_atIndex(Some(&g.args), g.args_offset as usize, 0);
+            encoder.setBuffer_offset_atIndex(Some(&g.argument), 0, 1);
+            encoder.setBytes_length_atIndex(
+                core::ptr::NonNull::from(&params).cast(),
+                size_of::<Params>(),
+                2,
+            );
+        }
+        encoder.useResource_usage(
+            ProtocolObject::<dyn MTLResource>::from_ref(&*g.icb),
+            MTLResourceUsage::Write,
+        );
+        let width = pipeline.threadExecutionWidth().max(1);
+        encoder.dispatchThreads_threadsPerThreadgroup(
+            MTLSize {
+                width: g.count as usize,
+                height: 1,
+                depth: 1,
+            },
+            MTLSize {
+                width,
+                height: 1,
+                depth: 1,
+            },
+        );
+    }
+    encoder.endEncoding();
+}

third_party/wgpu-hal/src/metal/mod.rs

@@ -24,6 +24,8 @@ mod adapter;
 mod command;
 mod conv;
 mod device;
+// Horizon patch: multi-draws through indirect command buffers.
+mod icb;
 mod library_from_metallib;
 mod surface;
 mod time;
@@ -637,6 +639,9 @@ pub struct Device {
     features: wgt::Features,
     counters: Arc<wgt::HalCounters>,
     limits: wgt::Limits,
+    /// Horizon patch: multi-draws through indirect command buffers, where
+    /// supported.
+    icb: Option<Arc<icb::Icbs>>,
 }
 
 #[derive(Debug)]
@@ -1029,6 +1034,10 @@ pub struct PipelineLayout {
     total_immediates: u32,
     per_stage_map: MultiStageResources,
     binding_array_length_map: FastHashMap<naga::ResourceBinding, u32>,
+    /// Horizon patch: no texture or sampler bound directly to the vertex or
+    /// fragment stage, which Metal requires of pipelines run from an indirect
+    /// command buffer.
+    icb_compatible: bool,
 }
 
 impl crate::DynPipelineLayout for PipelineLayout {}
@@ -1214,6 +1223,9 @@ pub struct RenderPipeline {
         Retained<ProtocolObject<dyn MTLDepthStencilState>>,
         wgt::DepthBiasState,
     )>,
+    /// Horizon patch: created with `supportIndirectCommandBuffers`, so its
+    /// multi-draws go through indirect command buffers.
+    icb: bool,
 }
 
 #[cfg(send_sync)]
@@ -1346,6 +1358,13 @@ struct CommandState {
 
     /// Timer query that should be executed when the next pass starts.
     pending_timer_queries: Vec<(QuerySet, u32)>,
+
+    /// Horizon patch: the current render pipeline can run from an indirect
+    /// command buffer.
+    icb_pipeline: bool,
+    /// Horizon patch: the index buffer the render pass last declared for
+    /// commands of an indirect command buffer.
+    icb_index_used: Option<*const ProtocolObject<dyn MTLBuffer>>,
 }
 
 // Any state in this struct that may be dirty after an abandoned encoding must
@@ -1357,6 +1376,16 @@ pub struct CommandEncoder {
     state: CommandState,
     temp: Temp,
     counters: Arc<wgt::HalCounters>,
+    /// Horizon patch: multi-draws through indirect command buffers.
+    icb: Option<Arc<icb::Icbs>>,
+    /// The ICBs the open command buffer's passes execute.
+    icb_leases: Vec<icb::Lease>,
+    /// Commands the open command buffer's passes execute, to generate in the
+    /// buffer that follows it (see `icb`).
+    icb_pending: Vec<icb::Generation>,
+    /// Commands of the previous command buffer, generated at the end of the
+    /// open one.
+    icb_carry: Vec<icb::Generation>,
 }
 
 impl fmt::Debug for CommandEncoder {
@@ -1374,6 +1403,8 @@ unsafe impl Sync for CommandEncoder {}
 pub struct CommandBuffer {
     raw: Retained<ProtocolObject<dyn MTLCommandBuffer>>,
     queue_shared: Arc<QueueShared>,
+    /// Horizon patch: kept until the GPU is done with this command buffer.
+    _icb_leases: Vec<icb::Lease>,
 }
 
 impl crate::DynCommandBuffer for CommandBuffer {}