Terrain: occlusion culling of solid sections on the GPU (Vulkan)

Two phases, as in GPU-driven renderers: the world pass draws the solid
sections in sight last frame, a Hi-Z pyramid of that depth tests every
solid draw's section box (grown by a block) in a compute pass, and a
second pass draws the ones in sight that phase 1 did not. The test's
result is the next frame's phase 1. Culled draws keep their indirect entry
with an instance count of 0, so this is on where multi-draw is native
(Vulkan); HORIZON_GPU_CULL=0/1 overrides, `section-cull on|off` switches
it in a running game for A/B checks.

A section is left out only when its box is behind phase 1's depth
everywhere it covers, so every fragment it has would fail the depth test.
Translucent draws are not culled.

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

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

@@ -183,6 +183,16 @@ impl Game<'_> {
                 }
                 Ok(Some(json!({ "ok": true, "flat": r.flat_draws() })))
             }
+            "section-cull" => {
+                let r = self.renderer.as_mut().ok_or("no renderer")?;
+                match words.get(1).copied() {
+                    None => {}
+                    Some("on") => r.set_section_cull(true)?,
+                    Some("off") => r.set_section_cull(false)?,
+                    Some(other) => return Err(format!("unknown section cull mode {other}").into()),
+                }
+                Ok(Some(json!({ "ok": true, "cull": r.section_cull() })))
+            }
             "sculk" => Ok(Some(self.sculk_command(&words)?)),
             "status" => Ok(Some(self.control_status())),
             "pad" => Ok(Some(self.pad_command(&words)?)),

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

@@ -12,6 +12,7 @@ pub const HELP: &[&str] = &[
     "particle NAME X Y Z [DX DY DZ] [COUNT]  spawn particles (vanilla /particle spread)",
     "signedit X Y Z [front|back]    open the sign editor for the sign at X Y Z",
     "section-draws [flat|indirect]  how terrain sections are drawn (A/B checks); replies the mode",
+    "section-cull [on|off]          occlusion culling of solid sections (A/B checks); replies the mode",
 ];
 
 impl Game<'_> {

crates/horizon-client/src/gfx.rs

@@ -58,6 +58,7 @@ pub fn outline_vertices(edges: &[[f64; 6]], camera: glam::DVec3) -> Vec<[f32; 6]
 }
 
 mod arena;
+mod cull;
 mod flat;
 mod java_stages;
 mod occlusion;
