Merge remote-tracking branch 'origin/perf/recipes-read-once-2' into perf/sync-data-once

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

Cargo.lock

@@ -1733,6 +1733,7 @@ dependencies = [
  "flate2",
  "horizon-core",
  "libm",
+ "log",
  "lz4_flex",
  "postcard",
  "serde",

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

@@ -112,7 +112,6 @@ mod cull;
 mod end_sky;
 mod flat;
 mod frame_gen;
-mod gpu_lists;
 mod hdr;
 mod hdr_icc;
 mod hdr_stats;
@@ -358,8 +357,6 @@ 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.
@@ -450,19 +447,12 @@ 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 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,
-);
+/// copies into `encoder` and, in the section arena, the ranges' flat draw ids
+/// (`flat`) along. False when the other pages have no room left.
 fn move_off_page<'a>(
     arena: &mut arena::Arena,
     encoder: &mut wgpu::CommandEncoder,
-    mut ids: Option<MeshIds<'_>>,
+    mut flat: Option<(&mut flat::Flat, &wgpu::Device, &mut upload::Uploads)>,
     meshes: impl Iterator<Item = &'a mut SectionGpu>,
     page: u32,
     part: Part,
@@ -484,14 +474,9 @@ fn move_off_page<'a>(
         };
         arena.copy(encoder, (page, start), to, vertices);
         arena.release(page, start, 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);
-            }
+        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));
         }
         part.set(s, to);
         *budget = budget.saturating_sub(bytes);
@@ -792,8 +777,6 @@ 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>,
@@ -3514,7 +3497,6 @@ 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,
@@ -3629,7 +3611,6 @@ impl Renderer {
             solid_cutout,
             translucent,
             flat_on: flat.is_some(),
-            gpu_lists,
             cull_on: cull.is_some(),
             sides_on: true,
             cull,
@@ -4176,11 +4157,6 @@ 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| {
@@ -4193,15 +4169,6 @@ 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 [
@@ -4232,7 +4199,11 @@ impl Renderer {
             start,
             translucent_at,
             solid,
-            sides,
+            sides: if mesh.sides.quads() == solid {
+                mesh.sides
+            } else {
+                Sides::unsorted(solid)
+            },
             translucent,
             sort: translucency::SortState::new(
                 mesh.translucent
@@ -4244,7 +4215,6 @@ impl Renderer {
             ),
             sorted_indices,
             flat_id,
-            list_id,
             fading: None,
             leaving: None,
             animated_layers,
@@ -4317,9 +4287,6 @@ 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);
@@ -4362,15 +4329,13 @@ impl Renderer {
             .sections
             .values_mut()
             .chain(self.lods.values_mut().flat_map(|tile| tile.meshes_mut()));
-        let ids = Some((
-            self.flat.as_mut(),
-            self.gpu_lists.as_mut(),
-            &self.device,
-            &mut self.section_uploads,
-        ));
+        let flat = self
+            .flat
+            .as_mut()
+            .map(|f| (f, &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, ids, meshes, page, part, &mut budget) {
+        if !move_off_page(arena, encoder, flat, meshes, page, part, &mut budget) {
             // The other pages have no room after all.
             arena.reopen();
         }
@@ -4676,11 +4641,6 @@ 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?;
@@ -6039,40 +5999,10 @@ 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 {
@@ -6682,8 +6612,7 @@ impl Renderer {
         let flat_frame = match &mut self.flat {
             Some(flat)
                 if self.flat_on
-                    && (!self.draws.is_empty()
-                        || self.gpu_lists.as_ref().is_some_and(|l| l.listed() > 0))
+                    && !self.draws.is_empty()
                     && (self.drawn_sections >= flat.min_sections
                         || pack_terrain
                         || shadow_view.is_some()) =>
@@ -6722,8 +6651,6 @@ 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(),
                     )
                 })
@@ -6731,8 +6658,7 @@ impl Renderer {
             _ => false,
         };
         if flat_frame {
-            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());
+            self.section_draw_calls = self.flat.as_ref().map_or(0, |f| f.shadow_calls[0]);
         } 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();
@@ -7333,26 +7259,19 @@ 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;
-                    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])],
+                    let passes: &[(&wgpu::RenderPipeline, std::ops::Range<usize>)] = match pack_flat
+                    {
+                        Some(p) => &[(&p.0, 0..flat.phase1_calls)],
                         None if split_solid => &[
-                            (&flat.solid_opaque, 0..opaque, &[0]),
-                            (&flat.solid_cutout, cutout..flat.phase1_calls, &[1]),
+                            (&flat.solid_opaque, 0..opaque),
+                            (&flat.solid_cutout, cutout..flat.phase1_calls),
                         ],
-                        None => &[(&flat.solid, 0..flat.phase1_calls, &[0, 1])],
+                        None => &[(&flat.solid, 0..flat.phase1_calls)],
                     };
                     let draw = |pass: &mut wgpu::RenderPass<'_>| {
-                        for (pipeline, calls, groups) in passes {
+                        for (pipeline, calls) 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);
                         }
                     };
@@ -7382,7 +7301,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);
                             }
