Flat draws: phase 1's solid lists made on the GPU

In frames drawn flat (Metal at render distance 32, Intel Gen11/12 Vulkan),
the render thread walked every phase 1 section's sides each frame to find the
runs of quads that face the eye, then the flat plan copied them into one index
list. Now each section uploads a row once (where its quads start, how Sides
sorted them), and each frame phase 1 pushes one record per section: the GPU
(gpu_lists.wgsl's flatten_sides) finds the same runs as Sides::visible_from and
writes their indices into a call per pass and arena page, drawn indirect.
Phase 2, translucent quads, tiles, mirrors, shader pack terrain and shadows
keep the CPU lists. HORIZON_GPU_LISTS=0 turns it off.

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

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

@@ -112,6 +112,7 @@ mod cull;
 mod end_sky;
 mod flat;
 mod frame_gen;
+mod gpu_lists;
 mod hdr;
 mod hdr_icc;
 mod hdr_stats;
@@ -357,6 +358,8 @@ struct SectionGpu {
     sorted_indices: Option<(u32, u64)>,
     /// Its id for flat draws (`flat::Flat::add`), when they are on.
     flat_id: Option<u32>,
+    /// Its row for phase 1's lists made on the GPU (`gpu_lists::GpuLists::add`).
+    list_id: Option<u32>,
     /// When it began to fade in over the far terrain it replaces, which
     /// stays whole behind it until then: a section new where far tiles are
     /// drawn. Meshed again, it stays as it was.
@@ -447,12 +450,19 @@ impl Part {
 /// hold on `page` of `arena` (`part`) goes to its other pages in video memory,
 /// until the next range would pass `budget` bytes (it keeps what is left; the
 /// first range moves whatever its size, so none keeps the page closed), the
-/// copies into `encoder` and, in the section arena, the ranges' flat draw ids
-/// (`flat`) along. False when the other pages have no room left.
+/// copies into `encoder` and, in the section arena, the ranges' flat draw ids and
+/// GPU list rows (`ids`) along. False when the other pages have no room left.
+/// The flat draws and GPU lists that name a mesh's range, with what writes to them.
+type MeshIds<'a> = (
+    Option<&'a mut flat::Flat>,
+    Option<&'a mut gpu_lists::GpuLists>,
+    &'a wgpu::Device,
+    &'a mut upload::Uploads,
+);
 fn move_off_page<'a>(
     arena: &mut arena::Arena,
     encoder: &mut wgpu::CommandEncoder,
-    mut flat: Option<(&mut flat::Flat, &wgpu::Device, &mut upload::Uploads)>,
+    mut ids: Option<MeshIds<'_>>,
     meshes: impl Iterator<Item = &'a mut SectionGpu>,
     page: u32,
     part: Part,
@@ -474,9 +484,14 @@ fn move_off_page<'a>(
         };
         arena.copy(encoder, (page, start), to, vertices);
         arena.release(page, start, vertices);
-        if let (Some((flat, device, uploads)), Some(id)) = (&mut flat, s.flat_id) {
-            flat.remove(id);
-            s.flat_id = Some(flat.add(device, encoder, uploads, to.0, to.1, vertices));
+        if let Some((flat, lists, device, uploads)) = &mut ids {
+            if let (Some(flat), Some(id)) = (flat, s.flat_id) {
+                flat.remove(id);
+                s.flat_id = Some(flat.add(device, encoder, uploads, to.0, to.1, vertices));
+            }
+            if let (Some(lists), Some(id)) = (lists, s.list_id) {
+                lists.set(device, encoder, uploads, id, to.1 as u32, &s.sides);
+            }
         }
         part.set(s, to);
         *budget = budget.saturating_sub(bytes);
@@ -777,6 +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`).
+    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).
     cull: Option<cull::Cull>,
@@ -3497,6 +3514,7 @@ impl Renderer {
                 .get_downlevel_capabilities()
                 .flags
                 .contains(wgpu::DownlevelFlags::INDIRECT_EXECUTION);
+        let gpu_lists = flat.as_ref().map(|_| gpu_lists::GpuLists::new(&device));
         let batchable = translucent_batchable(
             &adapter.get_info(),
             multi_draw,
@@ -3611,6 +3629,7 @@ impl Renderer {
             solid_cutout,
             translucent,
             flat_on: flat.is_some(),
+            gpu_lists,
             cull_on: cull.is_some(),
             sides_on: true,
             cull,
@@ -4157,6 +4176,11 @@ impl Renderer {
         } else {
             None
         };
+        let sides = if mesh.sides.quads() == solid {
+            mesh.sides
+        } else {
+            Sides::unsorted(solid)
+        };
         // Flat draws name the blocks of its range in `arena`: split, its solid
         // quads alone (no frame draws flat then: `translucent_split`).
         let flat_id = self.flat.as_mut().filter(|_| main > 0).map(|f| {
@@ -4169,6 +4193,15 @@ impl Renderer {
                 main,
             )
         });
+        let list_id = self.gpu_lists.as_mut().filter(|_| main > 0).map(|l| {
+            l.add(
+                &self.device,
+                encoder,
+                &mut self.section_uploads,
+                start as u32,
+                &sides,
+            )
+        });
         let stride = std::mem::size_of::<TerrainVertex>() as u64;
         let translucent_arena = self.translucent_arena.as_ref().unwrap_or(&self.arena);
         for (part, arena, (part_page, first)) in [
@@ -4199,11 +4232,7 @@ impl Renderer {
             start,
             translucent_at,
             solid,
-            sides: if mesh.sides.quads() == solid {
-                mesh.sides
-            } else {
-                Sides::unsorted(solid)
-            },
+            sides,
             translucent,
             sort: translucency::SortState::new(
                 mesh.translucent
@@ -4215,6 +4244,7 @@ impl Renderer {
             ),
             sorted_indices,
             flat_id,
+            list_id,
             fading: None,
             leaving: None,
             animated_layers,
@@ -4287,6 +4317,9 @@ impl Renderer {
         if let (Some(flat), Some(id)) = (&mut self.flat, s.flat_id) {
             flat.remove(id);
         }
+        if let (Some(lists), Some(id)) = (&mut self.gpu_lists, s.list_id) {
+            lists.remove(id);
+        }
         if let Some((page, start)) = s.sorted_indices {
             self.sorted_indices
                 .release(page, start, u64::from(s.translucent) * 6);
@@ -4329,13 +4362,15 @@ impl Renderer {
             .sections
             .values_mut()
             .chain(self.lods.values_mut().flat_map(|tile| tile.meshes_mut()));
-        let flat = self
-            .flat
-            .as_mut()
-            .map(|f| (f, &self.device, &mut self.section_uploads));
+        let ids = Some((
+            self.flat.as_mut(),
+            self.gpu_lists.as_mut(),
+            &self.device,
+            &mut self.section_uploads,
+        ));
         let split = self.translucent_arena.is_some();
         let (arena, part) = (&mut self.arena, Part::Section { split });
-        if !move_off_page(arena, encoder, flat, meshes, page, part, &mut budget) {
+        if !move_off_page(arena, encoder, ids, meshes, page, part, &mut budget) {
             // The other pages have no room after all.
             arena.reopen();
         }
@@ -4641,6 +4676,11 @@ impl Renderer {
                 flat.remove(id);
             }
         }
+        if let Some(lists) = &mut self.gpu_lists {
+            for id in released.iter().filter_map(|s| s.list_id) {
+                lists.remove(id);
+            }
+        }
         self.sorted_indices
             .release_all(released.iter().filter_map(|s| {
                 let (page, start) = s.sorted_indices?;
@@ -5999,10 +6039,40 @@ impl Renderer {
                 (true, true) => &[(1, Quads::All)],
             }
         };
+        // 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.
+        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);
+        if let Some(lists) = &mut self.gpu_lists {
+            lists.clear(&self.device);
+        }
+        let mut lists = self.gpu_lists.as_mut().filter(|_| use_lists);
         let mut opaque_draws = 0;
         for (i, v) in visible.iter().enumerate() {
             let s = self.sections.get(v.slot);
             let first_phase = cull.as_ref().is_none_or(|c| c.first(v.slot));
+            if first_phase
+                && let (Some(lists), Some(id)) = (&mut lists, s.list_id)
+                && lists.list(
+                    id,
+                    s.page,
+                    v.distance,
+                    &s.sides,
+                    parts(fade_code(s, &v.key) != 0),
+                    v.origin,
+                )
+            {
+                if let Some(c) = &mut cull {
+                    c.drawn(i, v.slot, s.sides.quads() * 6);
+                }
+                continue;
+            }
             let (draws, pages) = if first_phase {
                 (&mut self.draws, &mut self.draw_pages)
             } else {
@@ -6612,7 +6682,8 @@ impl Renderer {
         let flat_frame = match &mut self.flat {
             Some(flat)
                 if self.flat_on
-                    && !self.draws.is_empty()
+                    && (!self.draws.is_empty()
+                        || self.gpu_lists.as_ref().is_some_and(|l| l.listed() > 0))
                     && (self.drawn_sections >= flat.min_sections
                         || pack_terrain
                         || shadow_view.is_some()) =>
@@ -6651,6 +6722,8 @@ impl Renderer {
                                 Some((id, lod::instance(v.origin, v.fade)))
                             })),
                         |page| sorted.buffer(page).clone(),
+                        self.gpu_lists.as_mut().filter(|l| l.listed() > 0),
+                        eye.map(|e| e.to_array()),
                         self.own_timing.as_mut(),
                     )
                 })
@@ -6658,7 +6731,8 @@ impl Renderer {
             _ => false,
         };
         if flat_frame {
-            self.section_draw_calls = self.flat.as_ref().map_or(0, |f| f.shadow_calls[0]);
+            self.section_draw_calls = self.flat.as_ref().map_or(0, |f| f.shadow_calls[0])
+                + self.gpu_lists.as_ref().map_or(0, |l| l.calls().len());
         } else if self.multi_draw && !self.draws.is_empty() {
             if self.draws.len() as u64 > self.indirect_capacity {
                 self.indirect_capacity = (self.draws.len() as u64).next_power_of_two();
@@ -7259,19 +7333,26 @@ impl Renderer {
                     // Without a pack, 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 lists by group, too (`flat::Flat::list`).
                     let [opaque, cutout] = flat.pass_calls;
-                    let passes: &[(&wgpu::RenderPipeline, std::ops::Range<usize>)] = match pack_flat
-                    {
-                        Some(p) => &[(&p.0, 0..flat.phase1_calls)],
+                    type Passes<'p> =
+                        &'p [(&'p wgpu::RenderPipeline, std::ops::Range<usize>, &'p [u8])];
+                    let passes: Passes = match pack_flat {
+                        Some(p) => &[(&p.0, 0..flat.phase1_calls, &[0, 1])],
                         None if split_solid => &[
-                            (&flat.solid_opaque, 0..opaque),
-                            (&flat.solid_cutout, cutout..flat.phase1_calls),
+                            (&flat.solid_opaque, 0..opaque, &[0]),
+                            (&flat.solid_cutout, cutout..flat.phase1_calls, &[1]),
                         ],
-                        None => &[(&flat.solid, 0..flat.phase1_calls)],
+                        None => &[(&flat.solid, 0..flat.phase1_calls, &[0, 1])],
                     };
                     let draw = |pass: &mut wgpu::RenderPass<'_>| {
-                        for (pipeline, calls) in passes {
+                        for (pipeline, calls, groups) in passes {
                             pass.set_pipeline(pipeline);
+                            if let Some(lists) = &self.gpu_lists {
+                                for &group in *groups {
+                                    flat.draw_lists(pass, lists, group, vertices);
+                                }
+                            }
                             flat.draw(pass, &flat.calls[calls.clone()], vertices);
                         }
                     };
@@ -7301,7 +7382,7 @@ impl Renderer {
                             };
                             pass.set_bind_group(0, &self.group, &[]);
                             pack_frame(&mut pass);
-                            for (pipeline, _) in passes {
+                            for (pipeline, ..) in passes {
                                 pass.set_pipeline(pipeline);
                                 flat.draw_phase2(&mut pass, vertices);
                             }
@@ -8953,6 +9034,12 @@ fn projection_center(view_proj: Mat4) -> Option<Vec3> {
 fn facing_runs(sides: &Sides, eye: Option<Vec3>, origin: [f32; 3], draw: impl FnMut(u32, u32)) {
     facing_runs_from(sides, eye, None, origin, Quads::All, draw)
 }
+/// Whether phase 1's solid lists may be made on the GPU (`flat::Flat::list`);
+/// `HORIZON_GPU_LISTS=0` keeps them on the CPU.
+fn gpu_lists_on() -> bool {
+    static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
+    *ON.get_or_init(|| std::env::var("HORIZON_GPU_LISTS").as_deref() != Ok("0"))
+}
 /// `facing_runs` for the eye and, with `low`, every eye straight below it down to that
 /// camera-relative height (the water mirror's eye).
 fn facing_runs_from(

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

@@ -836,6 +836,7 @@ mod tests {
             sort: super::super::translucency::SortState::new(Vec::new()),
             sorted_indices: None,
             flat_id: None,
+            list_id: None,
             fading: None,
             leaving: None,
             animated_layers: Box::new([]),
@@ -1119,6 +1120,7 @@ mod tests {
                 sort: super::super::translucency::SortState::new(Vec::new()),
                 sorted_indices: (translucent > 0).then_some(if indices { at } else { (7, 7) }),
                 flat_id: None,
+                list_id: None,
                 fading: None,
                 leaving: None,
                 animated_layers: Box::new([]),

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

@@ -223,6 +223,7 @@ fn buffer(
 /// from `shadow[0]` (solid) and `shadow[1]` (translucent) on are a shader
 /// pack's shadow camera's, with calls of their own. Phase 1's draws before
 /// `passes[0]` and from `passes[1]` on, the two solid passes' on Metal, too.
+/// The list starts at index `base` (the GPU-made lists' rooms come first: `gpu_lists`).
 #[allow(clippy::too_many_arguments)]
 fn plan(
     draws: &[[u32; 5]],
@@ -235,6 +236,7 @@ fn plan(
     calls: &mut Vec<Call>,
     ranges: &mut Vec<Range>,
     segments: &mut Vec<(Source, u64, u64)>,
+    base: u64,
 ) -> Plan {
     calls.clear();
     ranges.clear();
@@ -242,7 +244,7 @@ fn plan(
     let (mut phase1_calls, mut solid_calls) = (None, None);
     let mut shadow_calls = [None; 2];
     let mut pass_calls = [None; 2];
-    let mut total = 0u64;
+    let mut total = base;
     let mut by_source: Vec<(Source, Vec<Range>)> = Vec::new();
     for (i, (d, &(page, sorted))) in draws.iter().zip(pages).enumerate() {
         if i == phase1_draws {
@@ -540,6 +542,8 @@ impl Flat {
         phase2: Option<&wgpu::Buffer>,
         origins: impl Iterator<Item = (u32, [f32; 4])>,
         sorted_page: impl Fn(u32) -> wgpu::Buffer,
+        mut lists: Option<&mut super::gpu_lists::GpuLists>,
+        eye: Option<[f32; 3]>,
         timing: Option<&mut super::pack_timing::PassTiming>,
     ) -> bool {
         self.phase2_group = None;
@@ -577,6 +581,8 @@ impl Flat {
                 bytemuck::cast_slice(&self.origin_data),
             );
         }
+        // The GPU-made lists' rooms come first in the index list.
+        let room = lists.as_mut().map_or(0, |l| l.rooms());
         let Plan {
             phase1_calls,
             solid_calls,
@@ -594,6 +600,7 @@ impl Flat {
             &mut self.calls,
             &mut self.range_data,
             &mut self.segments,
+            room,
         );
         self.phase1_calls = phase1_calls;
         self.solid_calls = solid_calls;
@@ -656,6 +663,9 @@ impl Flat {
                 bytemuck::cast_slice(&self.phase2_call_data),
             );
         }
+        if let Some(lists) = &mut lists {
+            lists.prepare(device, encoder, uploads, &self.indices, eye);
+        }
         let ranges_at = |first: u64, count: u64| {
             wgpu::BindingResource::Buffer(wgpu::BufferBinding {
                 buffer: &self.ranges,
@@ -725,13 +735,17 @@ impl Flat {
             label: Some("Flatten section draws"),
             timestamp_writes: timing.and_then(|t| t.compute_writes("Flatten")),
         });
+        if let Some(lists) = &lists {
+            lists.dispatch(&mut pass);
+        }
         pass.set_pipeline(&self.flatten);
         for (group, count) in &groups {
             pass.set_bind_group(0, group, &[]);
             pass.dispatch_workgroups((*count).min(GROUP_LIMIT), count.div_ceil(GROUP_LIMIT), 1);
         }
         drop(pass);
-        for page in self.calls.iter().map(|c| c.page) {
+        let list_pages = lists.iter().flat_map(|l| l.calls().iter().map(|c| c.page));
+        for page in self.calls.iter().map(|c| c.page).chain(list_pages) {
             let table = self.pages[page as usize]
                 .as_mut()
                 .expect("a drawn page has its block table");
@@ -772,6 +786,27 @@ impl Flat {
         pass.dispatch_workgroups((*count).min(GROUP_LIMIT), count.div_ceil(GROUP_LIMIT), 1);
     }
 
+    /// Draws pass `group`'s calls of `lists` (`gpu_lists::Out::Indices`, prepared with
+    /// this frame's `prepare`) with a flat pipeline set.
+    pub fn draw_lists<'a>(
+        &self,
+        pass: &mut wgpu::RenderPass<'_>,
+        lists: &super::gpu_lists::GpuLists,
+        group: u8,
+        vertices: impl Fn(u32) -> &'a wgpu::Buffer,
+    ) {
+        if !lists.prepared() {
+            return;
+        }
+        pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint32);
+        for (k, call) in lists.calls().iter().enumerate() {
+            if call.group == group {
+                self.bind_page(pass, call.page, &vertices);
+                lists.draw_indexed(pass, k);
+            }
+        }
+    }
+
     /// The camera's translucent calls.
     pub fn translucent_calls(&self) -> &[Call] {
         &self.calls[self.solid_calls..self.shadow_calls[0]]
@@ -897,6 +932,7 @@ mod tests {
             &mut calls,
             &mut ranges,
             &mut segments,
+            0,
         );
         assert_eq!(total, 6 * 13);
         assert_eq!(phase1_calls, solid_calls);
@@ -945,6 +981,7 @@ mod tests {
             &mut calls,
             &mut ranges,
             &mut segments,
+            0,
         );
         assert_eq!((plan.pass_calls, plan.solid_calls), ([3, 1], 4));
         let call = |page, first, count| Call { page, first, count };
@@ -980,6 +1017,7 @@ mod tests {
             &mut calls,
             &mut ranges,
             &mut segments,
+            0,
         );
         assert_eq!((plan.solid_calls, plan.shadow_calls), (3, [4, 5]));
         let call = |page, first, count| Call { page, first, count };
@@ -1008,6 +1046,7 @@ mod tests {
             &mut calls,
             &mut ranges,
             &mut segments,
+            0,
         );
         assert_eq!(
             plan,
@@ -1125,6 +1164,7 @@ mod tests {
                 &mut calls,
                 &mut ranges,
                 &mut segments,
+                0,
             );
             // Phase 2's draws as the test leaves them: instance counts of 0,
             // 1 where it saw the section.

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

@@ -0,0 +1,639 @@
+//! Phase 1's solid section draws made on the GPU (gpu_lists.wgsl).
+//!
+//! Each frame the render thread walked every section phase 1 draws through
+//! `mesher::Sides::visible_from` to find the runs of its quads that face the eye,
+//! 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.
+//!
+//! Solid draws write depth and are tested against it, so their order within a
+//! call does not change the picture.
+use crate::mesher::{Part, Sides};
+
+/// Words per row: first vertex, `Sides::gpu_row`, padding.
+const ROW: usize = 24;
+/// Most workgroups along one dimension of a dispatch.
+const GROUP_LIMIT: u32 = 65535;
+
+/// A part of a section in a record: `Part` as gpu_lists.wgsl names it.
+fn part_code(part: Part) -> u32 {
+    match part {
+        Part::All => 0,
+        Part::Opaque => 1,
+        Part::Cutout => 2,
+    }
+}
+
+/// One of the frame's calls: what pass `group` draws of arena `page`, from
+/// index `first` with room for `count`.
+#[derive(Clone, Copy, Debug, PartialEq, Eq)]
+pub(super) struct ListCall {
+    pub group: u8,
+    pub page: u32,
+    pub first: u32,
+    pub count: u32,
+}
+
+pub(super) struct GpuLists {
+    pipeline: wgpu::ComputePipeline,
+    layout: wgpu::BindGroupLayout,
+    /// Per id: its first vertex and `Sides::gpu_row` (`ROW` words).
+    rows: wgpu::Buffer,
+    row_capacity: u64,
+    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.
+    records: Vec<[u32; 8]>,
+    record_buffer: wgpu::Buffer,
+    record_capacity: u64,
+    calls: Vec<ListCall>,
+    /// `calls`' index by page and group in the band of distance the last
+    /// section listed is in (`u32::MAX`: none yet).
+    call_of: Vec<u32>,
+    band: u32,
+    call_data: Vec<[u32; 5]>,
+    /// Per call `DrawIndexedIndirectArgs`; the GPU adds what it writes to the count.
+    call_buffer: wgpu::Buffer,
+    call_capacity: u64,
+    view: wgpu::Buffer,
+    group: Option<(wgpu::BindGroup, u32)>,
+    /// The indices the frame's calls take, and the most they may.
+    room: u64,
+    cap: u64,
+}
+
+fn buffer(
+    device: &wgpu::Device,
+    label: &str,
+    size: u64,
+    usage: wgpu::BufferUsages,
+) -> wgpu::Buffer {
+    device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some(label),
+        size,
+        // COPY_SRC: rows move to a larger buffer; tests read calls back.
+        usage: usage | wgpu::BufferUsages::COPY_DST | wgpu::BufferUsages::COPY_SRC,
+        mapped_at_creation: false,
+    })
+}
+
+fn entry(binding: u32, ty: wgpu::BufferBindingType) -> wgpu::BindGroupLayoutEntry {
+    wgpu::BindGroupLayoutEntry {
+        binding,
+        visibility: wgpu::ShaderStages::COMPUTE,
+        ty: wgpu::BindingType::Buffer {
+            ty,
+            has_dynamic_offset: false,
+            min_binding_size: None,
+        },
+        count: None,
+    }
+}
+
+impl GpuLists {
+    pub fn new(device: &wgpu::Device) -> 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 {
+            label: Some("Section lists"),
+            entries: &[
+                entry(0, read),
+                entry(1, read),
+                entry(2, wgpu::BufferBindingType::Uniform),
+                entry(3, write),
+                entry(4, write),
+            ],
+        });
+        let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("Section lists"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("gpu_lists.wgsl").into()),
+        });
+        let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+            label: Some("Section lists"),
+            layout: Some(
+                &device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+                    label: Some("Section lists"),
+                    bind_group_layouts: &[Some(&layout)],
+                    ..Default::default()
+                }),
+            ),
+            module: &module,
+            entry_point: Some("flatten_sides"),
+            compilation_options: Default::default(),
+            cache: None,
+        });
+        let storage = wgpu::BufferUsages::STORAGE;
+        let indirect = wgpu::BufferUsages::INDIRECT | storage;
+        Self {
+            pipeline,
+            layout,
+            rows: buffer(device, "Section rows", ROW as u64 * 4, storage),
+            row_capacity: 1,
+            free_ids: Vec::new(),
+            next_id: 0,
+            records: Vec::new(),
+            record_buffer: buffer(device, "Sections listed", 32, storage),
+            record_capacity: 1,
+            calls: Vec::new(),
+            call_of: Vec::new(),
+            band: 0,
+            call_data: Vec::new(),
+            call_buffer: buffer(device, "Section list calls", 20, indirect),
+            call_capacity: 1,
+            view: buffer(
+                device,
+                "Section lists' eye",
+                16,
+                wgpu::BufferUsages::UNIFORM,
+            ),
+            group: None,
+            room: 0,
+            cap: 0,
+        }
+    }
+
+    /// A mesh uploaded from vertex `start` of its page, its solid quads sorted as
+    /// `sides` says: its id, whose row is written through `uploads` in `encoder`.
+    pub fn add(
+        &mut self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        uploads: &mut super::upload::Uploads,
+        start: u32,
+        sides: &Sides,
+    ) -> u32 {
+        let id = self.free_ids.pop().unwrap_or_else(|| {
+            self.next_id += 1;
+            self.next_id - 1
+        });
+        self.set(device, encoder, uploads, id, start, sides);
+        id
+    }
+
+    pub fn remove(&mut self, id: u32) {
+        self.free_ids.push(id);
+    }
+
+    /// Mesh `id` now starts at vertex `start` (moved to another page).
+    pub fn set(
+        &mut self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        uploads: &mut super::upload::Uploads,
+        id: u32,
+        start: u32,
+        sides: &Sides,
+    ) {
+        let need = u64::from(id) + 1;
+        if need > self.row_capacity {
+            let capacity = need.next_power_of_two().max(1024);
+            let rows = buffer(
+                device,
+                "Section rows",
+                capacity * ROW as u64 * 4,
+                wgpu::BufferUsages::STORAGE,
+            );
+            // Earlier rows: uploads record their copies in `encoder`, so this one follows them.
+            encoder.copy_buffer_to_buffer(
+                &self.rows,
+                0,
+                &rows,
+                0,
+                self.row_capacity * ROW as u64 * 4,
+            );
+            self.rows = rows;
+            self.row_capacity = capacity;
+        }
+        let mut row = [0u32; ROW];
+        row[0] = start;
+        row[1..23].copy_from_slice(&sides.gpu_row());
+        uploads.buffer(
+            encoder,
+            &self.rows,
+            u64::from(id) * ROW as u64 * 4,
+            bytemuck::cast_slice(&row),
+        );
+    }
+
+    /// Starts the frame's lists.
+    pub fn clear(&mut self, device: &wgpu::Device) {
+        self.records.clear();
+        self.calls.clear();
+        self.call_of.fill(u32::MAX);
+        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;
+    }
+
+    /// 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.
+    ///
+    /// 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
+    /// a call the GPU writes sections in no set order, and cutout quads, which
+    /// tile GPUs can't hide before shading, then draw mostly front to back.
+    pub fn list(
+        &mut self,
+        id: u32,
+        page: u32,
+        distance: u64,
+        sides: &Sides,
+        parts: &[(u8, Part)],
+        origin: [f32; 3],
+    ) -> 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;
+        }
+        let room = u64::from(counts[0]) + u64::from(counts[1]);
+        if self.room + room > self.cap {
+            return false;
+        }
+        self.room += room;
+        let band = super::page_runs::band(distance) as u32;
+        if band != self.band {
+            self.band = band;
+            self.call_of.fill(u32::MAX);
+        }
+        let mut record = [id, 0, 0, parts.len() as u32, 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 {
+                self.call_of.resize(key + 1, u32::MAX);
+            }
+            if self.call_of[key] == u32::MAX {
+                self.call_of[key] = self.calls.len() as u32;
+                self.calls.push(ListCall {
+                    group,
+                    page,
+                    first: 0,
+                    count: 0,
+                });
+            }
+            let call = self.call_of[key];
+            self.calls[call as usize].count += counts[k];
+            record[1 + k] = call | part_code(part) << 28;
+        }
+        for (to, o) in record[4..7].iter_mut().zip(origin) {
+            *to = o.to_bits();
+        }
+        self.records.push(record);
+        true
+    }
+
+    /// The sections listed this frame.
+    pub fn listed(&self) -> usize {
+        self.records.len()
+    }
+
+    /// The frame's calls, with their rooms once laid out (`rooms`).
+    pub fn calls(&self) -> &[ListCall] {
+        &self.calls
+    }
+
+    /// Lays the calls' rooms out one after the other from the start of the
+    /// index list: the indices they take.
+    pub fn rooms(&mut self) -> u64 {
+        let mut room = 0u64;
+        for call in &mut self.calls {
+            call.first = room as u32;
+            room += u64::from(call.count);
+        }
+        room
+    }
+
+    /// 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`.
+    pub fn prepare(
+        &mut self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        uploads: &mut super::upload::Uploads,
+        indices: &wgpu::Buffer,
+        eye: Option<[f32; 3]>,
+    ) {
+        self.group = None;
+        let records = self.records.len() as u64;
+        if records == 0 {
+            return;
+        }
+        if records > self.record_capacity {
+            self.record_capacity = records.next_power_of_two();
+            self.record_buffer = buffer(
+                device,
+                "Sections listed",
+                self.record_capacity * 32,
+                wgpu::BufferUsages::STORAGE,
+            );
+        }
+        uploads.buffer(
+            encoder,
+            &self.record_buffer,
+            0,
+            bytemuck::cast_slice(&self.records),
+        );
+        self.call_data.clear();
+        self.call_data
+            .extend(self.calls.iter().map(|c| [0, 1, c.first, 0, 0]));
+        let calls = self.call_data.len() as u64;
+        let indirect = wgpu::BufferUsages::INDIRECT | wgpu::BufferUsages::STORAGE;
+        if calls > self.call_capacity {
+            self.call_capacity = calls.next_power_of_two();
+            self.call_buffer = buffer(
+                device,
+                "Section list calls",
+                self.call_capacity * 20,
+                indirect,
+            );
+        }
+        uploads.buffer(
+            encoder,
+            &self.call_buffer,
+            0,
+            bytemuck::cast_slice(&self.call_data),
+        );
+        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 group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("Section lists"),
+            layout: &self.layout,
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding {
+                        buffer: &self.record_buffer,
+                        offset: 0,
+                        size: wgpu::BufferSize::new(records * 32),
+                    }),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: self.rows.as_entire_binding(),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 2,
+                    resource: self.view.as_entire_binding(),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 3,
+                    resource: self.call_buffer.as_entire_binding(),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 4,
+                    resource: indices.as_entire_binding(),
+                },
+            ],
+        });
+        self.group = Some((group, records as u32));
+    }
+
+    /// Writes the frame's lists (after `prepare`).
+    pub fn dispatch(&self, pass: &mut wgpu::ComputePass<'_>) {
+        let Some((group, records)) = &self.group else {
+            return;
+        };
+        let groups = *records;
+        pass.set_pipeline(&self.pipeline);
+        pass.set_bind_group(0, group, &[]);
+        pass.dispatch_workgroups(groups.min(GROUP_LIMIT), groups.div_ceil(GROUP_LIMIT), 1);
+    }
+
+    /// 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.
+    pub fn draw_indexed(&self, pass: &mut wgpu::RenderPass<'_>, k: usize) {
+        pass.draw_indexed_indirect(&self.call_buffer, k as u64 * 20);
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use crate::mesher::{Part, Sides};
+    use wgpu::util::DeviceExt;
+
+    struct Section {
+        sides: Sides,
+        origin: [f32; 3],
+        start: u32,
+        page: u32,
+        parts: &'static [(u8, Part)],
+    }
+
+    /// Random sorts: some parts empty, grouped or not, split or not, planes in
+    /// and out of a section, each listed in one of the three ways it can be.
+    fn sections(count: usize) -> Vec<Section> {
+        let mut seed = 0x2545_f491_4f6c_dd1du64;
+        let mut next = |n: u32| {
+            seed = seed
+                .wrapping_mul(6_364_136_223_846_793_005)
+                .wrapping_add(1_442_695_040_888_963_407);
+            ((seed >> 33) % u64::from(n)) as u32
+        };
+        (0..count)
+            .map(|_| {
+                let mut ends = [0; 14];
+                let mut end = 0;
+                for e in &mut ends {
+                    if next(3) > 0 {
+                        end += next(40);
+                    }
+                    *e = end;
+                }
+                let planes = std::array::from_fn(|_| next(17) as f32 + next(16) as f32 / 16.0);
+                Section {
+                    sides: Sides::from_parts(ends, next(2) == 0, next(4) > 0, planes),
+                    origin: [0, 1, 2].map(|_| next(129) as f32 - 64.0),
+                    start: next(10_000) * 4,
+                    page: next(3),
+                    parts: match next(3) {
+                        0 => &[(0, Part::All)],
+                        1 => &[(0, Part::Opaque), (1, Part::Cutout)],
+                        _ => &[(1, Part::All)],
+                    },
+                }
+            })
+            .collect()
+    }
+
+    /// The runs the CPU finds for a section's part, as `Gfx` drew them.
+    fn cpu_runs(s: &Section, part: Part, eye: Option<[f32; 3]>) -> Vec<(u32, u32)> {
+        let mut runs = Vec::new();
+        let draw = |first, quads| runs.push((first, quads));
+        match eye {
+            Some(eye) => {
+                let eye = [0, 1, 2].map(|a| eye[a] - s.origin[a]);
+                s.sides.visible_from(eye, eye[1], part, draw);
+            }
+            None => s.sides.every(part, draw),
+        }
+        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>)>> {
+        let (device, queue) = device()?;
+        let mut results = Vec::new();
+        for eye in [
+            Some([3.5f32, 70.25, -12.0]),
+            Some([-40.0, -3.0, 50.5]),
+            None,
+        ] {
+            let mut lists = GpuLists::new(&device);
+            lists.clear(&device);
+            let mut rows = vec![0u32; sections.len() * ROW];
+            for (id, s) in sections.iter().enumerate() {
+                rows[id * ROW] = s.start;
+                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 room = lists.rooms();
+            let init = |contents: &[u8], usage| {
+                device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+                    label: None,
+                    contents,
+                    usage,
+                })
+            };
+            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 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);
+            uploads.finish(&encoder);
+            {
+                let mut pass = encoder.begin_compute_pass(&Default::default());
+                lists.dispatch(&mut pass);
+            }
+            queue.submit([encoder.finish()]);
+            let read_back = |buffer: &wgpu::Buffer, words: u64| -> Vec<u32> {
+                let read = device.create_buffer(&wgpu::BufferDescriptor {
+                    label: None,
+                    size: words * 4,
+                    usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+                    mapped_at_creation: false,
+                });
+                let mut encoder = device.create_command_encoder(&Default::default());
+                encoder.copy_buffer_to_buffer(buffer, 0, &read, 0, words * 4);
+                queue.submit([encoder.finish()]);
+                read.slice(..).map_async(wgpu::MapMode::Read, |_| {});
+                device.poll(wgpu::PollType::wait_indefinitely()).unwrap();
+                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);
+            results.push((eye, lists, calls, words));
+        }
+        Some(results)
+    }
+
+    fn device() -> Option<(wgpu::Device, wgpu::Queue)> {
+        let instance = crate::gfx::test_gpu::instance();
+        let Ok(adapter) = pollster::block_on(instance.request_adapter(&Default::default())) else {
+            eprintln!("no GPU adapter: skipped");
+            return None;
+        };
+        let device = pollster::block_on(adapter.request_device(&Default::default())).ok();
+        if device.is_none() {
+            eprintln!("no GPU device: skipped");
+        }
+        device
+    }
+
+    /// The call section `i`'s `k`th part draws in, as listed.
+    fn call_of(lists: &GpuLists, i: usize, k: usize) -> usize {
+        (lists.records[i][1 + k] & 0x0fff_ffff) as usize
+    }
+
+    /// Each band of distance and page gets a pass's calls of its own, in the
+    /// order sections come (nearest first), and sections of one share them.
+    #[test]
+    fn calls_go_by_band_then_page_nearest_first() {
+        let Some((device, _)) = device() else {
+            return;
+        };
+        let mut lists = GpuLists::new(&device);
+        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]));
+        }
+        let calls: Vec<_> = lists
+            .calls()
+            .iter()
+            .map(|c| (c.group, c.page, c.count))
+            .collect();
+        assert_eq!(
+            calls,
+            [
+                (0, 1, 48),
+                (1, 1, 48),
+                (0, 0, 24),
+                (1, 0, 24),
+                (0, 1, 24),
+                (1, 1, 24),
+                (0, 0, 24),
+                (1, 0, 24)
+            ]
+        );
+    }
+
+    /// `flatten_sides` writes, in each call's room, the indices of exactly the
+    /// runs `Sides::visible_from` finds for each listed section's parts (in
+    /// some order: workgroups take room as they come), and counts them.
+    #[test]
+    fn flatten_sides_writes_the_runs_the_cpu_finds() {
+        let sections = sections(200);
+        let Some(results) = run(&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) {
+                        for k in 0..quads * 6 {
+                            let corner = [0, 1, 2, 0, 2, 3][(k % 6) as usize];
+                            expected[call].push(s.start + first * 4 + (k / 6) * 4 + corner);
+                        }
+                    }
+                }
+            }
+            for (k, call) in lists.calls.iter().enumerate() {
+                let count = calls[k * 5] as usize;
+                assert_eq!(
+                    count,
+                    expected[k].len(),
+                    "call {k}'s index count, eye {eye:?}"
+                );
+                assert_eq!(calls[k * 5 + 2], call.first);
+                let mut got = out[call.first as usize..call.first as usize + count].to_vec();
+                got.sort_unstable();
+                expected[k].sort_unstable();
+                assert_eq!(got, expected[k], "call {k}'s indices, eye {eye:?}");
+            }
+        }
+    }
+}

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