@@ -270,6 +271,9 @@ pub struct Renderer {
     flat: Option<flat::Flat>,
     /// `flat` draws this frame's sections (off: the per-section draws).
     flat_on: bool,
+    /// Occlusion culling of solid section draws, where multi-draw is native.
+    cull: Option<cull::Cull>,
+    cull_on: bool,
     sky: wgpu::RenderPipeline,
     dark_disc: wgpu::RenderPipeline,
     lines: wgpu::RenderPipeline,
@@ -1100,7 +1104,9 @@ impl Renderer {
         let instances = device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("Section origins"),
             size: instance_capacity * 16,
-            usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST,
+            usage: wgpu::BufferUsages::VERTEX
+                | wgpu::BufferUsages::COPY_DST
+                | wgpu::BufferUsages::STORAGE,
             mapped_at_creation: false,
         });
         let outline = device.create_buffer(&wgpu::BufferDescriptor {
@@ -1739,12 +1745,23 @@ impl Renderer {
         let indirect = device.create_buffer(&wgpu::BufferDescriptor {
             label: Some("Section draws"),
             size: 1024 * 20,
-            usage: wgpu::BufferUsages::INDIRECT | wgpu::BufferUsages::COPY_DST,
+            usage: wgpu::BufferUsages::INDIRECT
+                | wgpu::BufferUsages::COPY_DST
+                | wgpu::BufferUsages::STORAGE,
             mapped_at_creation: false,
         });
         let multi_draw = device
             .features()
             .contains(wgpu::Features::INDIRECT_FIRST_INSTANCE);
+        // HORIZON_GPU_CULL=0 turns occlusion culling off, 1 on where it is
+        // not by default (Metal replays multi-draws on the CPU).
+        let cull = match std::env::var("HORIZON_GPU_CULL").as_deref() {
+            Ok("0") => false,
+            Ok("1") => true,
+            _ => cull::Cull::supported(info.backend),
+        }
+        .then(|| cull::Cull::new(&device))
+        .filter(|_| multi_draw);
         let world_post = world_post::Post::new(&device, config.format.remove_srgb_suffix());
         let uploads = upload::Uploads::new(&device);
         let section_uploads = upload::Uploads::with_chunk(&device, SECTION_UPLOAD_CHUNK);
@@ -1772,6 +1789,8 @@ impl Renderer {
             solid,
             translucent,
             flat_on: flat.is_some(),
+            cull_on: cull.is_some(),
+            cull,
             flat,
             sky,
             dark_disc,
@@ -1933,7 +1952,9 @@ impl Renderer {
                 sample_count: 1,
                 dimension: wgpu::TextureDimension::D2,
                 format: wgpu::TextureFormat::Depth32Float,
-                usage: wgpu::TextureUsages::RENDER_ATTACHMENT,
+                // Occlusion culling reads the world's depth.
+                usage: wgpu::TextureUsages::RENDER_ATTACHMENT
+                    | wgpu::TextureUsages::TEXTURE_BINDING,
                 view_formats: &[],
             })
             .create_view(&Default::default())
@@ -1998,6 +2019,18 @@ impl Renderer {
         self.flat_on = on;
         Ok(())
     }
+    /// Whether solid sections are occlusion culled (`cull`).
+    pub fn section_cull(&self) -> bool {
+        self.cull_on
+    }
+    /// Switches occlusion culling, for A/B checks.
+    pub fn set_section_cull(&mut self, on: bool) -> Result<()> {
+        if on && self.cull.is_none() {
+            return Err("occlusion culling is off on this device".into());
+        }
+        self.cull_on = on;
+        Ok(())
+    }
     pub fn section_count(&self) -> usize {
         self.sections.len()
     }
@@ -2117,7 +2150,13 @@ impl Renderer {
     /// origins bound: one indirect multi-draw, or a loop without it.
     /// Consecutive draws in one arena page share a vertex buffer binding,
     /// keeping the list's order (translucent sections go back to front).
-    fn draw_sections(&self, pass: &mut wgpu::RenderPass<'_>, first: u32, count: u32) {
+    fn draw_sections(
+        &self,
+        pass: &mut wgpu::RenderPass<'_>,
+        indirect: &wgpu::Buffer,
+        first: u32,
+        count: u32,
+    ) {
         let (mut i, end) = (first as usize, (first + count) as usize);
         while i < end {
             let (page, index_page) = self.draw_pages[i];
@@ -2129,7 +2168,7 @@ impl Renderer {
             let indices = index_page.map_or(&self.indices, |p| self.sorted_indices.buffer(p));
             pass.set_index_buffer(indices.slice(..), wgpu::IndexFormat::Uint32);
             if self.multi_draw {
-                pass.multi_draw_indexed_indirect(&self.indirect, i as u64 * 20, run as u32);
+                pass.multi_draw_indexed_indirect(indirect, i as u64 * 20, run as u32);
             } else {
                 for d in &self.draws[i..i + run] {
                     pass.draw_indexed(d[2]..d[2] + d[0], d[3] as i32, d[4]..d[4] + 1);
@@ -2574,7 +2613,9 @@ impl Renderer {
             self.instances = self.device.create_buffer(&wgpu::BufferDescriptor {
                 label: Some("Section origins"),
                 size: self.instance_capacity * 16,
-                usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_DST,
+                usage: wgpu::BufferUsages::VERTEX
+                    | wgpu::BufferUsages::COPY_DST
+                    | wgpu::BufferUsages::STORAGE,
                 mapped_at_creation: false,
             });
         }
@@ -2833,6 +2874,8 @@ impl Renderer {
             frame.hud,
             [self.config.width, self.config.height],
         );
+        let cull_wanted = self.cull_on && !frame.title && solid_draws > 0;
+        let mut cull_staged = false;
         let flat_frame = match &mut self.flat {
             Some(flat)
                 if self.flat_on && !self.draws.is_empty() && self.draws.len() >= flat.min_draws =>
@@ -2864,7 +2907,9 @@ impl Renderer {
                 self.indirect = self.device.create_buffer(&wgpu::BufferDescriptor {
                     label: Some("Section draws"),
                     size: self.indirect_capacity * 20,
-                    usage: wgpu::BufferUsages::INDIRECT | wgpu::BufferUsages::COPY_DST,
+                    usage: wgpu::BufferUsages::INDIRECT
+                        | wgpu::BufferUsages::COPY_DST
+                        | wgpu::BufferUsages::STORAGE,
                     mapped_at_creation: false,
                 });
             }
@@ -2874,6 +2919,17 @@ impl Renderer {
                 0,
                 bytemuck::cast_slice(&self.draws),
             );
+            if cull_wanted && let Some(cull) = &mut self.cull {
+                cull.upload(
+                    &self.device,
+                    &mut self.uploads,
+                    &mut encoder,
+                    visible.iter().map(|v| v.slot),
+                    u64::from(solid_draws),
+                    [self.config.width, self.config.height],
+                );
+                cull_staged = true;
+            }
         }
         let view_target = target_texture.create_view(&wgpu::TextureViewDescriptor {
             format: Some(self.config.format.remove_srgb_suffix()),
@@ -2891,6 +2947,17 @@ impl Renderer {
             Some((color, depth)) => (color, depth),
             None => (&view_target, &self.depth),
         };
+        let cull_group = match &self.cull {
+            Some(cull) if cull_staged && scaled.is_none() => Some(cull.prepare(
+                &self.device,
+                &mut encoder,
+                &self.indirect,
+                &self.instances,
+                &self.globals,
+                solid_draws,
+            )),
+            _ => None,
+        };
         {
             let attachments = [Some(wgpu::RenderPassColorAttachment {
                 view: world_view,
@@ -2919,8 +2986,9 @@ impl Renderer {
                     view: world_depth,
                     depth_ops: Some(wgpu::Operations {
                         load: wgpu::LoadOp::Clear(0.0),
-                        // A shader pack reads the world's depth.
-                        store: if pack {
+                        // A shader pack reads the world's depth, and
+                        // so does occlusion culling after phase 1.
+                        store: if pack || cull_group.is_some() {
                             wgpu::StoreOp::Store
                         } else {
                             wgpu::StoreOp::Discard
@@ -2970,7 +3038,52 @@ impl Renderer {
                     pass.set_pipeline(&self.solid);
                     pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint32);
                     pass.set_vertex_buffer(1, self.instances.slice(..));
-                    self.draw_sections(&mut pass, 0, solid_draws);
+                    match (&self.cull, &cull_group) {
+                        (Some(cull), Some(group)) => {
+                            // Phase 1, then what the Hi-Z test of its depth
+                            // finds in sight, in a pass that loads both.
+                            self.draw_sections(&mut pass, &cull.phase1, 0, solid_draws);
+                            drop(pass);
+                            cull.test(&self.device, &mut encoder, group, world_depth, solid_draws);
+                            let attachments = [Some(wgpu::RenderPassColorAttachment {
+                                view: world_view,
+                                resolve_target: None,
+                                depth_slice: None,
+                                ops: wgpu::Operations {
+                                    load: wgpu::LoadOp::Load,
+                                    store: wgpu::StoreOp::Store,
+                                },
+                            })];
+                            pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor {
+                                label: Some("Horizon world, after culling"),
+                                color_attachments: &attachments,
+                                depth_stencil_attachment: Some(
+                                    wgpu::RenderPassDepthStencilAttachment {
+                                        view: world_depth,
+                                        depth_ops: Some(wgpu::Operations {
+                                            load: wgpu::LoadOp::Load,
+                                            store: if pack {
+                                                wgpu::StoreOp::Store
+                                            } else {
+                                                wgpu::StoreOp::Discard
+                                            },
+                                        }),
+                                        stencil_ops: None,
+                                    },
+                                ),
+                                ..Default::default()
+                            });
+                            pass.set_bind_group(0, &self.group, &[]);
+                            pass.set_pipeline(&self.solid);
+                            pass.set_index_buffer(
+                                self.indices.slice(..),
+                                wgpu::IndexFormat::Uint32,
+                            );
+                            pass.set_vertex_buffer(1, self.instances.slice(..));
+                            self.draw_sections(&mut pass, &cull.phase2, 0, solid_draws);
+                        }
+                        _ => self.draw_sections(&mut pass, &self.indirect, 0, solid_draws),
+                    }
                 }
             }
             if entity_quads > 0 {
@@ -3022,7 +3135,7 @@ impl Renderer {
                     pass.set_pipeline(&self.translucent);
                     pass.set_index_buffer(self.indices.slice(..), wgpu::IndexFormat::Uint32);
                     pass.set_vertex_buffer(1, self.instances.slice(..));
-                    self.draw_sections(&mut pass, solid_draws, translucent_draws);
+                    self.draw_sections(&mut pass, &self.indirect, solid_draws, translucent_draws);
                 }
             }
             // LevelRenderer.executeClassicTransparency: translucent targets'

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

@@ -0,0 +1,418 @@
+//! Occlusion culling of solid section draws, in two phases on the GPU.
+//!
+//! Frustum and cave culling leave many sections that terrain in front hides
+//! entirely (about half of the solid vertices drawn from the ground at render
+//! distance 16), and drawing fewer vertices is what helps an integrated GPU.
+//! Phase 1 draws the solid sections that were in sight last frame. A Hi-Z
+//! pyramid of its depth then tests every solid draw's section box: phase 2
+//! draws the ones in sight that phase 1 did not, and the result is the next
+//! frame's phase 1.
+//!
+//! The picture stays the one without culling: a section is left out only
+//! when its box (grown by a block, the farthest a block model reaches out of
+//! its cell) is behind phase 1's depth everywhere it covers, so all its
+//! fragments would fail the depth test against the final depth, which is
+//! only nearer. Translucent draws are not culled.
+//!
+//! Draws stay `DrawIndexedIndirectArgs` entries: the culled ones get an
+//! instance count of 0, so this is for backends with native multi-draw
+//! (Vulkan). Metal replays a multi-draw on the CPU, entry by entry.
+use std::num::NonZeroU64;
+
+const WORDS: u64 = 5 * 4;
+
+#[repr(C)]
+#[derive(Clone, Copy, bytemuck::Pod, bytemuck::Zeroable)]
+struct Params {
+    count: u32,
+    levels: u32,
+    screen: [u32; 2],
+}
+
+/// The Hi-Z pyramid for one depth buffer size.
+struct Pyramid {
+    size: [u32; 2],
+    view: wgpu::TextureView,
+    levels: u32,
+    /// Level l + 1 from level l, with its workgroup counts.
+    reduce_groups: Vec<(wgpu::BindGroup, [u32; 2])>,
+    level0: wgpu::TextureView,
+    level0_groups: [u32; 2],
+}
+
+pub struct Cull {
+    prepare: wgpu::ComputePipeline,
+    test: wgpu::ComputePipeline,
+    reduce_depth: wgpu::ComputePipeline,
+    reduce: wgpu::ComputePipeline,
+    layout: wgpu::BindGroupLayout,
+    pyramid: Option<Pyramid>,
+    params: wgpu::Buffer,
+    pub phase1: wgpu::Buffer,
+    pub phase2: wgpu::Buffer,
+    draw_capacity: u64,
+    slots: wgpu::Buffer,
+    slot_capacity: u64,
+    slot_data: Vec<u32>,
+    visibility: wgpu::Buffer,
+    visibility_capacity: u64,
+}
+
+fn buffer(
+    device: &wgpu::Device,
+    label: &str,
+    size: u64,
+    usage: wgpu::BufferUsages,
+) -> wgpu::Buffer {
+    device.create_buffer(&wgpu::BufferDescriptor {
+        label: Some(label),
+        size,
+        usage,
+        mapped_at_creation: false,
+    })
+}
+
+const PHASE_USAGE: wgpu::BufferUsages =
+    wgpu::BufferUsages::INDIRECT.union(wgpu::BufferUsages::STORAGE);
+
+impl Cull {
+    /// Whether the device draws an indirect multi-draw natively.
+    pub fn supported(backend: wgpu::Backend) -> bool {
+        backend == wgpu::Backend::Vulkan
+    }
+    pub fn new(device: &wgpu::Device) -> Self {
+        let storage = |binding, read_only| wgpu::BindGroupLayoutEntry {
+            binding,
+            visibility: wgpu::ShaderStages::COMPUTE,
+            ty: wgpu::BindingType::Buffer {
+                ty: wgpu::BufferBindingType::Storage { read_only },
+                has_dynamic_offset: false,
+                min_binding_size: None,
+            },
+            count: None,
+        };
+        let uniform = |binding| wgpu::BindGroupLayoutEntry {
+            binding,
+            visibility: wgpu::ShaderStages::COMPUTE,
+            ty: wgpu::BindingType::Buffer {
+                ty: wgpu::BufferBindingType::Uniform,
+                has_dynamic_offset: false,
+                min_binding_size: None,
+            },
+            count: None,
+        };
+        let layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+            label: Some("Section culling"),
+            entries: &[
+                storage(0, true),
+                storage(1, true),
+                storage(2, false),
+                storage(3, false),
+                storage(4, false),
+                storage(5, true),
+                uniform(6),
+                uniform(7),
+                wgpu::BindGroupLayoutEntry {
+                    binding: 8,
+                    visibility: wgpu::ShaderStages::COMPUTE,
+                    ty: wgpu::BindingType::Texture {
+                        sample_type: wgpu::TextureSampleType::Float { filterable: false },
+                        view_dimension: wgpu::TextureViewDimension::D2,
+                        multisampled: false,
+                    },
+                    count: None,
+                },
+            ],
+        });
+        let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+            label: Some("Section culling"),
+            bind_group_layouts: &[Some(&layout)],
+            ..Default::default()
+        });
+        let cull = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("Section culling"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("cull.wgsl").into()),
+        });
+        let hiz = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+            label: Some("Hi-Z"),
+            source: wgpu::ShaderSource::Wgsl(include_str!("hiz.wgsl").into()),
+        });
+        let pipeline = |module, layout: Option<&wgpu::PipelineLayout>, entry: &str| {
+            device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+                label: Some(entry),
+                layout,
+                module,
+                entry_point: Some(entry),
+                compilation_options: Default::default(),
+                cache: None,
+            })
+        };
+        let storage_usage = wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST;
+        Self {
+            prepare: pipeline(&cull, Some(&pipeline_layout), "prepare"),
+            test: pipeline(&cull, Some(&pipeline_layout), "test"),
+            reduce_depth: pipeline(&hiz, None, "reduce_depth"),
+            reduce: pipeline(&hiz, None, "reduce"),
+            layout,
+            pyramid: None,
+            params: buffer(
+                device,
+                "Section culling parameters",
+                std::mem::size_of::<Params>() as u64,
+                wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST,
+            ),
+            phase1: buffer(device, "Section draws, phase 1", WORDS, PHASE_USAGE),
+            phase2: buffer(device, "Section draws, phase 2", WORDS, PHASE_USAGE),
+            draw_capacity: 1,
+            slots: buffer(device, "Visible section slots", 4, storage_usage),
+            slot_capacity: 1,
+            slot_data: Vec::new(),
+            visibility: buffer(device, "Section visibility", 4, wgpu::BufferUsages::STORAGE),
+            visibility_capacity: 1,
+        }
+    }
+
+    /// The Hi-Z pyramid for a `size` depth buffer: level 0 is half its size.
+    fn pyramid(&mut self, device: &wgpu::Device, size: [u32; 2]) -> &Pyramid {
+        if self.pyramid.as_ref().is_some_and(|p| p.size == size) {
+            return self.pyramid.as_ref().expect("checked");
+        }
+        let base = size.map(|v| (v / 2).max(1));
+        let levels = 32 - base[0].max(base[1]).leading_zeros();
+        let texture = device.create_texture(&wgpu::TextureDescriptor {
+            label: Some("Hi-Z"),
+            size: wgpu::Extent3d {
+                width: base[0],
+                height: base[1],
+                depth_or_array_layers: 1,
+            },
+            mip_level_count: levels,
+            sample_count: 1,
+            dimension: wgpu::TextureDimension::D2,
+            format: wgpu::TextureFormat::R32Float,
+            usage: wgpu::TextureUsages::STORAGE_BINDING | wgpu::TextureUsages::TEXTURE_BINDING,
+            view_formats: &[],
+        });
+        let level = |l| {
+            texture.create_view(&wgpu::TextureViewDescriptor {
+                base_mip_level: l,
+                mip_level_count: Some(1),
+                ..Default::default()
+            })
+        };
+        let groups = |edge: u32| edge.div_ceil(8);
+        let layout = self.reduce.get_bind_group_layout(0);
+        let reduce_groups = (1..levels)
+            .map(|l| {
+                let dims = base.map(|v| (v >> l).max(1));
+                let group = device.create_bind_group(&wgpu::BindGroupDescriptor {
+                    label: Some("Hi-Z level"),
+                    layout: &layout,
+                    entries: &[
+                        wgpu::BindGroupEntry {
+                            binding: 2,
+                            resource: wgpu::BindingResource::TextureView(&level(l - 1)),
+                        },
+                        wgpu::BindGroupEntry {
+                            binding: 3,
+                            resource: wgpu::BindingResource::TextureView(&level(l)),
+                        },
+                    ],
+                });
+                (group, dims.map(groups))
+            })
+            .collect();
+        self.pyramid.insert(Pyramid {
+            size,
+            view: texture.create_view(&Default::default()),
+            levels,
+            reduce_groups,
+            level0: level(0),
+            level0_groups: base.map(groups),
+        })
+    }
+
+    /// Stage this frame's inputs: the resident slot of each visible section.
+    /// `draws` is the number of solid draws.
+    pub fn upload(
+        &mut self,
+        device: &wgpu::Device,
+        uploads: &mut super::upload::Uploads,
+        encoder: &mut wgpu::CommandEncoder,
+        slots: impl Iterator<Item = u32>,
+        draws: u64,
+        size: [u32; 2],
+    ) {
+        self.slot_data.clear();
+        self.slot_data.extend(slots);
+        let len = self.slot_data.len() as u64;
+        if len > self.slot_capacity {
+            self.slot_capacity = len.next_power_of_two();
+            self.slots = buffer(
+                device,
+                "Visible section slots",
+                self.slot_capacity * 4,
+                wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_DST,
+            );
+        }
+        let slots_needed = self.slot_data.iter().max().map_or(0, |&m| u64::from(m) + 1);
+        if slots_needed > self.visibility_capacity {
+            // A new buffer starts with nothing in sight: phase 2 tests all.
+            self.visibility_capacity = slots_needed.next_power_of_two();
+            self.visibility = buffer(
+                device,
+                "Section visibility",
+                self.visibility_capacity * 4,
+                wgpu::BufferUsages::STORAGE,
+            );
+        }
+        if draws > self.draw_capacity {
+            self.draw_capacity = draws.next_power_of_two();
+            self.phase1 = buffer(
+                device,
+                "Section draws, phase 1",
+                self.draw_capacity * WORDS,
+                PHASE_USAGE,
+            );
+            self.phase2 = buffer(
+                device,
+                "Section draws, phase 2",
+                self.draw_capacity * WORDS,
+                PHASE_USAGE,
+            );
+        }
+        uploads.buffer(
+            encoder,
+            &self.slots,
+            0,
+            bytemuck::cast_slice(&self.slot_data),
+        );
+        let levels = self.pyramid(device, size).levels;
+        let params = Params {
+            count: draws as u32,
+            levels,
+            screen: size,
+        };
+        uploads.buffer(encoder, &self.params, 0, bytemuck::bytes_of(&params));
+    }
+
+    fn group(
+        &self,
+        device: &wgpu::Device,
+        draws: &wgpu::Buffer,
+        origins: &wgpu::Buffer,
+        globals: &wgpu::Buffer,
+    ) -> wgpu::BindGroup {
+        fn whole(binding: u32, buffer: &wgpu::Buffer) -> wgpu::BindGroupEntry<'_> {
+            wgpu::BindGroupEntry {
+                binding,
+                resource: buffer.as_entire_binding(),
+            }
+        }
+        let pyramid = self.pyramid.as_ref().expect("uploaded first");
+        device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("Section culling"),
+            layout: &self.layout,
+            entries: &[
+                whole(0, draws),
+                whole(1, &self.slots),
+                whole(2, &self.visibility),
+                whole(3, &self.phase1),
+                whole(4, &self.phase2),
+                whole(5, origins),
+                wgpu::BindGroupEntry {
+                    binding: 6,
+                    resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding {
+                        buffer: globals,
+                        offset: 0,
+                        size: NonZeroU64::new(64),
+                    }),
+                },
+                whole(7, &self.params),
+                wgpu::BindGroupEntry {
+                    binding: 8,
+                    resource: wgpu::BindingResource::TextureView(&pyramid.view),
+                },
+            ],
+        })
+    }
+
+    /// Before the world pass: fill both phases' draws from `draws`.
+    pub fn prepare(
+        &self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        draws: &wgpu::Buffer,
+        origins: &wgpu::Buffer,
+        globals: &wgpu::Buffer,
+        count: u32,
+    ) -> wgpu::BindGroup {
+        let group = self.group(device, draws, origins, globals);
+        let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("Section culling, phase 1"),
+            timestamp_writes: None,
+        });
+        pass.set_pipeline(&self.prepare);
+        pass.set_bind_group(0, &group, &[]);
+        pass.dispatch_workgroups(count.div_ceil(64), 1, 1);
+        group
+    }
+
+    /// After phase 1: build the Hi-Z pyramid from `depth` and choose phase 2.
+    pub fn test(
+        &self,
+        device: &wgpu::Device,
+        encoder: &mut wgpu::CommandEncoder,
+        group: &wgpu::BindGroup,
+        depth: &wgpu::TextureView,
+        count: u32,
+    ) {
+        let pyramid = self.pyramid.as_ref().expect("uploaded first");
+        let level0 = device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("Hi-Z level 0"),
+            layout: &self.reduce_depth.get_bind_group_layout(0),
+            entries: &[
+                wgpu::BindGroupEntry {
+                    binding: 0,
+                    resource: wgpu::BindingResource::TextureView(depth),
+                },
+                wgpu::BindGroupEntry {
+                    binding: 1,
+                    resource: wgpu::BindingResource::TextureView(&pyramid.level0),
+                },
+            ],
+        });
+        let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+            label: Some("Section culling, phase 2"),
+            timestamp_writes: None,
+        });
+        pass.set_pipeline(&self.reduce_depth);
+        pass.set_bind_group(0, &level0, &[]);
+        let [x, y] = pyramid.level0_groups;
+        pass.dispatch_workgroups(x, y, 1);
+        pass.set_pipeline(&self.reduce);
+        for (level, [x, y]) in &pyramid.reduce_groups {
+            pass.set_bind_group(0, level, &[]);
+            pass.dispatch_workgroups(*x, *y, 1);
+        }
+        pass.set_pipeline(&self.test);
+        pass.set_bind_group(0, group, &[]);
+        pass.dispatch_workgroups(count.div_ceil(64), 1, 1);
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    #[test]
+    fn shaders_validate() {
+        for source in [include_str!("cull.wgsl"), include_str!("hiz.wgsl")] {
+            let module = naga::front::wgsl::parse_str(source).unwrap();
+            naga::valid::Validator::new(
+                naga::valid::ValidationFlags::all(),
+                naga::valid::Capabilities::all(),
+            )
+            .validate(&module)
+            .unwrap();
+        }
+    }
+}

