Vulkan, DX12: phase 1's solid draws made and counted on the GPU

Where multi-draws take their count from a buffer (MULTI_DRAW_INDIRECT_COUNT:
Vulkan 1.2+, DX12), the render thread no longer finds each phase 1 section's
facing runs nor uploads a draw per run: it lists one record per section, and
gpu_lists.wgsl's sides_draws writes one indirect draw per run the CPU would have
pushed (same count, base vertex and instance) into a call per pass and arena
page, drawn with multi_draw_indexed_indirect_count. Each call has room for
the most runs its sections' parts can have (Sides::part_runs). Same exclusions
as the flat path; HORIZON_GPU_LISTS=0 turns it off.

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

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

@@ -792,7 +792,8 @@ pub struct Renderer {
     flat: Option<flat::Flat>,
     /// `flat` draws this frame's sections (off: the per-section draws).
     flat_on: bool,
-    /// Phase 1's solid lists made on the GPU, in flat frames (with `flat`).
+    /// Phase 1's solid lists made on the GPU: in flat frames (indices), or as
+    /// draws counted on the GPU (`MULTI_DRAW_INDIRECT_COUNT`: Vulkan, DX12).
     gpu_lists: Option<gpu_lists::GpuLists>,
     /// Occlusion culling of solid section draws, where the GPU's instance
     /// counts reach a multi-draw's entries (Vulkan, Metal).
@@ -1689,6 +1690,7 @@ async fn open_backend(
         | oit::features(&adapter)
         | (adapter.features()
             & (wgpu::Features::INDIRECT_FIRST_INSTANCE
+                | wgpu::Features::MULTI_DRAW_INDIRECT_COUNT
                 | wgpu::Features::MAPPABLE_PRIMARY_BUFFERS
                 | wgpu::Features::TEXTURE_FORMAT_16BIT_NORM
                 | wgpu::Features::TEXTURE_ADAPTER_SPECIFIC_FORMAT_FEATURES
@@ -3514,7 +3516,18 @@ impl Renderer {
                 .get_downlevel_capabilities()
                 .flags
                 .contains(wgpu::DownlevelFlags::INDIRECT_EXECUTION);
-        let gpu_lists = flat.as_ref().map(|_| gpu_lists::GpuLists::new(&device));
+        let gpu_lists = if flat.is_some() {
+            Some(gpu_lists::Out::Indices)
+        } else if multi_draw
+            && device
+                .features()
+                .contains(wgpu::Features::MULTI_DRAW_INDIRECT_COUNT)
+        {
+            Some(gpu_lists::Out::Draws)
+        } else {
+            None
+        }
+        .map(|out| gpu_lists::GpuLists::new(&device, out));
         let batchable = translucent_batchable(
             &adapter.get_info(),
             multi_draw,
@@ -6040,15 +6053,18 @@ impl Renderer {
             }
         };
         // Phase 1's sections can leave finding their facing quads to the GPU
-        // (`gpu_lists`) in frames drawn flat with nothing but the camera's
-        // terrain: no mirror, shader pack terrain or shadow, nor tiles.
+        // (`gpu_lists`) in frames with nothing but the camera's terrain (no
+        // mirror, shader pack terrain or shadow, nor tiles): drawn flat, or
+        // where multi-draws take their count from the GPU.
         let use_lists = gpu_lists_on()
             && !pack_terrain
             && shadow_view.is_none()
             && mirror.is_none()
             && visible_lods.is_empty()
-            && self.flat_on
-            && (self.flat.as_ref()).is_some_and(|f| self.drawn_sections >= f.min_sections);
+            && match &self.flat {
+                Some(flat) => self.flat_on && self.drawn_sections >= flat.min_sections,
+                None => self.multi_draw,
+            };
         if let Some(lists) = &mut self.gpu_lists {
             lists.clear(&self.device);
         }
@@ -6066,6 +6082,7 @@ impl Renderer {
                     &s.sides,
                     parts(fade_code(s, &v.key) != 0),
                     v.origin,
+                    i as u32,
                 )
             {
                 if let Some(c) = &mut cull {
@@ -6750,6 +6767,27 @@ impl Renderer {
                 bytemuck::cast_slice(&self.draws),
             );
         }
+        // Without flat draws, the GPU writes the lists' draws and counts them.
+        if let Some(lists) = self
+            .gpu_lists
+            .as_mut()
+            .filter(|l| l.out() == gpu_lists::Out::Draws && l.listed() > 0)
+        {
+            lists.rooms();
+            lists.prepare(
+                &self.device,
+                &mut encoder,
+                &mut self.uploads,
+                None,
+                eye.map(|e| e.to_array()),
+            );
+            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some("Section lists"),
+                timestamp_writes: None,
+            });
+            lists.dispatch(&mut pass);
+            self.section_draw_calls += lists.calls().len();
+        }
         let view_target = target_texture.create_view(&wgpu::TextureViewDescriptor {
             format: Some(self.config.format.remove_srgb_suffix()),
             ..Default::default()
@@ -7397,13 +7435,31 @@ impl Renderer {
                     // Without a native material, on Metal, the opaque pass, then the
                     // cutout one (terrain.wgsl's `opaque_quad`), each over phase 1's
                     // lists with quads for it, and over all of phase 2's.
+                    // Each pass's GPU-made draws by group, too (`gpu_lists`).
                     let (opaque, cutout) = (opaque_end as u32, cutout_start as u32);
-                    let passes: &[(&wgpu::RenderPipeline, std::ops::Range<u32>)] = match material {
+                    type Passes<'p> =
+                        &'p [(&'p wgpu::RenderPipeline, std::ops::Range<u32>, &'p [u8])];
+                    let passes: Passes = match material {
                         None if split_solid => &[
-                            (&self.solid_opaque, 0..opaque),
-                            (&self.solid_cutout, cutout..phase1_draws),
+                            (&self.solid_opaque, 0..opaque, &[0]),
+                            (&self.solid_cutout, cutout..phase1_draws, &[1]),
                         ],
-                        _ => &[(solid_pipeline, 0..phase1_draws)],
+                        _ => &[(solid_pipeline, 0..phase1_draws, &[0, 1])],
+                    };
+                    let lists = self
+                        .gpu_lists
+                        .as_ref()
+                        .filter(|l| l.out() == gpu_lists::Out::Draws && l.prepared());
+                    let draw_lists = |pass: &mut wgpu::RenderPass<'_>, groups: &[u8]| {
+                        let Some(lists) = lists else {
+                            return;
+                        };
+                        for (k, call) in lists.calls().iter().enumerate() {
+                            if groups.contains(&call.group) {
+                                pass.set_vertex_buffer(0, self.arena.buffer(call.page).slice(..));
+                                lists.multi_draw(pass, k);
+                            }
+                        }
                     };
                     pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint32);
                     pass.set_vertex_buffer(1, self.instances.slice(..));
@@ -7413,8 +7469,9 @@ impl Renderer {
                             // its pass, or one of their own on Metal), then
                             // what they found in sight, in a pass that loads
                             // both targets (it reads the test's results).
-                            for (pipeline, draws) in passes {
+                            for (pipeline, draws, groups) in passes {
                                 pass.set_pipeline(pipeline);
+                                draw_lists(&mut pass, groups);
                                 self.draw_sections(
                                     &mut pass,
                                     &self.indirect,
@@ -7448,7 +7505,7 @@ impl Renderer {
                                 wgpu::IndexFormat::Uint32,
                             );
                             pass.set_vertex_buffer(1, self.instances.slice(..));
-                            for (pipeline, _) in passes {
+                            for (pipeline, ..) in passes {
                                 pass.set_pipeline(pipeline);
                                 self.draw_sections(
                                     &mut pass,
@@ -7461,8 +7518,9 @@ impl Renderer {
                         }
                         // Without culling, phase 1 is all the solid draws.
                         None => {
-                            for (pipeline, draws) in passes {
+                            for (pipeline, draws, groups) in passes {
                                 pass.set_pipeline(pipeline);
+                                draw_lists(&mut pass, groups);
                                 self.draw_sections(
                                     &mut pass,
                                     &self.indirect,

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

@@ -664,7 +664,7 @@ impl Flat {
             );
         }
         if let Some(lists) = &mut lists {
-            lists.prepare(device, encoder, uploads, &self.indices, eye);
+            lists.prepare(device, encoder, uploads, Some(&self.indices), eye);
         }
         let ranges_at = |first: u64, count: u64| {
             wgpu::BindingResource::Buffer(wgpu::BufferBinding {

crates/horizon-client-render/src/gfx/gpu_lists.rs

@@ -5,10 +5,11 @@
 //! and pushed one draw per run. Here it pushes one record per section instead:
 //! its row (where its quads start and how `Sides` sorted them, uploaded with its
 //! mesh), the call each of its parts draws in and its camera-relative origin. The
-//! GPU finds the same runs and writes their indices, per call (one per pass
-//! group and arena page), into the flat draws' index list, drawn with
-//! `draw_indexed_indirect`. Calls take room for all the quads of their sections'
-//! parts; the GPU counts the indices it wrote.
+//! GPU finds the same runs and writes, per call (one per pass group and arena
+//! page), either the runs' indices into the flat draws' index list
+//! (`Out::Indices`, drawn with `draw_indexed_indirect`), or one indirect draw per
+//! run (`Out::Draws`, drawn with `multi_draw_indexed_indirect_count`). Calls take
+//! room for the most their sections may write; the GPU counts what they wrote.
 //!
 //! Solid draws write depth and are tested against it, so their order within a
 //! call does not change the picture.
@@ -16,9 +17,21 @@ use crate::mesher::{Part, Sides};
 
 /// Words per row: first vertex, `Sides::gpu_row`, padding.
 const ROW: usize = 24;
+/// Records each workgroup of `sides_draws` takes.
+const GROUP: u32 = 64;
 /// Most workgroups along one dimension of a dispatch.
 const GROUP_LIMIT: u32 = 65535;
 
+/// What the GPU writes for a frame's calls.
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+pub(super) enum Out {
+    /// Indices into the flat draws' list (`flat::Flat`), one indirect draw per call.
+    Indices,
+    /// `DrawIndexedIndirectArgs` per run, over the quad index pattern, with the
+    /// section's instance; per call a count for `multi_draw_indexed_indirect_count`.
+    Draws,
+}
+
 /// A part of a section in a record: `Part` as gpu_lists.wgsl names it.
 fn part_code(part: Part) -> u32 {
     match part {
@@ -29,7 +42,7 @@ fn part_code(part: Part) -> u32 {
 }
 
 /// One of the frame's calls: what pass `group` draws of arena `page`, from
-/// index `first` with room for `count`.
+/// `first` (an index, or a draw) with room for `count`.
 #[derive(Clone, Copy, Debug, PartialEq, Eq)]
 pub(super) struct ListCall {
     pub group: u8,
@@ -39,6 +52,7 @@ pub(super) struct ListCall {
 }
 
 pub(super) struct GpuLists {
+    out: Out,
     pipeline: wgpu::ComputePipeline,
     layout: wgpu::BindGroupLayout,
     /// Per id: its first vertex and `Sides::gpu_row` (`ROW` words).
@@ -47,7 +61,7 @@ pub(super) struct GpuLists {
     free_ids: Vec<u32>,
     next_id: u32,
     /// The frame's sections: id, the call and part of each of its parts, its
-    /// number of parts; then its origin.
+    /// number of parts plus its instance times 4; then its origin.
     records: Vec<[u32; 8]>,
     record_buffer: wgpu::Buffer,
     record_capacity: u64,
@@ -57,12 +71,16 @@ pub(super) struct GpuLists {
     call_of: Vec<u32>,
     band: u32,
     call_data: Vec<[u32; 5]>,
-    /// Per call `DrawIndexedIndirectArgs`; the GPU adds what it writes to the count.
+    /// Per call `DrawIndexedIndirectArgs` (`Out::Indices`), or its count of
+    /// draws first (`Out::Draws`); the GPU adds what it writes to the first word.
     call_buffer: wgpu::Buffer,
     call_capacity: u64,
     view: wgpu::Buffer,
+    /// `Out::Draws`' draws.
+    draws: wgpu::Buffer,
+    draw_capacity: u64,
     group: Option<(wgpu::BindGroup, u32)>,
-    /// The indices the frame's calls take, and the most they may.
+    /// The room the frame's calls take (indices, or draws), and the most they may.
     room: u64,
     cap: u64,
 }
@@ -76,7 +94,7 @@ fn buffer(
     device.create_buffer(&wgpu::BufferDescriptor {
         label: Some(label),
         size,
-        // COPY_SRC: rows move to a larger buffer; tests read calls back.
+        // COPY_SRC: rows move to a larger buffer; tests read calls and draws back.
         usage: usage | wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::COPY_SRC,
         mapped_at_creation: false,
     })
@@ -96,7 +114,7 @@ fn entry(binding: u32, ty: wgpu::BufferBindingType) -> wgpu::BindGroupLayoutEntr
 }
 
 impl GpuLists {
-    pub fn new(device: &wgpu::Device) -> Self {
+    pub fn new(device: &wgpu::Device, out: Out) -> Self {
         let read = wgpu::BufferBindingType::Storage { read_only: true };
         let write = wgpu::BufferBindingType::Storage { read_only: false };
         let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
@@ -123,13 +141,17 @@ impl GpuLists {
                 }),
             ),
             module: &module,
-            entry_point: Some("flatten_sides"),
+            entry_point: Some(match out {
+                Out::Indices => "flatten_sides",
+                Out::Draws => "sides_draws",
+            }),
             compilation_options: Default::default(),
             cache: None,
         });
         let storage = wgpu::BufferUsages::STORAGE;
         let indirect = wgpu::BufferUsages::INDIRECT | storage;
         Self {
+            out,
             pipeline,
             layout,
             rows: buffer(device, "Section rows", ROW as u64 * 4, storage),
@@ -151,6 +173,8 @@ impl GpuLists {
                 16,
                 wgpu::BufferUsages::UNIFORM,
             ),
+            draws: buffer(device, "Section list draws", 20, indirect),
+            draw_capacity: 1,
             group: None,
             room: 0,
             cap: 0,
@@ -228,14 +252,19 @@ impl GpuLists {
         self.band = 0;
         self.room = 0;
         self.group = None;
-        // A quarter of the binding they share with the CPU's draws.
-        self.cap = device.limits().max_storage_buffer_binding_size / 16;
+        // A quarter of a binding: indices share theirs with the CPU's draws.
+        let binding = device.limits().max_storage_buffer_binding_size;
+        self.cap = match self.out {
+            Out::Indices => binding / 16,
+            Out::Draws => binding / 20 / 4,
+        };
     }
 
     /// Section `id`, in `page`, at `origin` from the camera and `distance`
-    /// (squared, as bits) from it: each of `parts` (pass group, part) draws the
-    /// quads of the part that face the eye. False, listing nothing, when the
-    /// frame's lists hold all the room they may: the caller draws it.
+    /// (squared, as bits) from it, instance `instance` of the frame's list:
+    /// each of `parts` (pass group, part) draws the quads of the part that face
+    /// the eye. False, listing nothing, when the frame's lists hold all the room
+    /// they may: the caller draws it.
     ///
     /// Sections come nearest first. Each band of distance gets calls of its own,
     /// drawn in that order, as the CPU's lists keep them (`page_runs`): within
@@ -249,14 +278,18 @@ impl GpuLists {
         sides: &Sides,
         parts: &[(u8, Part)],
         origin: [f32; 3],
+        instance: u32,
     ) -> bool {
         let parts = &parts[..parts.len().min(2)];
         let mut counts = [0; 2];
         for (count, &(_, part)) in counts.iter_mut().zip(parts) {
-            *count = sides.part_quads(part) * 6;
+            *count = match self.out {
+                Out::Indices => sides.part_quads(part) * 6,
+                Out::Draws => sides.part_runs(part),
+            };
         }
         let room = u64::from(counts[0]) + u64::from(counts[1]);
-        if self.room + room > self.cap {
+        if self.room + room > self.cap || instance >= 1 << 30 {
             return false;
         }
         self.room += room;
@@ -265,7 +298,7 @@ impl GpuLists {
             self.band = band;
             self.call_of.fill(u32::MAX);
         }
-        let mut record = [id, 0, 0, parts.len() as u32, 0, 0, 0, 0];
+        let mut record = [id, 0, 0, parts.len() as u32 | instance << 2, 0, 0, 0, 0];
         for (k, &(group, part)) in parts.iter().enumerate() {
             let key = page as usize * 2 + usize::from(group.min(1));
             if self.call_of.len() <= key {
@@ -301,8 +334,8 @@ impl GpuLists {
         &self.calls
     }
 
-    /// Lays the calls' rooms out one after the other from the start of the
-    /// index list: the indices they take.
+    /// Lays the calls' rooms out one after the other from the start: the
+    /// room they take (indices, or draws).
     pub fn rooms(&mut self) -> u64 {
         let mut room = 0u64;
         for call in &mut self.calls {
@@ -313,14 +346,14 @@ impl GpuLists {
     }
 
     /// Stages the frame's sections, calls and eye (camera-relative; none: every
-    /// quad faces it), and binds them for `dispatch` writing into `indices`.
-    /// After `rooms`.
+    /// quad faces it), and binds them for `dispatch` writing into `indices`
+    /// (`Out::Indices`) or into the lists' own draws. After `rooms`.
     pub fn prepare(
         &mut self,
         device: &wgpu::Device,
         encoder: &mut wgpu::CommandEncoder,
         uploads: &mut super::upload::Uploads,
-        indices: &wgpu::Buffer,
+        indices: Option<&wgpu::Buffer>,
         eye: Option<[f32; 3]>,
     ) {
         self.group = None;
@@ -363,9 +396,26 @@ impl GpuLists {
             0,
             bytemuck::cast_slice(&self.call_data),
         );
+        let room = self
+            .calls
+            .last()
+            .map_or(0, |c| u64::from(c.first + c.count));
+        if self.out == Out::Draws && room > self.draw_capacity {
+            self.draw_capacity = room.next_power_of_two();
+            self.draws = buffer(
+                device,
+                "Section list draws",
+                self.draw_capacity * 20,
+                indirect,
+            );
+        }
         let [x, y, z] = eye.unwrap_or_default();
         let view = [x, y, z, if eye.is_some() { 1.0 } else { 0.0 }];
         uploads.buffer(encoder, &self.view, 0, bytemuck::cast_slice(&view));
+        let out = match self.out {
+            Out::Indices => indices.expect("flat draws' index list"),
+            Out::Draws => &self.draws,
+        };
         let group = device.create_bind_group(&wgpu::BindGroupDescriptor {
             label: Some("Section lists"),
             layout: &self.layout,
@@ -392,7 +442,7 @@ impl GpuLists {
                 },
                 wgpu::BindGroupEntry {
                     binding: 4,
-                    resource: indices.as_entire_binding(),
+                    resource: out.as_entire_binding(),
                 },
             ],
         });
@@ -404,21 +454,42 @@ impl GpuLists {
         let Some((group, records)) = &self.group else {
             return;
         };
-        let groups = *records;
+        let groups = match self.out {
+            Out::Indices => *records,
+            Out::Draws => records.div_ceil(GROUP),
+        };
         pass.set_pipeline(&self.pipeline);
         pass.set_bind_group(0, group, &[]);
         pass.dispatch_workgroups(groups.min(GROUP_LIMIT), groups.div_ceil(GROUP_LIMIT), 1);
     }
 
+    pub fn out(&self) -> Out {
+        self.out
+    }
+
     /// Whether `dispatch` writes this frame's lists.
     pub fn prepared(&self) -> bool {
         self.group.is_some()
     }
 
-    /// Draws call `k` with a pipeline, its page's vertices and the index list bound.
+    /// Draws call `k` (`Out::Indices`) with a pipeline, its page's vertices and
+    /// the index list bound.
     pub fn draw_indexed(&self, pass: &mut wgpu::RenderPass<'_>, k: usize) {
         pass.draw_indexed_indirect(&self.call_buffer, k as u64 * 20);
     }
+
+    /// Draws call `k` (`Out::Draws`) with a pipeline, its page's vertices, the
+    /// quad indices and the frame's instances bound.
+    pub fn multi_draw(&self, pass: &mut wgpu::RenderPass<'_>, k: usize) {
+        let call = self.calls[k];
+        pass.multi_draw_indexed_indirect_count(
+            &self.draws,
+            u64::from(call.first) * 20,
+            &self.call_buffer,
+            k as u64 * 20,
+            call.count,
+        );
+    }
 }
 
 #[cfg(test)]
@@ -485,9 +556,12 @@ mod tests {
         runs
     }
 
-    /// Lists `sections` from each of a few eyes; per eye, the lists, the
-    /// calls' words and the indices written.
-    fn run(sections: &[Section]) -> Option<Vec<(Option<[f32; 3]>, GpuLists, Vec<u32>, Vec<u32>)>> {
+    /// Lists `sections` for `out` from each of a few eyes; per eye, the
+    /// lists, the calls' words and the written words.
+    fn run(
+        out: Out,
+        sections: &[Section],
+    ) -> Option<Vec<(Option<[f32; 3]>, GpuLists, Vec<u32>, Vec<u32>)>> {
         let (device, queue) = device()?;
         let mut results = Vec::new();
         for eye in [
@@ -495,7 +569,7 @@ mod tests {
             Some([-40.0, -3.0, 50.5]),
             None,
         ] {
-            let mut lists = GpuLists::new(&device);
+            let mut lists = GpuLists::new(&device, out);
             lists.clear(&device);
             let mut rows = vec![0u32; sections.len() * ROW];
             for (id, s) in sections.iter().enumerate() {
@@ -503,7 +577,10 @@ mod tests {
                 rows[id * ROW + 1..id * ROW + 23].copy_from_slice(&s.sides.gpu_row());
                 // Nearest first, over several bands.
                 let distance = (id as f64 * 2.0).powi(2).to_bits();
-                assert!(lists.list(id as u32, s.page, distance, &s.sides, s.parts, s.origin));
+                let instance = id as u32 + 7;
+                assert!(lists.list(
+                    id as u32, s.page, distance, &s.sides, s.parts, s.origin, instance
+                ));
             }
             let room = lists.rooms();
             let init = |contents: &[u8], usage| {
@@ -515,16 +592,24 @@ mod tests {
             };
             let storage = wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC;
             lists.rows = init(bytemuck::cast_slice(&rows), storage);
-            let indices = init(&vec![0xffu8; room as usize * 4], storage);
+            let words = match out {
+                Out::Indices => room,
+                Out::Draws => room * 5,
+            };
+            let indices = init(&vec![0xffu8; words as usize * 4], storage);
             let mut encoder = device.create_command_encoder(&Default::default());
             let mut uploads = super::super::upload::Uploads::new(&device);
-            lists.prepare(&device, &mut encoder, &mut uploads, &indices, eye);
+            lists.prepare(&device, &mut encoder, &mut uploads, Some(&indices), eye);
             uploads.finish(&encoder);
             {
                 let mut pass = encoder.begin_compute_pass(&Default::default());
                 lists.dispatch(&mut pass);
             }
             queue.submit([encoder.finish()]);
+            let written = match out {
+                Out::Indices => &indices,
+                Out::Draws => &lists.draws,
+            };
             let read_back = |buffer: &wgpu::Buffer, words: u64| -> Vec<u32> {
                 let read = device.create_buffer(&wgpu::BufferDescriptor {
                     label: None,
@@ -540,7 +625,7 @@ mod tests {
                 bytemuck::cast_slice(&read.slice(..).get_mapped_range().unwrap()).to_vec()
             };
             let calls = read_back(&lists.call_buffer, lists.calls.len() as u64 * 5);
-            let words = read_back(&indices, room);
+            let words = read_back(written, words);
             results.push((eye, lists, calls, words));
         }
         Some(results)
@@ -571,13 +656,13 @@ mod tests {
         let Some((device, _)) = device() else {
             return;
         };
-        let mut lists = GpuLists::new(&device);
+        let mut lists = GpuLists::new(&device, Out::Indices);
         lists.clear(&device);
         let sides = Sides::unsorted(4);
         let parts = &[(0, Part::Opaque), (1, Part::Cutout)];
         for (page, blocks) in [(1, 10.0), (0, 20.0), (1, 30.0), (1, 100.0), (0, 130.0)] {
             let distance = f64::powi(blocks, 2).to_bits();
-            assert!(lists.list(0, page, distance, &sides, parts, [0.0; 3]));
+            assert!(lists.list(0, page, distance, &sides, parts, [0.0; 3], 0));
         }
         let calls: Vec<_> = lists
             .calls()
@@ -605,7 +690,7 @@ mod tests {
     #[test]
     fn flatten_sides_writes_the_runs_the_cpu_finds() {
         let sections = sections(200);
-        let Some(results) = run(&sections) else {
+        let Some(results) = run(Out::Indices, &sections) else {
             return;
         };
         for (eye, lists, calls, out) in results {
@@ -636,4 +721,44 @@ mod tests {
             }
         }
     }
+
+    /// `sides_draws` writes, in each call's room, one draw per run the CPU finds
+    /// (as the CPU lists them: the run's indices over the quad pattern, from its
+    /// first vertex, with the section's instance), and counts them; the room
+    /// holds them all.
+    #[test]
+    fn sides_draws_writes_a_draw_per_run_the_cpu_finds() {
+        let sections = sections(200);
+        let Some(results) = run(Out::Draws, &sections) else {
+            return;
+        };
+        for (eye, lists, calls, out) in results {
+            let mut expected = vec![Vec::new(); lists.calls.len()];
+            for (i, s) in sections.iter().enumerate() {
+                for (k, &(_, part)) in s.parts.iter().enumerate() {
+                    let call = call_of(&lists, i, k);
+                    for (first, quads) in cpu_runs(s, part, eye) {
+                        expected[call].push([quads * 6, 1, 0, s.start + first * 4, i as u32 + 7]);
+                    }
+                }
+            }
+            for (k, call) in lists.calls.iter().enumerate() {
+                let count = calls[k * 5] as usize;
+                assert_eq!(
+                    count,
+                    expected[k].len(),
+                    "call {k}'s draw count, eye {eye:?}"
+                );
+                assert!(count as u32 <= call.count, "call {k} wrote past its room");
+                let first = call.first as usize * 5;
+                let mut got: Vec<[u32; 5]> = out[first..first + count * 5]
+                    .chunks(5)
+                    .map(|d| d.try_into().unwrap())
+                    .collect();
+                got.sort_unstable();
+                expected[k].sort_unstable();
+                assert_eq!(got, expected[k], "call {k}'s draws, eye {eye:?}");
+            }
+        }
+    }
 }

crates/horizon-client-render/src/gfx/gpu_lists.wgsl

@@ -2,9 +2,11 @@
 // listed, the runs of its quads that face the eye, found as
 // `mesher::Sides::visible_from` finds them on the CPU, written into the call
 // each of its parts draws in (one per pass group and arena page) after taking
-// room there: their indices, in the flat draws' index list, as `flatten_seen`
-// does for occlusion culling's phase 2. Solid draws write depth and are tested
-// against it, so their order within a call does not change the picture.
+// room there: their indices (`flatten_sides`, for the flat draws' index list,
+// as `flatten_seen` does for occlusion culling's phase 2), or one indirect draw
+// per run (`sides_draws`, for `multi_draw_indexed_indirect_count`). Solid draws
+// write depth and are tested against it, so their order within a call does not
+// change the picture.
 struct View {
     // The eye relative to the camera, and 1.0 when there is one (`sides_on`);
     // without, every bucket is drawn (`Sides::every`).
@@ -12,8 +14,8 @@ struct View {
     has_eye: f32,
 }
 // Per section listed: its row's id; for each of its parts (up to two) the call
-// it draws in, with the part (`PART_*`) in the top bits; its number of parts.
-// Then its origin relative to the camera.
+// it draws in, with the part (`PART_*`) in the top bits; its number of parts
+// plus 4 times its instance. Then its origin relative to the camera.
 struct Record {
     words: vec4<u32>,
     origin: vec4<f32>,
@@ -23,10 +25,11 @@ struct Record {
 // (14 ends, 7 planes' bits, flags).
 @group(0) @binding(1) var<storage, read> rows: array<u32>;
 @group(0) @binding(2) var<uniform> view: View;
-// Per call `DrawIndexedIndirectArgs`: its index count starts at 0, its first
-// index is where its room starts.
+// Per call five words: its count (indices, or draws), which starts at 0, then
+// for `flatten_sides` the rest of its `DrawIndexedIndirectArgs`; word 2 is
+// where its room starts.
 @group(0) @binding(3) var<storage, read_write> calls: array<atomic<u32>>;
-// The flat draws' index list.
+// Indices (`flatten_sides`) or `DrawIndexedIndirectArgs` (`sides_draws`).
 @group(0) @binding(4) var<storage, read_write> out: array<u32>;
 
 const ROW: u32 = 24u;
@@ -146,7 +149,7 @@ fn flatten_sides(@builtin(workgroup_id) group: vec3<u32>,
     }
     let record = records[r];
     let row = record.words.x * ROW;
-    let parts = record.words.w;
+    let parts = record.words.w & 3u;
     if lane == 0u {
         let eye = eye_of(record);
         for (var p = 0u; p < parts; p++) {
@@ -176,3 +179,38 @@ fn flatten_sides(@builtin(workgroup_id) group: vec3<u32>,
         }
     }
 }
+
+// One record per invocation: a draw per run, as the CPU lists them
+// (`[quads * 6, 1, 0, first vertex, instance]` over the quad index pattern).
+@compute @workgroup_size(64)
+fn sides_draws(@builtin(workgroup_id) group: vec3<u32>,
+    @builtin(local_invocation_index) lane: u32) {
+    let r = record_index(group) * LANES + lane;
+    if r >= arrayLength(&records) {
+        return;
+    }
+    let record = records[r];
+    let row = record.words.x * ROW;
+    let parts = record.words.w & 3u;
+    let instance = record.words.w >> 2u;
+    let eye = eye_of(record);
+    let base = rows[row];
+    for (var p = 0u; p < parts; p++) {
+        let word = part_word(record, p);
+        let runs = part_runs(row, word >> 28u, eye);
+        if runs.n == 0u {
+            continue;
+        }
+        let call = (word & CALL_MASK) * 5u;
+        var to = (atomicLoad(&calls[call + 2u]) + atomicAdd(&calls[call], runs.n)) * 5u;
+        for (var k = 0u; k < runs.n; k++) {
+            let run = runs.runs[k];
+            out[to] = run.y * 6u;
+            out[to + 1u] = 1u;
+            out[to + 2u] = 0u;
+            out[to + 3u] = base + run.x * 4u;
+            out[to + 4u] = instance;
+            to += 5u;
+        }
+    }
+}

crates/horizon-client-render/src/mesher/sides.rs

@@ -169,6 +169,31 @@ impl Sides {
         self.runs(|_| true, part, draw);
     }
 
+    /// Whether `part` draws the cutout (or opaque) quads of a bucket.
+    fn wants(&self, part: Part, cutout: bool) -> bool {
+        match part {
+            Part::All => true,
+            _ if !self.split => true,
+            Part::Opaque => !cutout,
+            Part::Cutout => cutout,
+        }
+    }
+
+    /// The most runs `visible_from` finds for `part`, whatever the eye: one per
+    /// part of a bucket it draws that holds quads, and at most 7, as runs are
+    /// apart (gpu_lists.wgsl's room for its draws).
+    pub fn part_runs(&self, part: Part) -> u32 {
+        let mut start = 0;
+        let mut runs = 0;
+        for (n, &end) in self.ends.iter().enumerate() {
+            if end != start && self.wants(part, slot(self.grouped, n).1) {
+                runs += 1;
+            }
+            start = end;
+        }
+        runs.min(7)
+    }
+
     /// Calls `draw` for each run of `part`'s quads in the buckets `seen` keeps.
     fn runs(&self, seen: impl Fn(usize) -> bool, part: Part, mut draw: impl FnMut(u32, u32)) {
         let mut start = 0;
@@ -178,13 +203,7 @@ impl Sides {
                 continue;
             }
             let (bucket, cutout) = slot(self.grouped, n);
-            let wanted = match part {
-                Part::All => true,
-                _ if !self.split => true,
-                Part::Opaque => !cutout,
-                Part::Cutout => cutout,
-            };
-            match (wanted && seen(bucket), run) {
+            match (self.wants(part, cutout) && seen(bucket), run) {
                 (true, None) => run = Some(start),
                 (false, Some(first)) => {
                     draw(first, start - first);