@@ -9034,12 +8953,6 @@ 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,7 +836,6 @@ 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([]),
@@ -1120,7 +1119,6 @@ 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,7 +223,6 @@ 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]],
@@ -236,7 +235,6 @@ fn plan(
     calls: &mut Vec<Call>,
     ranges: &mut Vec<Range>,
     segments: &mut Vec<(Source, u64, u64)>,
-    base: u64,
 ) -> Plan {
     calls.clear();
     ranges.clear();
@@ -244,7 +242,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 = base;
+    let mut total = 0u64;
     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 {
@@ -542,8 +540,6 @@ 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;
@@ -581,8 +577,6 @@ 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,
@@ -600,7 +594,6 @@ impl Flat {
             &mut self.calls,
             &mut self.range_data,
             &mut self.segments,
-            room,
         );
         self.phase1_calls = phase1_calls;
         self.solid_calls = solid_calls;
@@ -663,9 +656,6 @@ 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,
@@ -735,17 +725,13 @@ 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);
-        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) {
+        for page in self.calls.iter().map(|c| c.page) {
             let table = self.pages[page as usize]
                 .as_mut()
                 .expect("a drawn page has its block table");
@@ -786,27 +772,6 @@ 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]]
@@ -932,7 +897,6 @@ mod tests {
             &mut calls,
             &mut ranges,
             &mut segments,
-            0,
         );
         assert_eq!(total, 6 * 13);
         assert_eq!(phase1_calls, solid_calls);
@@ -981,7 +945,6 @@ 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 };
@@ -1017,7 +980,6 @@ 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 };
@@ -1046,7 +1008,6 @@ mod tests {
             &mut calls,
             &mut ranges,
             &mut segments,
-            0,
         );
         assert_eq!(
             plan,
@@ -1164,7 +1125,6 @@ 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

@@ -1,639 +0,0 @@
-//! 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

@@ -1,178 +0,0 @@
-// 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,11 +29,6 @@ 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>,
@@ -63,7 +58,7 @@ impl PageRuns {
         tag: [u32; 3],
     ) {
         if len > 0 {
-            let band = band(distance);
+            let band = (f64::from_bits(distance).sqrt() / BAND) as u64;
             self.runs.push(Run {
                 key: u64::from(group) << GROUP_SHIFT | band << 32 | u64::from(page),
                 start: start as u32,

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

@@ -49,7 +49,7 @@ pub mod hud_overlays;
 pub mod inventory_effects;
 pub mod java_text;
 pub mod jigsaw_block;
-pub mod json_folder;
+pub use horizon_world::json_folder;
 pub mod level_loading;
 pub mod light;
 pub mod lod;

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

@@ -12,7 +12,7 @@ use crate::{
     assets::{Assets, RenderAssets},
     biome::Biomes,
     blocks::{BlockFacts, RenderShape},
-    chunk::{Chunk, Light, MIN_Y},
+    chunk::{Chunk, MIN_Y},
     graphics::LeafCullingMode,
     models::catalog::BakedFace,
 };
@@ -505,10 +505,6 @@ 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,
@@ -677,7 +673,6 @@ impl MeshTables {
             });
         }
         Ok(Self {
-            solid: info.iter().map(|i| i.solid).collect(),
             info,
             visibility,
             water_still: layer("minecraft:block/water_still")?,
@@ -847,66 +842,43 @@ 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 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.
+        // The section itself is unpacked in one pass; the border cells of the
+        // neighbouring sections are read one by one.
         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 {
-                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]);
+                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);
                     } else {
-                        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()]),
+                        light[i] = 0xF0;
                     }
                 }
             }