crates/horizon-client/src/gfx/cull.wgsl

@@ -0,0 +1,117 @@
+// Occlusion culling of solid section draws (gfx/cull.rs).
+
+struct Params {
+    // Solid section draws.
+    count: u32,
+    // Hi-Z mip levels.
+    levels: u32,
+    // The depth buffer's size in pixels.
+    screen: vec2<u32>,
+}
+// The leading member of terrain.wgsl's Globals.
+struct Globals {
+    view_proj: mat4x4<f32>,
+}
+
+// The frame's indirect draws, 5 words each (index count, instance count,
+// first index, base vertex, first instance = the visible section).
+@group(0) @binding(0) var<storage, read> draws: array<u32>;
+// Resident slot of each visible section.
+@group(0) @binding(1) var<storage, read> slots: array<u32>;
+// Per resident slot: 1 when its box was in sight at the last test.
+@group(0) @binding(2) var<storage, read_write> visibility: array<u32>;
+@group(0) @binding(3) var<storage, read_write> phase1: array<u32>;
+@group(0) @binding(4) var<storage, read_write> phase2: array<u32>;
+// Camera-relative origin of each visible section.
+@group(0) @binding(5) var<storage, read> origins: array<vec4<f32>>;
+@group(0) @binding(6) var<uniform> globals: Globals;
+@group(0) @binding(7) var<uniform> params: Params;
+// Farthest depth (reverse Z: the smallest value) of each texel's pixels.
+@group(0) @binding(8) var hiz: texture_2d<f32>;
+
+// Phase 1 draws the sections in sight last frame; phase 2 starts empty.
+@compute @workgroup_size(64)
+fn prepare(@builtin(global_invocation_id) id: vec3<u32>) {
+    let j = id.x;
+    if j >= params.count {
+        return;
+    }
+    let b = j * 5u;
+    for (var k = 0u; k < 5u; k++) {
+        phase1[b + k] = draws[b + k];
+        phase2[b + k] = draws[b + k];
+    }
+    phase1[b + 1u] = min(visibility[slots[draws[b + 4u]]], 1u);
+    phase2[b + 1u] = 0u;
+}
+
+// After phase 1's depth: phase 2 draws what is in sight and phase 1 did not
+// draw; the result is the next frame's phase 1.
+@compute @workgroup_size(64)
+fn test(@builtin(global_invocation_id) id: vec3<u32>) {
+    let j = id.x;
+    if j >= params.count {
+        return;
+    }
+    let b = j * 5u;
+    let section = draws[b + 4u];
+    let seen = box_in_sight(origins[section].xyz);
+    visibility[slots[section]] = select(0u, 1u, seen);
+    phase2[b + 1u] = select(0u, 1u, seen && phase1[b + 1u] == 0u);
+}
+
+// Whether the section's box, grown by a block (models reach one block out
+// of their cell), may have a pixel in front of the depth buffer: false only
+// when the box's nearest depth is behind the farthest depth everywhere it
+// covers, so every fragment of the section would fail the depth test.
+fn box_in_sight(origin: vec3<f32>) -> bool {
+    var lo = vec2(1e30);
+    var hi = vec2(-1e30);
+    var front = 0.0;
+    for (var c = 0u; c < 8u; c++) {
+        let corner = origin + vec3(
+            select(-1.0, 17.0, (c & 1u) != 0u),
+            select(-1.0, 17.0, (c & 2u) != 0u),
+            select(-1.0, 17.0, (c & 4u) != 0u),
+        );
+        let clip = globals.view_proj * vec4(corner, 1.0);
+        if clip.w <= 1e-4 {
+            // Crosses the camera plane.
+            return true;
+        }
+        let ndc = clip.xyz / clip.w;
+        lo = min(lo, ndc.xy);
+        hi = max(hi, ndc.xy);
+        front = max(front, ndc.z);
+    }
+    let size = vec2<f32>(params.screen);
+    // Pixels whose centres the box can cover (y down).
+    let p0 = floor(vec2((lo.x + 1.0) * 0.5, (1.0 - hi.y) * 0.5) * size);
+    let p1 = floor(vec2((hi.x + 1.0) * 0.5, (1.0 - lo.y) * 0.5) * size);
+    if p1.x < 0.0 || p1.y < 0.0 || p0.x >= size.x || p0.y >= size.y {
+        return false;
+    }
+    let a = vec2<u32>(clamp(p0, vec2(0.0), size - 1.0));
+    let z = vec2<u32>(clamp(p1, vec2(0.0), size - 1.0));
+    // Level 0 texel t covers pixels 2t and 2t + 1, level l + 1 texel t its
+    // level-l texels 2t and 2t + 1; each level's last texel also covers an
+    // odd remainder, hence the clamps.
+    var level = 0u;
+    var t0 = a >> vec2(1u);
+    var t1 = z >> vec2(1u);
+    while level + 1u < params.levels && (t1.x - t0.x > 3u || t1.y - t0.y > 3u) {
+        level++;
+        t0 >>= vec2(1u);
+        t1 >>= vec2(1u);
+    }
+    let last = textureDimensions(hiz, level) - 1u;
+    t1 = min(t1, last);
+    t0 = min(t0, t1);
+    var far = 1.0;
+    for (var y = t0.y; y <= t1.y; y++) {
+        for (var x = t0.x; x <= t1.x; x++) {
+            far = min(far, textureLoad(hiz, vec2(x, y), i32(level)).r);
+        }
+    }
+    return front >= far;
+}

