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.rs

@@ -12,7 +12,7 @@ use crate::{
     assets::{Assets, RenderAssets},
     biome::Biomes,
     blocks::{BlockFacts, RenderShape},
-    chunk::{Chunk, MIN_Y},
+    chunk::{Chunk, Light, MIN_Y},
     graphics::LeafCullingMode,
     models::catalog::BakedFace,
 };
@@ -505,6 +505,10 @@ pub struct MeshFacts {
 }
 pub struct MeshTables {
     info: Vec<Info>,
+    /// `Info::solid` by state, packed apart: the opacity mask of a section
+    /// reads 8,000 states, and a few bytes each stay in cache where `Info`
+    /// does not.
+    solid: Vec<bool>,
     pub visibility: FaceVisibility,
     water_still: u16,
     water_flow: u16,
@@ -673,6 +677,7 @@ impl MeshTables {
             });
         }
         Ok(Self {
+            solid: info.iter().map(|i| i.solid).collect(),
             info,
             visibility,
             water_still: layer("minecraft:block/water_still")?,
@@ -842,43 +847,66 @@ impl Padded {
     }
     fn fill(&mut self, region: &[[Option<&Chunk>; 3]; 3], sy: i32, own: &mut [u32; 4096]) {
         self.blocks.resize(P * P * P, 0);
-        self.blocks.fill(0);
         self.light.resize(P * P * P, 0);
         let (blocks, light) = (&mut self.blocks, &mut self.light);
         let base_y = region[1][1].map_or(MIN_Y, Chunk::min_y) + sy * 16;
-        // The section itself is unpacked in one pass; the border cells of the
-        // neighbouring sections are read one by one.
+        // The section itself is unpacked in one pass and copied row by row;
+        // the border cells of the neighbouring sections are read one by one,
+        // each run of them from a section found once.
         own.fill(0);
         let center = region[1][1].map(|c| &c.sections[sy as usize]);
         if let Some(section) = center {
             section.unpack(own);
         }
+        // The x runs of a padded row by the chunk column they fall in.
+        const RUNS: [(usize, i32, i32); 3] = [(0, -B, 0), (1, 0, 16), (2, 16, 16 + B)];
         for y in -B..16 + B {
             let wy = base_y + y;
             for z in -B..16 + B {
-                for x in -B..16 + B {
-                    let chunk = region[((z + 16) >> 4) as usize][((x + 16) >> 4) as usize];
-                    let i = pi(x, y, z);
-                    if let Some(c) = chunk {
-                        blocks[i] =
-                            if (0..16).contains(&x) && (0..16).contains(&y) && (0..16).contains(&z)
-                            {
-                                own[(y << 8 | z << 4 | x) as usize]
-                            } else {
-                                c.get(x & 15, wy, z & 15)
-                            };
-                        light[i] = c.light(x & 15, wy, z & 15);
-                    } else if let Some(c) = region[1][1] {
-                        // A neighbour not loaded continues this chunk's edge
-                        // column. As lit air, the outermost chunks' water drew
-                        // a wall toward it (fluids show both sides), a light
-                        // sheet in the water where far terrain carries on.
-                        // Vanilla draws no chunk without its neighbours.
-                        let (ex, ez) = (x.clamp(0, 15), z.clamp(0, 15));
-                        blocks[i] = c.get(ex, wy, ez);
-                        light[i] = c.light(ex, wy, ez);
+                let row = pi(-B, y, z);
+                let rz = ((z + 16) >> 4) as usize;
+                for (rx, x0, x1) in RUNS {
+                    let (start, end) = (row + (x0 + B) as usize, row + (x1 + B) as usize);
+                    let (blocks, light) = (&mut blocks[start..end], &mut light[start..end]);
+                    let Some(c) = region[rz][rx] else {
+                        match region[1][1] {
+                            // A neighbour not loaded continues this chunk's edge
+                            // column. As lit air, the outermost chunks' water drew
+                            // a wall toward it (fluids show both sides), a light
+                            // sheet in the water where far terrain carries on.
+                            // Vanilla draws no chunk without its neighbours.
+                            Some(c) => {
+                                let ez = z.clamp(0, 15);
+                                for (k, x) in (x0..x1).enumerate() {
+                                    let ex = x.clamp(0, 15);
+                                    blocks[k] = c.get(ex, wy, ez);
+                                    light[k] = c.light(ex, wy, ez);
+                                }
+                            }
+                            None => {
+                                blocks.fill(0);
+                                light.fill(0xF0);
+                            }
+                        }
+                        continue;
+                    };
+                    // `Chunk::get` and `Chunk::light` with the section found once.
+                    let s = (wy - c.min_y()) >> 4;
+                    let Some(section) = c.sections.get(s as usize).filter(|_| s >= 0) else {
+                        blocks.fill(0);
+                        light.fill(if s < 0 { 0 } else { 0xF0 });
+                        continue;
+                    };
+                    let at = crate::chunk::index(x0, wy, z);
+                    if rx == 1 && rz == 1 && (0..16).contains(&y) {
+                        let own_row = (y << 8 | z << 4) as usize;
+                        blocks.copy_from_slice(&own[own_row..own_row + 16]);
                     } else {
-                        light[i] = 0xF0;
+                        section.blocks.get_run(at, blocks);
+                    }
+                    match &section.light {
+                        Light::Uniform(v) => light.fill(*v),
+                        Light::Full(all) => light.copy_from_slice(&all[at..at + light.len()]),
                     }
                 }
             }
@@ -894,6 +922,28 @@ pub struct MeshScratch {
     own: Option<Box<[u32; 4096]>>,
     visibility: Vec<usize>,
     leaf: leaf_culling::LeafMask,
+    /// Per padded row (`pi(-B, y, z) / P`), bit `x + B` set where the block
+    /// is `isSolidRender`.
+    opaque: Vec<u32>,
+}
+/// Fills `rows` with the padded section's opacity, one bit per block (`MeshScratch::opaque`).
+fn opaque_rows(pad: &Padded, solid: &[bool], rows: &mut Vec<u32>) {
+    rows.clear();
+    rows.extend(pad.blocks.as_chunks::<P>().0.iter().map(|row| {
+        row.iter().enumerate().fold(0u32, |bits, (x, &id)| {
+            bits | u32::from(solid[id as usize]) << x
+        })
+    }));
+}
+/// Whether the block at `(x, y, z)` (section-local) and its six neighbours
+/// are all `isSolidRender`: none of its faces can show.
+#[inline]
+fn hidden(rows: &[u32], x: i32, y: i32, z: i32) -> bool {
+    let row = |y: i32, z: i32| rows[((y + B) as usize) * P + (z + B) as usize];
+    let r = row(y, z);
+    let around =
+        r & r << 1 & r >> 1 & row(y - 1, z) & row(y + 1, z) & row(y, z - 1) & row(y, z + 1);
+    around >> (x + B) & 1 != 0
 }
 #[cfg(test)]
 mod leaf_culling_tests;
@@ -990,6 +1040,12 @@ enum Resolver {
 /// layer plus the blend radius around it, and the blended colours per column.
 struct Layer {
     zoomed: Vec<u16>,
+    /// Each resolver's colour of every cell of `zoomed`, filled on first use:
+    /// a cell is in the blend square of up to `(2r+1)^2` columns.
+    cells: [Vec<u32>; 4],
+    /// `BiomeZoom::corner_jitter` of the quart cells `zoomed` reaches, by
+    /// their offset from the first: each is shared by up to 16 cells of the layer.
+    jitter: Vec<Option<[[f64; 3]; 8]>>,
     colors: [[u32; 256]; 4],
 }
 /// `ClientLevel.calculateBlockTint` behind a per-layer `BlockTintCache`: each
@@ -1041,6 +1097,16 @@ impl<'a> Tints<'a> {
         [x, y, z]: [i32; 3],
     ) -> u16 {
         let q = zoom.quart_at_block(BlockPos::new(cx * 16 + x, base_y + y, cz * 16 + z));
+        Self::quart_biome(region, plains, [cx, cz], q)
+    }
+    /// The biome of quart cell `q` in the loaded chunks around chunk `(cx, cz)`, plains where
+    /// there is none.
+    fn quart_biome(
+        region: &[[Option<&Chunk>; 3]; 3],
+        plains: u16,
+        [cx, cz]: [i32; 2],
+        q: [i32; 3],
+    ) -> u16 {
         let (rx, rz) = ((q[0] >> 2) - cx + 1, (q[2] >> 2) - cz + 1);
         if (0..3).contains(&rx) && (0..3).contains(&rz) {
             region[rz as usize][rx as usize].map_or(plains, |c| {
@@ -1057,6 +1123,8 @@ impl<'a> Tints<'a> {
         let layer = self.layers[y as usize].get_or_insert_with(|| {
             Box::new(Layer {
                 zoomed: vec![NO_BIOME; dim * dim],
+                cells: Default::default(),
+                jitter: Vec::new(),
                 colors: [[UNSET; 256]; 4],
             })
         });
@@ -1065,21 +1133,45 @@ impl<'a> Tints<'a> {
         if cached != UNSET {
             return cached;
         }
-        let [cx, cz, _] = self.origin;
-        let (bx, bz) = (cx * 16, cz * 16);
+        let [cx, cz, base_y] = self.origin;
+        let (bx, by, bz) = (cx * 16, base_y + y, cz * 16);
         let (mut red, mut green, mut blue) = (0, 0, 0);
+        let cells = &mut layer.cells[resolver as usize];
+        if cells.is_empty() {
+            cells.resize(dim * dim, UNSET);
+        }
         for dz in -r..=r {
             for dx in -r..=r {
                 let cell = ((z + dz + r) as usize) * dim + (x + dx + r) as usize;
+                let known = cells[cell];
+                if known != UNSET {
+                    red += known >> 16 & 255;
+                    green += known >> 8 & 255;
+                    blue += known & 255;
+                    continue;
+                }
                 let mut biome = layer.zoomed[cell];
                 if biome == NO_BIOME {
-                    biome = Self::biome_id(
-                        &self.zoom,
-                        self.region,
-                        self.plains,
-                        self.origin,
-                        [x + dx, y, z + dz],
-                    );
+                    // BiomeManager.getBiome over the loaded chunks (plains
+                    // where there is none): `quart_at_block`, its cell's
+                    // corner jitter mixed once per layer.
+                    let block = BlockPos::new(bx + x + dx, by, bz + z + dz);
+                    let base = [block.x, block.y, block.z].map(|v| v.wrapping_sub(2) >> 2);
+                    let first = [bx - r, bz - r].map(|v| v.wrapping_sub(2) >> 2);
+                    let span = dim / 4 + 2;
+                    let at = (base[2] - first[1]) as usize * span + (base[0] - first[0]) as usize;
+                    if layer.jitter.is_empty() {
+                        layer.jitter.resize(span * span, None);
+                    }
+                    let jitter =
+                        layer.jitter[at].get_or_insert_with(|| self.zoom.corner_jitter(base));
+                    let corner = BiomeZoom::select_corner(block, jitter);
+                    let q = [
+                        base[0].wrapping_add((corner >> 2 & 1) as i32),
+                        base[1].wrapping_add((corner >> 1 & 1) as i32),
+                        base[2].wrapping_add((corner & 1) as i32),
+                    ];
+                    biome = Self::quart_biome(self.region, self.plains, [cx, cz], q);
                     layer.zoomed[cell] = biome;
                 }
                 let c = match resolver {
@@ -1091,6 +1183,7 @@ impl<'a> Tints<'a> {
                     Resolver::Dry => self.biomes.get(biome).dry_foliage,
                     Resolver::Water => self.biomes.get(biome).water,
                 };
+                cells[cell] = c;
                 red += c >> 16 & 255;
                 green += c >> 8 & 255;
                 blue += c & 255;
@@ -1163,6 +1256,8 @@ impl Mesher<'_> {
         scratch.padded.fill(region, sy, own);
         let pad = &scratch.padded;
         out.visibility = visibility_reusing(pad, &self.tables.info, &mut scratch.visibility);
+        opaque_rows(pad, &self.tables.solid, &mut scratch.opaque);
+        let opaque = &scratch.opaque;
         let base_y = center.min_y() + sy * 16;
         if self.leaf_culling.is_interior() {
             scratch
@@ -1195,12 +1290,7 @@ impl Mesher<'_> {
                         out.lights
                             .push(([center.x * 16 + x, base_y + y, center.z * 16 + z], id));
                     }
-                    if state.solid
-                        && Direction::ALL.iter().all(|d| {
-                            let [dx, dy, dz] = d.normal();
-                            info[pad.blocks[pi(x + dx, y + dy, z + dz)] as usize].solid
-                        })
-                    {
+                    if hidden(opaque, x, y, z) {
                         continue;
                     }
                     if state.fluid.kind() != FluidKind::Empty {

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);
@@ -210,6 +246,29 @@ pub fn sort(solid: &mut Vec<TerrainVertex>) -> Sides {
 
 /// `sort`, with all the opaque parts first when `grouped`.
 fn sort_parts(solid: &mut Vec<TerrainVertex>, grouped: bool) -> Sides {
+    SORT_SCRATCH.with_borrow_mut(|scratch| sort_with(solid, grouped, scratch))
+}
+
+/// `sort`'s buffers, kept by each meshing thread: a section's quads went
+/// through fresh vectors, and the sorted copy replaced the mesh's own
+/// (recycled) buffer.
+#[derive(Default)]
+struct SortScratch {
+    slots: Vec<u8>,
+    keys: Vec<u64>,
+    sorted: Vec<[TerrainVertex; 4]>,
+}
+thread_local! {
+    static SORT_SCRATCH: std::cell::RefCell<SortScratch> = Default::default();
+}
+
+/// `f32::total_cmp`'s order as an unsigned integer's.
+fn total_order_key(v: f32) -> u32 {
+    let bits = v.to_bits() as i32;
+    (bits ^ (((bits >> 31) as u32) >> 1) as i32) as u32 ^ 0x8000_0000
+}
+
+fn sort_with(solid: &mut Vec<TerrainVertex>, grouped: bool, scratch: &mut SortScratch) -> Sides {
     let (quads, rest) = solid.as_chunks::<4>();
     debug_assert!(rest.is_empty(), "solid vertices come in quads");
     let cutout = |q: &[TerrainVertex; 4]| q[0].layer & OPAQUE_LAYER == 0;
@@ -218,15 +277,14 @@ fn sort_parts(solid: &mut Vec<TerrainVertex>, grouped: bool) -> Sides {
         true => b + 7 * usize::from(cutout),
         false => b * 2 + usize::from(cutout),
     };
-    let slots: Vec<u8> = quads
-        .iter()
-        .map(|q| position(bucket(q), cutout(q)) as u8)
-        .collect();
+    let slots = &mut scratch.slots;
+    slots.clear();
+    slots.extend(quads.iter().map(|q| position(bucket(q), cutout(q)) as u8));
     let has_cutout = quads.iter().any(cutout);
     let mut counts = [0u32; 7];
     let mut slot_counts = [0u32; 14];
     let mut planes = [0f32; 7];
-    for (quad, &part) in quads.iter().zip(&slots) {
+    for (quad, &part) in quads.iter().zip(slots.iter()) {
         slot_counts[usize::from(part)] += 1;
         let b = slot(grouped, usize::from(part)).0;
         if let Bucket::Side(axis, positive) = BUCKETS[b] {
@@ -246,33 +304,34 @@ fn sort_parts(solid: &mut Vec<TerrainVertex>, grouped: bool) -> Sides {
         *end = total;
     }
     if has_cutout || counts.iter().filter(|&&c| c > 0).count() > 1 {
-        let mut next: [u32; 14] = std::array::from_fn(|s| slot_ends[s] - slot_counts[s]);
-        let mut sorted = vec![quads[0]; quads.len()];
-        for (quad, &part) in quads.iter().zip(&slots) {
-            let at = &mut next[usize::from(part)];
-            sorted[*at as usize] = *quad;
-            *at += 1;
-        }
-        if has_cutout {
-            let mut start = 0;
-            for (s, &end) in slot_ends.iter().enumerate() {
-                let end = end as usize;
-                let slice = &mut sorted[start..end];
-                if let Bucket::Side(axis, positive) = BUCKETS[slot(grouped, s).0] {
-                    slice.sort_by(|a, b| {
-                        let a = a[0].pos[axis];
-                        let b = b[0].pos[axis];
-                        if positive {
-                            b.total_cmp(&a)
-                        } else {
-                            a.total_cmp(&b)
+        // One integer key per quad: its part, then (cutout meshes, side
+        // buckets) its plane from the outside in, then its index, which keeps
+        // the order of equal quads as a stable sort does.
+        let keys = &mut scratch.keys;
+        keys.clear();
+        keys.extend(
+            quads
+                .iter()
+                .zip(slots.iter())
+                .enumerate()
+                .map(|(i, (quad, &part))| {
+                    let plane = match BUCKETS[slot(grouped, usize::from(part)).0] {
+                        Bucket::Side(axis, positive) if has_cutout => {
+                            let key = total_order_key(quad[0].pos[axis]);
+                            if positive { !key } else { key }
                         }
-                    });
-                }
-                start = end;
-            }
-        }
-        *solid = sorted.into_flattened();
+                        _ => 0,
+                    };
+                    u64::from(part) << 56 | u64::from(plane) << 24 | i as u64
+                }),
+        );
+        debug_assert!(quads.len() < 1 << 24);
+        keys.sort_unstable();
+        let sorted = &mut scratch.sorted;
+        sorted.clear();
+        sorted.extend(keys.iter().map(|&k| quads[(k & 0xFF_FFFF) as usize]));
+        solid.clear();
+        solid.extend_from_slice(sorted.as_flattened());
     }
     Sides {
         ends: slot_ends,
@@ -286,6 +345,74 @@ fn sort_parts(solid: &mut Vec<TerrainVertex>, grouped: bool) -> Sides {
 mod tests {
     use super::*;
 
+    /// The sort as it was before the integer keys: quads moved to their
+    /// parts in order, then a stable sort by plane in each side bucket.
+    fn reference(solid: &[TerrainVertex], grouped: bool) -> Vec<TerrainVertex> {
+        let quads = solid.as_chunks::<4>().0;
+        let has_cutout = quads.iter().any(|q| q[0].layer & OPAQUE_LAYER == 0);
+        let mut sorted = Vec::new();
+        for n in 0..14 {
+            let (b, cutout) = slot(grouped, n);
+            let mut part: Vec<[TerrainVertex; 4]> = quads
+                .iter()
+                .filter(|q| bucket(q) == b && (q[0].layer & OPAQUE_LAYER == 0) == cutout)
+                .copied()
+                .collect();
+            if let (Bucket::Side(axis, positive), true) = (BUCKETS[b], has_cutout) {
+                part.sort_by(|a, b| {
+                    let (a, b) = (a[0].pos[axis], b[0].pos[axis]);
+                    if positive {
+                        b.total_cmp(&a)
+                    } else {
+                        a.total_cmp(&b)
+                    }
+                });
+            }
+            sorted.extend(part);
+        }
+        sorted.into_flattened()
+    }
+
+    #[test]
+    fn integer_keys_sort_as_the_stable_sort_did() {
+        let mut state = 0x2545_F491_4F6C_DD1Du64;
+        let mut next = move || {
+            state ^= state << 13;
+            state ^= state >> 7;
+            state ^= state << 17;
+            state
+        };
+        for round in 0..100 {
+            let grouped = round % 2 == 1;
+            let mut solid = Vec::new();
+            for _ in 0..(next() % 300) {
+                let origin = [0, 1, 2].map(|_| (next() % 34) as f32 / 2.0 - 0.5);
+                // Some negative and signed-zero planes too.
+                let origin = if next() % 9 == 0 {
+                    origin.map(|v| -v)
+                } else {
+                    origin
+                };
+                let layer = if round % 3 == 0 {
+                    OPAQUE_LAYER
+                } else {
+                    (next() % 2) as u16 * OPAQUE_LAYER
+                };
+                let faces = cube(origin, layer);
+                let face = (next() % 6) as usize;
+                solid.extend_from_slice(&faces[face * 4..face * 4 + 4]);
+            }
+            let expected = reference(&solid, grouped);
+            let sides = sort_parts(&mut solid, grouped);
+            assert_eq!(
+                bytemuck::cast_slice::<_, u8>(&solid),
+                bytemuck::cast_slice::<_, u8>(&expected),
+                "round {round}"
+            );
+            assert_eq!(sides.quads() as usize * 4, solid.len());
+        }
+    }
+
     fn quad(corners: [[f32; 3]; 4], layer: u16) -> [TerrainVertex; 4] {
         corners.map(|p| TerrainVertex::new(p, [0.0; 2], layer, 0, 0))
     }

crates/horizon-core/src/paletted.rs

@@ -272,6 +272,23 @@ impl PalettedContainer {
             Palette::Global => self.raw(i) as u32,
         }
     }
+    /// `get` of the entries `start..start + out.len()`, the palette matched once.
+    #[inline]
+    pub fn get_run(&self, start: usize, out: &mut [u32]) {
+        match &self.palette {
+            Palette::Single(value) => out.fill(*value),
+            Palette::List(list) => {
+                for (k, cell) in out.iter_mut().enumerate() {
+                    *cell = list[self.raw(start + k)];
+                }
+            }
+            Palette::Global => {
+                for (k, cell) in out.iter_mut().enumerate() {
+                    *cell = self.raw(start + k) as u32;
+                }
+            }
+        }
+    }
     /// The value when every entry is equal, without the palette having to be
     /// minimal.
     pub fn uniform(&self) -> Option<u32> {
@@ -736,6 +753,26 @@ mod tests {
         *state
     }
     #[test]
+    fn runs_read_what_get_reads() {
+        let mut state = 0x9E37_79B9_7F4A_7C15;
+        // One value, a list palette, and the global palette.
+        for distinct in [1u64, 5, 3000] {
+            let values: Vec<u32> = (0..4096)
+                .map(|_| (rng(&mut state) % distinct) as u32)
+                .collect();
+            let container = PalettedContainer::from_values(BLOCKS, &values);
+            for _ in 0..200 {
+                let start = (rng(&mut state) % 4096) as usize;
+                let len = (rng(&mut state) % 20) as usize;
+                let len = len.min(4096 - start);
+                let mut run = vec![u32::MAX; len];
+                container.get_run(start, &mut run);
+                let expected: Vec<u32> = (start..start + len).map(|i| container.get(i)).collect();
+                assert_eq!(run, expected, "{distinct} values, {start}+{len}");
+            }
+        }
+    }
+    #[test]
     fn containers_stored_alike_are_the_same() {
         let values: Vec<u32> = (0..4096).map(|i| (i % 7) as u32).collect();
         let a = PalettedContainer::from_values(BLOCKS, &values);

crates/horizon-world/src/block_collision_scan.rs

@@ -79,6 +79,48 @@ pub struct BlockCollisionHit {
     pub shape: VoxelShape,
 }
 
+#[derive(Clone, Copy, Debug)]
+enum CellError {
+    InvalidBounds,
+    TooLarge,
+}
+
+/// The cursor a scan of `bounds` walks: its padded start, the box's extents in
+/// cells and the cell count. `new` and [`scan_cells`] share it, so a caller
+/// that sizes a budget from a box counts exactly what the scan will visit.
+fn cells(bounds: Aabb) -> Result<(BlockPos, [i32; 3], usize), CellError> {
+    let min = bounds.min_position();
+    let max = bounds.max_position();
+    if !min.is_finite() || !max.is_finite() {
+        return Err(CellError::InvalidBounds);
+    }
+    let start = BlockPos::containing(Vec3::new(min.x - 1.0e-7, min.y - 1.0e-7, min.z - 1.0e-7))
+        .offset(-1, -1, -1);
+    let end = BlockPos::containing(Vec3::new(max.x + 1.0e-7, max.y + 1.0e-7, max.z + 1.0e-7))
+        .offset(1, 1, 1);
+    let sizes = [
+        end.x.wrapping_sub(start.x).wrapping_add(1),
+        end.y.wrapping_sub(start.y).wrapping_add(1),
+        end.z.wrapping_sub(start.z).wrapping_add(1),
+    ];
+    if sizes.iter().any(|size| *size <= 0) {
+        return Err(CellError::TooLarge);
+    }
+    let count = sizes
+        .into_iter()
+        .try_fold(1u64, |n, size| {
+            n.checked_mul(size as u64).filter(|n| *n <= i32::MAX as u64)
+        })
+        .ok_or(CellError::TooLarge)? as usize;
+    Ok((start, sizes, count))
+}
+
+/// The cells a scan of `bounds` visits, or `None` for a box `new` refuses
+/// (`InvalidBounds`, `TooLarge`): the most a budget can usefully be.
+pub fn scan_cells(bounds: Aabb) -> Option<usize> {
+    cells(bounds).ok().map(|(_, _, count)| count)
+}
+
 pub struct BlockCollisionScan<'w, 'c, W: BlockCollisionWorld> {
     world: &'w mut W,
     context: CollisionContext<'c>,
@@ -110,29 +152,10 @@ impl<'w, 'c, W: BlockCollisionWorld> BlockCollisionScan<'w, 'c, W> {
         only_suffocating: bool,
         candidate_budget: usize,
     ) -> Result<Self, CollisionScanError<W::Error>> {
-        let min = bounds.min_position();
-        let max = bounds.max_position();
-        if !min.is_finite() || !max.is_finite() {
-            return Err(CollisionScanError::InvalidBounds);
-        }
-        let start = BlockPos::containing(Vec3::new(min.x - 1.0e-7, min.y - 1.0e-7, min.z - 1.0e-7))
-            .offset(-1, -1, -1);
-        let end = BlockPos::containing(Vec3::new(max.x + 1.0e-7, max.y + 1.0e-7, max.z + 1.0e-7))
-            .offset(1, 1, 1);
-        let sizes = [
-            end.x.wrapping_sub(start.x).wrapping_add(1),
-            end.y.wrapping_sub(start.y).wrapping_add(1),
-            end.z.wrapping_sub(start.z).wrapping_add(1),
-        ];
-        if sizes.iter().any(|size| *size <= 0) {
-            return Err(CollisionScanError::TooLarge);
-        }
-        let count = sizes
-            .into_iter()
-            .try_fold(1u64, |n, size| {
-                n.checked_mul(size as u64).filter(|n| *n <= i32::MAX as u64)
-            })
-            .ok_or(CollisionScanError::TooLarge)? as usize;
+        let (start, sizes, count) = cells(bounds).map_err(|error| match error {
+            CellError::InvalidBounds => CollisionScanError::InvalidBounds,
+            CellError::TooLarge => CollisionScanError::TooLarge,
+        })?;
         Ok(Self {
             world,
             context,

crates/horizon-world/src/world_movement.rs

@@ -1,7 +1,9 @@
 //! Entity.collide collection order and step resolution over a collision world.
 //! Travel, velocity, grounded flags and inside-block effects remain caller work.
 use crate::{
-    block_collision_scan::{BlockCollisionScan, BlockCollisionWorld, CollisionScanError},
+    block_collision_scan::{
+        BlockCollisionScan, BlockCollisionWorld, CollisionScanError, scan_cells,
+    },
     collision_context::CollisionContext,
     entity_obstruction::EntityIdentity,
 };
@@ -56,6 +58,19 @@ where
             entity_collision_bounds(request.movement, request.bounds, request.max_step),
         )
         .map_err(CollisionScanError::World)?;
+    // Vanilla walks the whole box a movement expands to with no candidate
+    // counter, so a movement big enough to pass the caller's budget is slow
+    // there, never fatal (#6124: a legal movement_speed of 400 crossed more
+    // than a million cells on a diagonal and aborted the game). The budget
+    // only has to cover the cells both queries of a step can visit: the
+    // sweep, and the step's own box raised by max_step.
+    let sweep = request
+        .bounds
+        .expand_towards(request.movement)
+        .expand_towards(Vec3::new(0.0, f64::from(request.max_step.max(0.0)), 0.0));
+    let candidate_budget = scan_cells(sweep).map_or(request.candidate_budget, |cells| {
+        cells.max(request.candidate_budget)
+    });
     try_resolve_step_movement(
         request.movement,
         request.bounds,
@@ -69,13 +84,9 @@ where
             {
                 shapes.push(border);
             }
-            for hit in BlockCollisionScan::new(
-                world,
-                request.context,
-                bounds,
-                false,
-                request.candidate_budget,
-            )? {
+            for hit in
+                BlockCollisionScan::new(world, request.context, bounds, false, candidate_budget)?
+            {
                 shapes.push(hit?.shape);
             }
             Ok(shapes)

crates/horizon-world/tests/world_movement.rs

@@ -0,0 +1,131 @@
+//! The candidate budget bounds a scan, never a movement: a movement big enough
+//! to pass the caller's budget is walked in full, as vanilla's uncounted
+//! collision walk does (#6124).
+use horizon_core::{
+    geometry::{Aabb, Vec3, VoxelShape},
+    position::{BlockPos, ChunkPos},
+};
+use horizon_world::{
+    block_collision_scan::{BlockCollisionWorld, CollisionBlock, scan_cells},
+    collision_context::CollisionContext,
+    entity_obstruction::EntityIdentity,
+    registry::BlockStateId,
+    world_movement::{EntityMovementWorld, MovementRequest, resolve_world_movement},
+};
+use std::io;
+
+/// Open air: every cell is reachable and has no shape, so a scan visits all of
+/// them and the element count is the walk it performed.
+struct OpenAir {
+    cells: usize,
+}
+impl BlockCollisionWorld for OpenAir {
+    type Error = io::Error;
+    type Chunk = ();
+    fn collision_chunk(&mut self, _: ChunkPos) -> io::Result<Option<()>> {
+        Ok(Some(()))
+    }
+    fn collision_block(&mut self, _: (), _: BlockPos) -> io::Result<CollisionBlock> {
+        self.cells += 1;
+        Ok(CollisionBlock {
+            state: BlockStateId(0),
+            large_collision_shape: false,
+            moving_piston: false,
+            empty_shape: true,
+        })
+    }
+    fn is_suffocating(&mut self, _: (), _: CollisionBlock, _: BlockPos) -> io::Result<bool> {
+        unreachable!()
+    }
+    fn context_shape(
+        &mut self,
+        _: CollisionBlock,
+        _: BlockPos,
+        _: &CollisionContext<'_>,
+    ) -> io::Result<VoxelShape> {
+        unreachable!()
+    }
+}
+impl EntityMovementWorld for OpenAir {
+    type Error = io::Error;
+    fn entity_collisions(&mut self, _: EntityIdentity, _: Aabb) -> io::Result<Vec<VoxelShape>> {
+        Ok(Vec::new())
+    }
+    fn movement_border(&mut self, _: EntityIdentity, _: Aabb) -> io::Result<Option<VoxelShape>> {
+        Ok(None)
+    }
+}
+
+fn player() -> Aabb {
+    Aabb::new(Vec3::new(0.5, 64.0, 0.5), Vec3::new(1.1, 65.8, 1.1))
+}
+
+fn resolve(movement: Vec3, candidate_budget: usize) -> (Vec3, OpenAir) {
+    let mut world = OpenAir { cells: 0 };
+    let resolved = resolve_world_movement(
+        &mut world,
+        MovementRequest {
+            entity: EntityIdentity(0),
+            context: CollisionContext::empty(false),
+            bounds: player(),
+            movement,
+            max_step: 0.6,
+            on_ground: true,
+            candidate_budget,
+        },
+    )
+    .expect("a movement is never aborted by the candidate budget");
+    (resolved, world)
+}
+
+#[test]
+fn a_diagonal_sweep_past_the_budget_is_walked_in_full() {
+    // 40 blocks each way pads to about 42 x 4 x 42 cells, far past the budget.
+    let movement = Vec3::new(40.0, 0.0, 40.0);
+    let sweep = scan_cells(player().expand_towards(movement)).expect("a scannable box");
+    assert!(sweep > 1_000, "the sweep is {sweep} cells");
+    let (resolved, world) = resolve(movement, 10);
+    assert_eq!(resolved, movement);
+    assert!(
+        world.cells > 0,
+        "the whole sweep was walked: {} cells",
+        world.cells
+    );
+}
+
+#[test]
+fn scan_cells_covers_what_the_scan_visits() {
+    // The count includes the corners the cursor skips, so it is an upper bound
+    // on the walk that is never short of it by more than the eight corners.
+    for movement in [
+        Vec3::new(0.0, 0.0, 0.0),
+        Vec3::new(1.0, -0.5, 2.0),
+        Vec3::new(40.0, 3.0, 40.0),
+    ] {
+        let (_, world) = resolve(movement, 0);
+        let sweep = scan_cells(player().expand_towards(movement)).expect("a scannable box");
+        if movement.length_squared() == 0.0 {
+            // A zero movement is never scanned.
+            assert_eq!(world.cells, 0);
+        } else {
+            assert!(world.cells <= sweep, "movement {movement:?}");
+            assert!(sweep - world.cells <= 8, "movement {movement:?}");
+        }
+    }
+}
+
+#[test]
+fn ordinary_movements_are_unchanged() {
+    // Resolved movement does not depend on the budget once it covers the sweep.
+    for movement in [
+        Vec3::new(0.2, -0.0784, 0.1),
+        Vec3::new(1.0, 0.0, -3.0),
+        Vec3::new(0.0, -0.5, 0.0),
+    ] {
+        assert_eq!(
+            resolve(movement, 1_000_000).0,
+            resolve(movement, 1).0,
+            "movement {movement:?}"
+        );
+    }
+}