@@ -0,0 +1,178 @@
+// Phase 1's solid section draws made on the GPU (gfx/gpu_lists.rs): per section
+// 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.
+struct View {
+    // The eye relative to the camera, and 1.0 when there is one (`sides_on`);
+    // without, every bucket is drawn (`Sides::every`).
+    eye: vec3<f32>,
+    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.
+struct Record {
+    words: vec4<u32>,
+    origin: vec4<f32>,
+}
+@group(0) @binding(0) var<storage, read> records: array<Record>;
+// Per id, `ROW` words: its first vertex in its page, then `Sides::gpu_row`
+// (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.
+@group(0) @binding(3) var<storage, read_write> calls: array<atomic<u32>>;
+// The flat draws' index list.
+@group(0) @binding(4) var<storage, read_write> out: array<u32>;
+
+const ROW: u32 = 24u;
+const LANES: u32 = 64u;
+const PART_ALL: u32 = 0u;
+const PART_OPAQUE: u32 = 1u;
+const PART_CUTOUT: u32 = 2u;
+const CALL_MASK: u32 = 0x0fffffffu;
+// `sides::MARGIN`.
+const MARGIN: f32 = 0.0625;
+
+// Index `k` of a solid draw: quad k / 6's vertices 0 1 2 0 2 3.
+fn quad(k: u32) -> u32 {
+    let c = k % 6u;
+    return (k / 6u) * 4u + select(select(c - 2u, 0u, c == 3u), c, c < 3u);
+}
+
+fn plane(row: u32, bucket: u32) -> f32 {
+    return bitcast<f32>(rows[row + 15u + bucket]);
+}
+
+// `Sides::visible_from`'s `seen`, buckets in mesh order: +X, +Z, other, +Y, -Z, -X, -Y.
+fn seen(row: u32, bucket: u32, eye: vec3<f32>) -> bool {
+    var faces = true;
+    if view.has_eye != 0.0 {
+        switch bucket {
+            case 0u: { faces = eye.x > plane(row, 0u) - MARGIN; }
+            case 1u: { faces = eye.z > plane(row, 1u) - MARGIN; }
+            case 3u: { faces = eye.y > plane(row, 3u) - MARGIN; }
+            case 4u: { faces = eye.z < plane(row, 4u) + MARGIN; }
+            case 5u: { faces = eye.x < plane(row, 5u) + MARGIN; }
+            case 6u: { faces = eye.y < plane(row, 6u) + MARGIN; }
+            default: {}
+        }
+    }
+    return faces;
+}
+
+// A part's runs (first quad, quads): at most 7, as a run ends at a part left
+// out and there are 14.
+struct Runs {
+    n: u32,
+    quads: u32,
+    runs: array<vec2<u32>, 7>,
+}
+
+// `Sides::runs` for `part` of the mesh at `row`, seen from `eye` (relative to
+// the section): empty parts do not break a run.
+fn part_runs(row: u32, part: u32, eye: vec3<f32>) -> Runs {
+    var r: Runs;
+    let grouped = (rows[row + 22u] & 1u) != 0u;
+    let split = (rows[row + 22u] & 2u) != 0u;
+    var start = 0u;
+    var open = false;
+    var first = 0u;
+    for (var n = 0u; n < 14u; n++) {
+        let end = rows[row + 1u + n];
+        if end == start {
+            continue;
+        }
+        var bucket: u32;
+        var cutout: bool;
+        if grouped {
+            bucket = n % 7u;
+            cutout = n >= 7u;
+        } else {
+            bucket = n / 2u;
+            cutout = n % 2u == 1u;
+        }
+        let wanted = part == PART_ALL || !split
+            || (part == PART_OPAQUE && !cutout) || (part == PART_CUTOUT && cutout);
+        let take = wanted && seen(row, bucket, eye);
+        if take && !open {
+            open = true;
+            first = start;
+        } else if !take && open {
+            r.runs[r.n] = vec2(first, start - first);
+            r.quads += start - first;
+            r.n++;
+            open = false;
+        }
+        start = end;
+    }
+    if open {
+        r.runs[r.n] = vec2(first, start - first);
+        r.quads += start - first;
+        r.n++;
+    }
+    return r;
+}
+
+fn record_index(group: vec3<u32>) -> u32 {
+    return group.y * 65535u + group.x;
+}
+
+// The eye relative to `record`'s section.
+fn eye_of(record: Record) -> vec3<f32> {
+    return view.eye - record.origin.xyz;
+}
+
+fn part_word(record: Record, p: u32) -> u32 {
+    return select(record.words.z, record.words.y, p == 0u);
+}
+
+// Lane 0's runs per part, and where each part's indices start.
+var<workgroup> found: array<Runs, 2>;
+var<workgroup> at: array<u32, 2>;
+
+// One workgroup per record: lane 0 finds the runs and takes room for their
+// indices, then every lane writes them.
+@compute @workgroup_size(64)
+fn flatten_sides(@builtin(workgroup_id) group: vec3<u32>,
+    @builtin(local_invocation_index) lane: u32) {
+    let r = record_index(group);
+    if r >= arrayLength(&records) {
+        return;
+    }
+    let record = records[r];
+    let row = record.words.x * ROW;
+    let parts = record.words.w;
+    if lane == 0u {
+        let eye = eye_of(record);
+        for (var p = 0u; p < parts; p++) {
+            let word = part_word(record, p);
+            let runs = part_runs(row, word >> 28u, eye);
+            let call = (word & CALL_MASK) * 5u;
+            var begin = 0u;
+            if runs.quads > 0u {
+                begin = atomicLoad(&calls[call + 2u]) + atomicAdd(&calls[call], runs.quads * 6u);
+            }
+            found[p] = runs;
+            at[p] = begin;
+        }
+    }
+    workgroupBarrier();
+    let base = rows[row];
+    for (var p = 0u; p < parts; p++) {
+        var to = at[p];
+        for (var k = 0u; k < found[p].n; k++) {
+            let run = found[p].runs[k];
+            let first = base + run.x * 4u;
+            let count = run.y * 6u;
+            for (var i = lane; i < count; i += LANES) {
+                out[to + i] = first + quad(i);
+            }
+            to += count;
+        }
+    }
+}

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

@@ -29,6 +29,11 @@ pub(super) struct Run {
     pub tag: [u32; 3],
 }
 
+/// The band of a section `distance` from the camera (squared, as bits).
+pub(super) fn band(distance: u64) -> u64 {
+    (f64::from_bits(distance).sqrt() / BAND) as u64
+}
+
 #[derive(Default)]
 pub(super) struct PageRuns {
     runs: Vec<Run>,
@@ -58,7 +63,7 @@ impl PageRuns {
         tag: [u32; 3],
     ) {
         if len > 0 {
-            let band = (f64::from_bits(distance).sqrt() / BAND) as u64;
+            let band = band(distance);
             self.runs.push(Run {
                 key: u64::from(group) << GROUP_SHIFT | band << 32 | u64::from(page),
                 start: start as u32,

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

@@ -103,6 +103,22 @@ impl Sides {
         }
     }
 
+    /// A sort with these parts' `ends` and `planes`, for tests.
+    #[cfg(test)]
+    pub(crate) fn from_parts(
+        ends: [u32; 14],
+        grouped: bool,
+        split: bool,
+        planes: [f32; 7],
+    ) -> Self {
+        Self {
+            ends,
+            grouped,
+            split,
+            planes,
+        }
+    }
+
     /// The number of solid quads the buckets hold.
     pub fn quads(&self) -> u32 {
         self.ends[13]
@@ -128,6 +144,26 @@ impl Sides {
         self.runs(seen, part, draw);
     }
 
+    /// How many quads `part` draws facing every way (`every`).
+    pub fn part_quads(&self, part: Part) -> u32 {
+        let mut quads = 0;
+        self.every(part, |_, count| quads += count);
+        quads
+    }
+
+    /// The facts `visible_from` reads, for the GPU to find the same runs
+    /// (gpu_lists.wgsl): the parts' ends, the planes' bits,
+    /// then 1 when grouped plus 2 when split.
+    pub fn gpu_row(&self) -> [u32; 22] {
+        let mut row = [0; 22];
+        row[..14].copy_from_slice(&self.ends);
+        for (to, plane) in row[14..21].iter_mut().zip(self.planes) {
+            *to = plane.to_bits();
+        }
+        row[21] = u32::from(self.grouped) | u32::from(self.split) << 1;
+        row
+    }
+
     /// The runs of `part`'s quads facing every way.
     pub fn every(&self, part: Part, draw: impl FnMut(u32, u32)) {
         self.runs(|_| true, part, draw);