@@ -922,28 +894,6 @@ 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;
@@ -1040,12 +990,6 @@ 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
@@ -1097,16 +1041,6 @@ 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| {
@@ -1123,8 +1057,6 @@ 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],
             })
         });
@@ -1133,45 +1065,21 @@ impl<'a> Tints<'a> {
         if cached != UNSET {
             return cached;
         }
-        let [cx, cz, base_y] = self.origin;
-        let (bx, by, bz) = (cx * 16, base_y + y, cz * 16);
+        let [cx, cz, _] = self.origin;
+        let (bx, bz) = (cx * 16, 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 {
-                    // 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);
+                    biome = Self::biome_id(
+                        &self.zoom,
+                        self.region,
+                        self.plains,
+                        self.origin,
+                        [x + dx, y, z + dz],
+                    );
                     layer.zoomed[cell] = biome;
                 }
                 let c = match resolver {
@@ -1183,7 +1091,6 @@ 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;
@@ -1256,8 +1163,6 @@ 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
@@ -1290,7 +1195,12 @@ impl Mesher<'_> {
                         out.lights
                             .push(([center.x * 16 + x, base_y + y, center.z * 16 + z], id));
                     }
-                    if hidden(opaque, x, y, z) {
+                    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
+                        })
+                    {
                         continue;
                     }
                     if state.fluid.kind() != FluidKind::Empty {

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

@@ -103,22 +103,6 @@ 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]
@@ -144,26 +128,6 @@ 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);
@@ -246,29 +210,6 @@ 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;
@@ -277,14 +218,15 @@ fn sort_with(solid: &mut Vec<TerrainVertex>, grouped: bool, scratch: &mut SortSc
         true => b + 7 * usize::from(cutout),
         false => b * 2 + usize::from(cutout),
     };