crates/horizon-client/src/gfx/hiz.wgsl

@@ -0,0 +1,48 @@
+// Hi-Z pyramid for occlusion culling (gfx/cull.rs): each texel holds the
+// farthest depth (reverse Z: the smallest value) of the pixels it covers.
+
+@group(0) @binding(0) var depth: texture_depth_2d;
+@group(0) @binding(1) var level0: texture_storage_2d<r32float, write>;
+
+// Hi-Z level 0 from the depth buffer.
+@compute @workgroup_size(8, 8)
+fn reduce_depth(@builtin(global_invocation_id) id: vec3<u32>) {
+    let size = textureDimensions(level0);
+    if id.x >= size.x || id.y >= size.y {
+        return;
+    }
+    let source = textureDimensions(depth);
+    let first = id.xy * 2u;
+    let odd = (id.xy == size - 1u) & ((source & vec2(1u)) == vec2(1u));
+    let to = min(first + select(vec2(1u), vec2(2u), odd), source - 1u);
+    var m = 1.0;
+    for (var y = first.y; y <= to.y; y++) {
+        for (var x = first.x; x <= to.x; x++) {
+            m = min(m, textureLoad(depth, vec2(x, y), 0));
+        }
+    }
+    textureStore(level0, id.xy, vec4(m));
+}
+
+@group(0) @binding(2) var source_level: texture_2d<f32>;
+@group(0) @binding(3) var next_level: texture_storage_2d<r32float, write>;
+
+// Hi-Z level l + 1 from level l.
+@compute @workgroup_size(8, 8)
+fn reduce(@builtin(global_invocation_id) id: vec3<u32>) {
+    let size = textureDimensions(next_level);
+    if id.x >= size.x || id.y >= size.y {
+        return;
+    }
+    let source = textureDimensions(source_level);
+    let first = id.xy * 2u;
+    let odd = (id.xy == size - 1u) & ((source & vec2(1u)) == vec2(1u));
+    let to = min(first + select(vec2(1u), vec2(2u), odd), source - 1u);
+    var m = 1.0;
+    for (var y = first.y; y <= to.y; y++) {
+        for (var x = first.x; x <= to.x; x++) {
+            m = min(m, textureLoad(source_level, vec2(x, y), 0).r);
+        }
+    }
+    textureStore(next_level, id.xy, vec4(m));
+}