-    let slots = &mut scratch.slots;
-    slots.clear();
-    slots.extend(quads.iter().map(|q| position(bucket(q), cutout(q)) as u8));
+    let slots: Vec<u8> = quads
+        .iter()
+        .map(|q| position(bucket(q), cutout(q)) as u8)
+        .collect();
     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.iter()) {
+    for (quad, &part) in quads.iter().zip(&slots) {
         slot_counts[usize::from(part)] += 1;
         let b = slot(grouped, usize::from(part)).0;
         if let Bucket::Side(axis, positive) = BUCKETS[b] {
@@ -304,34 +246,33 @@ fn sort_with(solid: &mut Vec<TerrainVertex>, grouped: bool, scratch: &mut SortSc
         *end = total;
     }
     if has_cutout || counts.iter().filter(|&&c| c > 0).count() > 1 {
-        // 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 }
+        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)
                         }
-                        _ => 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());
+                    });
+                }
+                start = end;
+            }
+        }
+        *solid = sorted.into_flattened();
     }
     Sides {
         ends: slot_ends,
@@ -345,74 +286,6 @@ fn sort_with(solid: &mut Vec<TerrainVertex>, grouped: bool, scratch: &mut SortSc
 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-client/src/game/lan.rs

@@ -111,9 +111,14 @@ impl Game<'_> {
 pub(super) struct SyncPreload(std::thread::JoinHandle<std::result::Result<SyncData, String>>);
 
 /// Start reading the server's data beside the world's preparation; `None`
-/// when it is read already, or being read (`preload_world_data`).
+/// when it is read already, or being read: by `sync_data` (the lock is
+/// held), or soon by a world data preload still running, which ends with
+/// it. A second read would be thrown away by `adopt`.
 pub(super) fn preload_sync_data(root: &std::path::Path) -> Option<SyncPreload> {
-    if !SYNC.try_lock().is_ok_and(|sync| sync.is_none()) {
+    // PRELOAD before SYNC, never both: `preload_world_data` takes SYNC
+    // under PRELOAD.
+    if super::world_loading::preload_running() || !SYNC.try_lock().is_ok_and(|sync| sync.is_none())
+    {
         return None;
     }
     let root = root.to_owned();

crates/horizon-client/src/game/world_loading.rs

@@ -388,6 +388,10 @@ pub(super) fn preload_world_data(
                         // seed shares: about 140 ms off opening the world.
                         let _ = data.prepare_decoration("minecraft:normal", "minecraft:overworld");
                     }
+                    // The loot data's item tables and the server's data both
+                    // read the recipe folder: held here, it is read once.
+                    let _recipe_files =
+                        horizon_client::json_folder::read(&resources.join("data/minecraft/recipe"));
                     // The preload is the game's own data; the world's packs'
                     // tables are built when it opens (finish_open_world).
                     let _ = horizon_client::structure_blocks::loot_data(
@@ -459,6 +463,10 @@ pub(super) fn preload_launch_world(
                 let preset = horizon_client::save::WorldGen::preset(&world_root);
                 let _ = data.prepare_decoration(preset, "minecraft:overworld");
             }
+            // As in `preload_world_data`: the loot data and the server's
+            // data share one read of the recipe folder.
+            let _recipe_files =
+                horizon_client::json_folder::read(&resources.join("data/minecraft/recipe"));
             let _ = horizon_client::structure_blocks::loot_data(
                 &resources,
                 &horizon_world::loot::PackOverlay::NONE,
@@ -475,6 +483,15 @@ pub(super) fn preload_launch_world(
     }
 }
 
+/// Whether the "World data preload" thread is still running: it reads the
+/// server's data (`sync_data`) last, so nothing else needs to.
+pub(super) fn preload_running() -> bool {
+    matches!(
+        &*PRELOAD.lock().unwrap_or_else(|e| e.into_inner()),
+        Preload::Started { thread, .. } if !thread.is_finished()
+    )
+}
+
 /// A world opens: what `preload_world_data` read is the world's from now on.
 pub(super) fn world_data_used() {
     *PRELOAD.lock().unwrap_or_else(|e| e.into_inner()) = Preload::Used;

crates/horizon-client/src/lan/data.rs

@@ -208,6 +208,10 @@ impl SyncData {
         resources: &Path,
         overlay: &horizon_world::loot::PackOverlay,
     ) -> Result<Self, String> {
+        // The recipe book, the recipe sync and the advancements read the
+        // recipe folder: held here, they share one read of it.
+        let _recipe_files =
+            horizon_world::json_folder::read(&resources.join("data/minecraft/recipe"));
         let builtin = Builtin::find(resources)
             .and_then(|p| Builtin::load(&p).ok())
             .ok_or("registries.json not found next to the resources")?;

crates/horizon-client/src/lan/recipes.rs

@@ -52,17 +52,16 @@ fn ingredient(v: &Value) -> Option<Ingredient> {
     }
 }
 
-/// Every recipe file, by recipe id (the registry's order).
-pub(super) fn recipes(dir: &Path, prefix: &str, out: &mut BTreeMap<String, Value>) {
-    for entry in std::fs::read_dir(dir).into_iter().flatten().flatten() {
-        let path = entry.path();
-        let name = entry.file_name().to_string_lossy().into_owned();
-        if path.is_dir() {
-            recipes(&path, &format!("{prefix}{name}/"), out);
-        } else if let Some(stem) = name.strip_suffix(".json")
-            && let Ok(v) = serde_json::from_slice(&std::fs::read(&path).unwrap_or_default())
-        {
-            out.insert(format!("minecraft:{prefix}{stem}"), v);
+/// Every recipe file, by recipe id (the registry's order): read through
+/// `json_folder`, so a load holding the folder shares one read of it. A file
+/// that can't be read or doesn't parse is left out.
+pub(super) fn recipes(dir: &Path, out: &mut BTreeMap<String, Value>) {
+    let Ok(folder) = horizon_world::json_folder::read(dir) else {
+        return;
+    };
+    for file in &folder.files {
+        if let Ok(v) = serde_json::from_slice(file.bytes().unwrap_or_default()) {
+            out.insert(format!("minecraft:{}", file.name), v);
         }
     }
 }
@@ -94,7 +93,7 @@ pub(super) fn network_recipes(
     overlay: &horizon_world::loot::PackOverlay,
 ) -> (BTreeMap<String, Box<str>>, HashMap<String, i32>) {
     let mut all = BTreeMap::new();
-    recipes(&resources.join("data/minecraft/recipe"), "", &mut all);
+    recipes(&resources.join("data/minecraft/recipe"), &mut all);
     overlay_recipes(overlay, &mut all);
     // The data's recipe filters, as the server's own tables apply them
     // (#5910), once, before the ingredient normalization, so a swapped

crates/horizon-core/src/paletted.rs

@@ -272,23 +272,6 @@ 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> {
@@ -753,26 +736,6 @@ 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/Cargo.toml

@@ -8,6 +8,7 @@ description = "Horizon vanilla world data, NBT, regions and chunk storage"
 
 [dependencies]
 libm = "=0.2.16"
+log = "0.4"
 serde_json = { version = "1.0.151", features = ["raw_value"] }
 serde = { version = "1", features = ["derive"] }
 sha2 = "0.10"

crates/horizon-world/src/block_collision_scan.rs

@@ -79,48 +79,6 @@ 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>,
@@ -152,10 +110,29 @@ impl<'w, 'c, W: BlockCollisionWorld> BlockCollisionScan<'w, 'c, W> {
         only_suffocating: bool,
         candidate_budget: usize,
     ) -> Result<Self, CollisionScanError<W::Error>> {
-        let (start, sizes, count) = cells(bounds).map_err(|error| match error {
-            CellError::InvalidBounds => CollisionScanError::InvalidBounds,
-            CellError::TooLarge => CollisionScanError::TooLarge,
-        })?;
+        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;
         Ok(Self {
             world,
             context,

crates/horizon-world/src/lib.rs

@@ -77,6 +77,7 @@ pub mod item_use_catalog;
 pub mod item_use_lifecycle;
 pub mod item_use_start;
 pub mod java_hash_map;
+pub mod json_folder;
 pub mod lazy_json;
 pub mod level;
 pub mod liquid_collision;

crates/horizon-world/src/loot.rs

@@ -635,20 +635,15 @@ fn load_dir(
     prefix: &str,
     out: &mut HashMap<String, Value>,
 ) -> io::Result<()> {
-    let Ok(entries) = fs::read_dir(dir) else {
+    // Through `json_folder`: a load holding the folder (the recipes, read by
+    // the item tables, the recipe book and the server's sync too) shares one
+    // read of it.
+    let Ok(folder) = crate::json_folder::read(dir) else {
         return Ok(());
     };
-    for entry in entries.flatten() {
-        let path = entry.path();
-        let Some(stem) = path.file_stem().and_then(|s| s.to_str()) else {
-            continue;
-        };
-        if path.is_dir() {
-            load_dir(&path, namespace, &format!("{prefix}{stem}/"), out)?;
-        } else if path.extension().is_some_and(|e| e == "json") {
-            let value = serde_json::from_slice(&fs::read(&path)?)?;
-            out.insert(format!("{namespace}:{prefix}{stem}"), value);
-        }
+    for file in &folder.files {
+        let value = serde_json::from_slice(file.bytes()?)?;
+        out.insert(format!("{namespace}:{prefix}{}", file.name), value);
     }
     Ok(())
 }

crates/horizon-world/src/world_movement.rs

@@ -1,9 +1,7 @@
 //! 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, scan_cells,
-    },
+    block_collision_scan::{BlockCollisionScan, BlockCollisionWorld, CollisionScanError},
     collision_context::CollisionContext,
     entity_obstruction::EntityIdentity,
 };
@@ -58,19 +56,6 @@ 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,
@@ -84,9 +69,13 @@ where
             {
                 shapes.push(border);
             }
-            for hit in
-                BlockCollisionScan::new(world, request.context, bounds, false, candidate_budget)?
-            {
+            for hit in BlockCollisionScan::new(
+                world,
+                request.context,
+                bounds,
+                false,
+                request.candidate_budget,
+            )? {
                 shapes.push(hit?.shape);
             }
             Ok(shapes)

crates/horizon-world/tests/world_movement.rs

@@ -1,131 +0,0 @@
-//! 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:?}"
-        );
-    }
-}

launcher/Cargo.lock

@@ -2817,6 +2817,7 @@ dependencies = [
  "flate2",
  "horizon-core",
  "libm",
+ "log",
  "lz4_flex",
  "postcard",
  "serde",