Add GPU-accelerated rendering: cpu/gpu/hybrid RenderBackend (Phase 8)

Splits RenderBackend's per-voxel face-visibility extraction and tile
shading out as GPU-offloadable steps (wgpu compute shaders), while
keeping greedy-mesh merge/compaction CPU-only per the plan's "partial
GPU rendering" design. RENDER_BACKEND=cpu|gpu|hybrid selects the
strategy, falling back to cpu automatically (logged) if no compatible
GPU adapter is found. Verified against a real GPU: all tests pass,
including ones asserting byte-identical output between the cpu and
gpu backends; a new benchmark example honestly shows cpu currently
outperforming gpu/hybrid at realistic batch sizes since each call is
its own dispatch/readback round trip rather than batched across a
whole render batch (documented as a follow-up optimization).

Also fixes two real, pre-existing gaps found while validating the
worker's actual `docker build`: a missing .dockerignore was sending
the local multi-GB target/ dir into the build context, and the
Dockerfile's rust:1.80 pin was already too old for current
transitive dependency MSRVs (bumped to rust:1.97).

Co-Authored-By: Claude Sonnet 5 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_015tKdPZt78zbPUZMXWzKEKt

README.md

@@ -11,12 +11,19 @@ Three independently-deployable services, each its own docker-compose service:
 
 - `api/` — ElysiaJS on Bun. The I/O layer: WS gateway for mod connections, chat relay, chunk
   store, marker/waypoint sharing, admin config, region export, tile/mesh serving.
-- `worker/` — Rust. CPU-bound rendering: tile rasterization and chunk meshing, consumed off a
-  Redis dirty-chunk stream. Stateless — scale it with `docker compose up --scale worker=N`, or
-  run instances on separate hardware pointed at the same Postgres/Redis/MinIO over a private
-  network. Rendering strategy (`cpu`/`gpu`/`hybrid`) is config-selectable behind a
-  `RenderBackend` trait; only `cpu` (rayon) exists so far — `gpu`/`hybrid` (wgpu) land in
-  Phase 8.
+- `worker/` — Rust. Tile rasterization and chunk meshing, consumed off a Redis dirty-chunk
+  stream. Stateless — scale it with `docker compose up --scale worker=N`, or run instances on
+  separate hardware pointed at the same Postgres/Redis/MinIO over a private network. Rendering
+  strategy is config-selectable (`RENDER_BACKEND=cpu`/`gpu`/`hybrid`, see `worker/.env.example`)
+  behind a `RenderBackend` trait (Phase 8): `cpu` (rayon, default) parallelizes across a whole
+  batch of dirty chunks; `gpu` additionally offloads tile shading and per-voxel
+  face-visibility extraction to a `wgpu` compute shader per chunk/section (greedy-mesh
+  merge/compaction stays CPU-only regardless — sequential/branchy, not GPU-parallel-friendly);
+  `hybrid` offloads only tile shading, keeping meshing on CPU. `gpu`/`hybrid` fall back to `cpu`
+  automatically (logged) if no compatible GPU adapter is found — see `worker/src/render/gpu.rs`'s
+  doc comment for the current known limitation (each tile/section is its own GPU dispatch, so at
+  small batch sizes `cpu` currently outruns `gpu`/`hybrid` — see `cargo run --release --example
+  benchmark`, in `worker/`, for real numbers on your own hardware).
 - `frontend/` — ElysiaJS + Pug + Tailwind 4 + Alpine.js. The public-facing pages (map viewer,
   chat, admin panel). Stateless — no DB access, calls `api` for anything server-rendered; the
   browser's own live map/chat/tile traffic talks to `api` directly, not proxied through here.
@@ -117,7 +124,9 @@ map markers with a `show`/hide toggle and an online count, right above the regio
 ## Running tests
 
 `worker`'s tests (`cargo test`, in `worker/`) are pure unit tests (greedy mesher, tile
-rasterizer, block-color palette) and need nothing running. `api`'s and `frontend`'s (`bun test`,
+rasterizer, block-color palette) and need nothing running — including the `gpu`/`hybrid` backend
+tests (Phase 8), which request a real GPU adapter and skip themselves (rather than failing) if
+none is found, so `cargo test` still passes on a machine with no GPU. `api`'s and `frontend`'s (`bun test`,
 in each directory) are integration tests against real infra — start it first:
 
 ```

docker-compose.yml

@@ -36,6 +36,17 @@ services:
     # Scale render capacity locally with, e.g.: docker compose up --scale worker=3
     # A `worker` can also run standalone on a separate machine, pointed at the same
     # postgres/redis/minio over a private network (VPN/Tailscale/LAN) — see README.
+    #
+    # RENDER_BACKEND=gpu/hybrid (worker/.env.example, Phase 8) needs the container to actually
+    # see a GPU — the nvidia-container-toolkit must be installed on the host, then uncomment:
+    # deploy:
+    #   resources:
+    #     reservations:
+    #       devices:
+    #         - driver: nvidia
+    #           count: 1
+    #           capabilities: [gpu]
+    # Default RENDER_BACKEND=cpu needs none of this — GPU rendering is opt-in, not required.
 
   postgres:
     image: postgres:17-alpine

worker/.dockerignore

@@ -0,0 +1,2 @@
+target/
+.env

worker/.env.example

@@ -9,7 +9,10 @@ MINIO_USE_SSL=false
 MINIO_ACCESS_KEY=mcmapper
 MINIO_SECRET_KEY=changeme-set-in-untracked-env
 
-# cpu | gpu | hybrid — see render::backend. Only `cpu` exists so far (Phase 8 adds gpu/hybrid).
+# cpu | gpu | hybrid (Phase 8) — see render/mod.rs's RenderBackend doc comment. `gpu` offloads
+# tile shading and per-voxel face-visibility extraction to a wgpu compute shader; `hybrid` offloads
+# only tile shading and keeps meshing on CPU. Unrecognized values, and `gpu`/`hybrid` on a machine
+# with no compatible GPU adapter, fall back to `cpu` with a logged warning rather than crashing.
 RENDER_BACKEND=cpu
 
 # Thread count for the rayon pool that renders a batch of dirty chunks in parallel (see

worker/Cargo.toml

@@ -15,6 +15,8 @@ aws-sdk-s3 = "1"
 aws-config = "1"
 aws-credential-types = "1"
 base64 = "0.22"
+wgpu = "23"
+bytemuck = { version = "1", features = ["derive"] }
 
 [profile.release]
 lto = true

worker/Dockerfile

@@ -1,4 +1,4 @@
-FROM rust:1.80 AS build
+FROM rust:1.97 AS build
 WORKDIR /app
 COPY . .
 RUN cargo build --release

worker/examples/benchmark.rs

@@ -0,0 +1,103 @@
+//! `cargo run --release --example benchmark` — compares tile-rasterization and section-meshing
+//! throughput across the `cpu`/`gpu`/`hybrid` `RenderBackend`s on a synthetic workload sized to
+//! match `main.rs`'s real batching (one `RENDER_PROFILE=server` xread batch's worth of chunks —
+//! see `config::batch_size_for`), run through the same rayon-parallel-batch shape `main.rs` uses.
+//! This is the Phase 8 plan's "measure throughput difference" verification step; `render::gpu`'s
+//! own tests already cover byte-for-byte output equivalence, so this only measures timing.
+use std::time::{Duration, Instant};
+
+use mcmapper_worker::config::{batch_size_for, Profile};
+use mcmapper_worker::mesh;
+use mcmapper_worker::render::{
+    ColumnPixel, CpuRenderBackend, GpuRenderBackend, HybridRenderBackend, RenderBackend,
+};
+use rayon::prelude::*;
+
+fn synthetic_columns() -> Vec<ColumnPixel> {
+    let mut columns = Vec::with_capacity(256);
+    for x in 0..16u8 {
+        for z in 0..16u8 {
+            let block_id = ((x as u16 + z as u16) % 20) + 1;
+            columns.push(ColumnPixel { local_x: x, local_z: z, block_id, block_meta: 0 });
+        }
+    }
+    columns
+}
+
+/// A wavy, partially-solid section — not fully solid (which greedy-merges to almost nothing) or
+/// fully empty (no work at all), so meshing does a realistic amount of face-mask/merge work.
+fn synthetic_section(seed: i32) -> [u16; 4096] {
+    let mut blocks = [0u16; 4096];
+    for y in 0..16i32 {
+        for z in 0..16i32 {
+            for x in 0..16i32 {
+                let height = 6 + (x + z + seed) % 6;
+                if y <= height {
+                    let block_id = 1 + ((x * 3 + z * 7 + y * 11 + seed) % 12) as u16;
+                    blocks[((y as usize) * 16 + z as usize) * 16 + x as usize] = block_id << 4;
+                }
+            }
+        }
+    }
+    blocks
+}
+
+fn bench_tiles(backend: &dyn RenderBackend, columns: &[ColumnPixel], batch: usize) -> Duration {
+    let start = Instant::now();
+    (0..batch).into_par_iter().for_each(|_| {
+        backend.rasterize_tile(columns).expect("rasterize_tile must not fail on synthetic input");
+    });
+    start.elapsed()
+}
+
+fn bench_meshing(backend: &dyn RenderBackend, sections: &[[u16; 4096]]) -> Duration {
+    let start = Instant::now();
+    sections.par_iter().for_each(|blocks| {
+        let _ = mesh::mesh_section(blocks, backend);
+    });
+    start.elapsed()
+}
+
+fn run_backend(name: &str, backend: &dyn RenderBackend, columns: &[ColumnPixel], sections: &[[u16; 4096]], batch: usize) {
+    // One warmup pass per stage so pipeline/shader compilation and first-dispatch driver
+    // overhead (real on GPU, one-time) doesn't get counted as steady-state throughput.
+    backend.rasterize_tile(columns).unwrap();
+    let _ = mesh::mesh_section(&sections[0], backend);
+
+    let tiles = bench_tiles(backend, columns, batch);
+    let meshes = bench_meshing(backend, sections);
+    println!(
+        "{name:<8} tiles: {batch} in {tiles:>9.2?} ({:>8.1}/s)   sections: {} in {meshes:>9.2?} ({:>8.1}/s)",
+        batch as f64 / tiles.as_secs_f64(),
+        sections.len(),
+        sections.len() as f64 / meshes.as_secs_f64(),
+    );
+}
+
+#[tokio::main]
+async fn main() {
+    let batch = batch_size_for(Profile::Server);
+    let columns = synthetic_columns();
+    let sections: Vec<[u16; 4096]> = (0..batch as i32).map(synthetic_section).collect();
+
+    println!("Synthetic workload: {batch} tiles + {} sections per backend (RENDER_PROFILE=server batch size)\n", sections.len());
+
+    let cpu = CpuRenderBackend::new();
+    run_backend("cpu", &cpu, &columns, &sections, batch);
+
+    match GpuRenderBackend::try_new().await {
+        Some(gpu) => {
+            println!("GPU adapter: {}", gpu.adapter_name());
+            run_backend("gpu", &gpu, &columns, &sections, batch);
+        }
+        None => println!("gpu      (skipped — no compatible GPU adapter found)"),
+    }
+
+    match GpuRenderBackend::try_new().await {
+        Some(gpu) => {
+            let hybrid = HybridRenderBackend::new(gpu);
+            run_backend("hybrid", &hybrid, &columns, &sections, batch);
+        }
+        None => println!("hybrid   (skipped — no compatible GPU adapter found)"),
+    }
+}

worker/src/lib.rs

@@ -0,0 +1,6 @@
+pub mod config;
+pub mod db;
+pub mod mesh;
+pub mod palette;
+pub mod render;
+pub mod storage;

worker/src/main.rs

@@ -1,18 +1,14 @@
-mod config;
-mod db;
-mod mesh;
-mod palette;
-mod render;
-mod storage;
-
 use std::collections::hash_map::DefaultHasher;
 use std::hash::{Hash, Hasher};
 
 use base64::{engine::general_purpose::STANDARD, Engine as _};
+use mcmapper_worker::{config, db, mesh, storage};
 use rayon::prelude::*;
 use redis::streams::{StreamReadOptions, StreamReadReply};
 use redis::AsyncCommands;
-use render::{ColumnPixel, CpuRenderBackend, RenderBackend};
+use mcmapper_worker::render::{
+    ColumnPixel, CpuRenderBackend, GpuRenderBackend, HybridRenderBackend, RenderBackend,
+};
 use uuid::Uuid;
 
 const DIRTY_CHUNK_STREAM: &str = "mcmapper:dirty-chunks";
@@ -50,15 +46,34 @@ async fn main() -> anyhow::Result<()> {
         .xgroup_create_mkstream(DIRTY_CHUNK_STREAM, CONSUMER_GROUP, "$")
         .await;
 
-    // Only `cpu` exists so far — `RENDER_BACKEND=gpu`/`hybrid` fall back with a warning until
-    // Phase 8 adds the wgpu path.
     let requested_backend = std::env::var("RENDER_BACKEND").unwrap_or_else(|_| "cpu".into());
-    if requested_backend != "cpu" {
-        eprintln!(
-            "[worker] RENDER_BACKEND={requested_backend} not implemented yet, falling back to cpu"
-        );
-    }
-    let backend: Box<dyn RenderBackend> = Box::new(CpuRenderBackend::new());
+    let backend: Box<dyn RenderBackend> = match requested_backend.as_str() {
+        "gpu" => match GpuRenderBackend::try_new().await {
+            Some(gpu) => {
+                println!("[worker] GPU adapter: {}", gpu.adapter_name());
+                Box::new(gpu)
+            }
+            None => {
+                eprintln!("[worker] RENDER_BACKEND=gpu requested but no compatible GPU adapter was found, falling back to cpu");
+                Box::new(CpuRenderBackend::new())
+            }
+        },
+        "hybrid" => match GpuRenderBackend::try_new().await {
+            Some(gpu) => {
+                println!("[worker] GPU adapter: {}", gpu.adapter_name());
+                Box::new(HybridRenderBackend::new(gpu))
+            }
+            None => {
+                eprintln!("[worker] RENDER_BACKEND=hybrid requested but no compatible GPU adapter was found, falling back to cpu");
+                Box::new(CpuRenderBackend::new())
+            }
+        },
+        "cpu" => Box::new(CpuRenderBackend::new()),
+        other => {
+            eprintln!("[worker] RENDER_BACKEND={other} not recognized (expected cpu/gpu/hybrid), falling back to cpu");
+            Box::new(CpuRenderBackend::new())
+        }
+    };
 
     // A single 16x16 tile is too small for rayon to help within itself, so the parallelism
     // target is "many chunks in flight at once" — see config.rs's doc comment. RENDER_THREADS/
@@ -240,7 +255,7 @@ fn render_chunk(data: &ChunkData, backend: &dyn RenderBackend) -> anyhow::Result
                 continue;
             }
         };
-        let mesh_buf = mesh::mesh_section(&blocks);
+        let mesh_buf = mesh::mesh_section(&blocks, backend);
         if mesh_buf.is_empty() {
             continue;
         }

worker/src/mesh.rs

@@ -1,4 +1,5 @@
 use crate::palette::color_for;
+use crate::render::RenderBackend;
 
 const SIZE: i32 = 16;
 
@@ -57,13 +58,6 @@ impl MeshBuffers {
     }
 }
 
-fn block_at(blocks: &[u16; 4096], x: i32, y: i32, z: i32) -> u16 {
-    if x < 0 || x >= SIZE || y < 0 || y >= SIZE || z < 0 || z >= SIZE {
-        return 0; // section boundary — treated as air, so boundary faces are always drawn
-    }
-    blocks[((y as usize) * 16 + z as usize) * 16 + x as usize]
-}
-
 /// Maps (axis, layer, u, v) to a 3D voxel coordinate. axis 0 fixes x, 1 fixes y, 2 fixes z.
 fn axis_pos(axis: usize, layer: i32, u: i32, v: i32) -> (i32, i32, i32) {
     match axis {
@@ -73,33 +67,38 @@ fn axis_pos(axis: usize, layer: i32, u: i32, v: i32) -> (i32, i32, i32) {
     }
 }
 
-pub fn mesh_section(blocks: &[u16; 4096]) -> MeshBuffers {
+/// Greedy-merges each of the 6 backend-computed face-visibility masks (see
+/// `render::RenderBackend::compute_face_masks`'s doc comment for what "face_index" means) into
+/// quads. The per-voxel visibility extraction may have run on GPU; this merge/compaction step is
+/// always CPU — it's sequential and branchy (each cell's fate depends on what its neighbors in
+/// the same pass already claimed), not a good GPU-parallel fit.
+pub fn mesh_section(blocks: &[u16; 4096], backend: &dyn RenderBackend) -> MeshBuffers {
+    let face_masks = backend.compute_face_masks(blocks);
     let mut buf = MeshBuffers::default();
+    let mut face_index = 0;
     for axis in 0..3 {
         for &dir in &[-1i32, 1i32] {
-            mesh_axis(blocks, axis, dir, &mut buf);
+            mesh_axis_from_visibility(&face_masks[face_index], axis, dir, &mut buf);
+            face_index += 1;
         }
     }
     buf
 }
 
-fn mesh_axis(blocks: &[u16; 4096], axis: usize, dir: i32, buf: &mut MeshBuffers) {
-    let mut mask = [[0u16; SIZE as usize]; SIZE as usize];
+fn voxel_at(visibility: &[u16; 4096], x: i32, y: i32, z: i32) -> u16 {
+    if x < 0 || x >= SIZE || y < 0 || y >= SIZE || z < 0 || z >= SIZE {
+        return 0;
+    }
+    visibility[((y as usize) * 16 + z as usize) * 16 + x as usize]
+}
 
+fn mesh_axis_from_visibility(visibility: &[u16; 4096], axis: usize, dir: i32, buf: &mut MeshBuffers) {
     for layer in 0..SIZE {
-        // Build the visibility mask for this layer: mask[u][v] = blockId if a face should be
-        // drawn there (the voxel is solid and the neighbor in `dir` along `axis` is air/boundary).
+        let mut mask = [[0u16; SIZE as usize]; SIZE as usize];
         for u in 0..SIZE {
             for v in 0..SIZE {
                 let (x, y, z) = axis_pos(axis, layer, u, v);
-                let block = block_at(blocks, x, y, z);
-                mask[u as usize][v as usize] = if block == 0 {
-                    0
-                } else {
-                    let (ox, oy, oz) = offset_along_axis(axis, dir);
-                    let neighbor = block_at(blocks, x + ox, y + oy, z + oz);
-                    if neighbor == 0 { block } else { 0 }
-                };
+                mask[u as usize][v as usize] = voxel_at(visibility, x, y, z);
             }
         }
 
@@ -108,14 +107,6 @@ fn mesh_axis(blocks: &[u16; 4096], axis: usize, dir: i32, buf: &mut MeshBuffers)
     }
 }
 
-fn offset_along_axis(axis: usize, dir: i32) -> (i32, i32, i32) {
-    match axis {
-        0 => (dir, 0, 0),
-        1 => (0, dir, 0),
-        _ => (0, 0, dir),
-    }
-}
-
 fn greedy_merge_and_emit(
     mask: &mut [[u16; SIZE as usize]; SIZE as usize],
     axis: usize,
@@ -207,11 +198,16 @@ fn emit_quad(
 #[cfg(test)]
 mod tests {
     use super::*;
+    use crate::render::CpuRenderBackend;
+
+    fn cpu() -> CpuRenderBackend {
+        CpuRenderBackend::new()
+    }
 
     #[test]
     fn empty_section_produces_no_geometry() {
         let blocks = [0u16; 4096];
-        let mesh = mesh_section(&blocks);
+        let mesh = mesh_section(&blocks, &cpu());
         assert!(mesh.is_empty());
         assert_eq!(mesh.positions.len(), 0);
     }
@@ -220,7 +216,7 @@ mod tests {
     fn single_voxel_produces_six_unmerged_quads() {
         let mut blocks = [0u16; 4096];
         blocks[((0 * 16 + 0) * 16 + 0) as usize] = (2 << 4) | 0; // grass at local (0,0,0)
-        let mesh = mesh_section(&blocks);
+        let mesh = mesh_section(&blocks, &cpu());
         assert_eq!(mesh.positions.len(), 6 * 4, "6 faces x 4 verts");
         assert_eq!(mesh.indices.len(), 6 * 6, "6 faces x 2 tris x 3 indices");
     }
@@ -234,7 +230,7 @@ mod tests {
         for b in blocks.iter_mut() {
             *b = (1 << 4) | 0; // stone everywhere
         }
-        let mesh = mesh_section(&blocks);
+        let mesh = mesh_section(&blocks, &cpu());
         assert_eq!(mesh.positions.len(), 6 * 4, "6 merged outer faces x 4 verts");
         assert_eq!(mesh.indices.len(), 6 * 6);
     }
@@ -254,10 +250,14 @@ mod tests {
                 }
             }
         }
-        let mesh = mesh_section(&blocks);
+        let backend = cpu();
+        let mesh = mesh_section(&blocks, &backend);
         // Just check the +y (top) face count via a targeted single-axis call.
+        // face_index 3 = axis 1 (y), dir +1 — see compute_face_masks's doc comment for the
+        // face_index = axis*2 + (dir==1) convention.
+        let face_masks = backend.compute_face_masks(&blocks);
         let mut buf = MeshBuffers::default();
-        mesh_axis(&blocks, 1, 1, &mut buf);
+        mesh_axis_from_visibility(&face_masks[3], 1, 1, &mut buf);
         assert_eq!(buf.positions.len(), solid_count * 4);
         let _ = mesh; // silence unused warning if full mesh isn't otherwise inspected
     }
@@ -266,7 +266,7 @@ mod tests {
     fn encode_round_trip_header() {
         let mut blocks = [0u16; 4096];
         blocks[0] = (2 << 4) | 0;
-        let mesh = mesh_section(&blocks);
+        let mesh = mesh_section(&blocks, &cpu());
         let bytes = mesh.encode();
         let vertex_count = u32::from_le_bytes(bytes[0..4].try_into().unwrap());
         let index_count = u32::from_le_bytes(bytes[4..8].try_into().unwrap());

worker/src/render/cpu.rs

@@ -1,15 +1,4 @@
-use image::{ImageEncoder, Rgb, RgbImage};
-use std::io::Cursor;
-
-use super::{ColumnPixel, RenderBackend};
-use crate::palette::color_for;
-
-/// Native tile resolution (one pixel per block within a chunk) before upscaling for display.
-const CHUNK_SIZE: u32 = 16;
-/// Upscale factor applied so tiles are a reasonable size for a Leaflet `tileSize: 256` layer —
-/// Phase 1 has one native zoom level (see api's tile route), so this is purely cosmetic, not a
-/// multi-resolution pyramid (that's Phase 2's job).
-const UPSCALE: u32 = 16;
+use super::{base_colors, upscale_and_encode, ColumnPixel, FaceMasks, RenderBackend};
 
 /// CPU rendering backend, parallelized with rayon. Thread count/task-batch granularity is
 /// config-tunable (`auto` via `std::thread::available_parallelism`, or a `server`/`consumer`
@@ -29,31 +18,53 @@ impl RenderBackend for CpuRenderBackend {
     }
 
     fn rasterize_tile(&self, columns: &[ColumnPixel]) -> anyhow::Result<Vec<u8>> {
-        let mut base = RgbImage::from_pixel(CHUNK_SIZE, CHUNK_SIZE, Rgb([30, 30, 40]));
-        for col in columns {
-            if col.local_x as u32 >= CHUNK_SIZE || col.local_z as u32 >= CHUNK_SIZE {
-                continue;
-            }
-            let color = color_for(col.block_id, col.block_meta);
-            base.put_pixel(col.local_x as u32, col.local_z as u32, Rgb(color));
-        }
+        upscale_and_encode(&base_colors(columns))
+    }
 
-        let mut upscaled = RgbImage::new(CHUNK_SIZE * UPSCALE, CHUNK_SIZE * UPSCALE);
-        for y in 0..upscaled.height() {
-            for x in 0..upscaled.width() {
-                let src = base.get_pixel(x / UPSCALE, y / UPSCALE);
-                upscaled.put_pixel(x, y, *src);
+    fn compute_face_masks(&self, blocks: &[u16; 4096]) -> FaceMasks {
+        let mut out: FaceMasks = [[0u16; 4096]; 6];
+        let mut face_index = 0;
+        for axis in 0..3 {
+            for &dir in &[-1i32, 1i32] {
+                let (ox, oy, oz) = offset_along_axis(axis, dir);
+                for y in 0..16i32 {
+                    for z in 0..16i32 {
+                        for x in 0..16i32 {
+                            let block = block_at(blocks, x, y, z);
+                            let value = if block == 0 {
+                                0
+                            } else if block_at(blocks, x + ox, y + oy, z + oz) == 0 {
+                                block
+                            } else {
+                                0
+                            };
+                            out[face_index][voxel_index(x, y, z)] = value;
+                        }
+                    }
+                }
+                face_index += 1;
             }
         }
+        out
+    }
+}
 
-        let mut bytes = Vec::new();
-        image::codecs::png::PngEncoder::new(&mut Cursor::new(&mut bytes)).write_image(
-            upscaled.as_raw(),
-            upscaled.width(),
-            upscaled.height(),
-            image::ExtendedColorType::Rgb8,
-        )?;
-        Ok(bytes)
+pub(crate) fn voxel_index(x: i32, y: i32, z: i32) -> usize {
+    ((y as usize) * 16 + z as usize) * 16 + x as usize
+}
+
+pub(crate) fn block_at(blocks: &[u16; 4096], x: i32, y: i32, z: i32) -> u16 {
+    if x < 0 || x >= 16 || y < 0 || y >= 16 || z < 0 || z >= 16 {
+        return 0; // section boundary — treated as air, so boundary faces are always drawn
+    }
+    blocks[voxel_index(x, y, z)]
+}
+
+pub(crate) fn offset_along_axis(axis: usize, dir: i32) -> (i32, i32, i32) {
+    match axis {
+        0 => (dir, 0, 0),
+        1 => (0, dir, 0),
+        _ => (0, 0, dir),
     }
 }
 
@@ -61,8 +72,9 @@ impl RenderBackend for CpuRenderBackend {
 mod tests {
     use super::*;
     use crate::palette::color_for;
+    use image::Rgb;
 
-    fn decode(bytes: &[u8]) -> RgbImage {
+    fn decode(bytes: &[u8]) -> image::RgbImage {
         image::load_from_memory(bytes).expect("worker must always produce a decodable PNG").to_rgb8()
     }
 
@@ -116,4 +128,56 @@ mod tests {
         let result = backend.rasterize_tile(&columns);
         assert!(result.is_ok());
     }
+
+    #[test]
+    fn empty_section_has_no_visible_faces_in_any_direction() {
+        let backend = CpuRenderBackend::new();
+        let blocks = [0u16; 4096];
+        let masks = backend.compute_face_masks(&blocks);
+        for mask in masks {
+            assert!(mask.iter().all(|&v| v == 0));
+        }
+    }
+
+    #[test]
+    fn a_lone_voxel_exposes_all_six_faces() {
+        let backend = CpuRenderBackend::new();
+        let mut blocks = [0u16; 4096];
+        let block = (2 << 4) | 0; // grass
+        blocks[voxel_index(3, 3, 3)] = block;
+        let masks = backend.compute_face_masks(&blocks);
+        for (face_index, mask) in masks.iter().enumerate() {
+            let exposed = mask.iter().filter(|&&v| v != 0).count();
+            assert_eq!(exposed, 1, "face {face_index} should expose exactly the lone voxel");
+            assert_eq!(mask[voxel_index(3, 3, 3)], block);
+        }
+    }
+
+    #[test]
+    fn two_adjacent_voxels_hide_the_faces_between_them() {
+        let backend = CpuRenderBackend::new();
+        let mut blocks = [0u16; 4096];
+        let stone = (1 << 4) | 0;
+        blocks[voxel_index(4, 0, 0)] = stone;
+        blocks[voxel_index(5, 0, 0)] = stone;
+        let masks = backend.compute_face_masks(&blocks);
+        // face_index 0 = axis 0 (x), dir -1: voxel at x=5 sees a solid neighbor at x=4, so hidden.
+        assert_eq!(masks[0][voxel_index(5, 0, 0)], 0);
+        // face_index 1 = axis 0 (x), dir +1: voxel at x=4 sees a solid neighbor at x=5, so hidden.
+        assert_eq!(masks[1][voxel_index(4, 0, 0)], 0);
+        // The outward-facing sides of the pair are still exposed.
+        assert_eq!(masks[0][voxel_index(4, 0, 0)], stone);
+        assert_eq!(masks[1][voxel_index(5, 0, 0)], stone);
+    }
+
+    #[test]
+    fn a_voxel_on_the_section_boundary_treats_the_boundary_as_air() {
+        let backend = CpuRenderBackend::new();
+        let mut blocks = [0u16; 4096];
+        let stone = (1 << 4) | 0;
+        blocks[voxel_index(0, 0, 0)] = stone;
+        let masks = backend.compute_face_masks(&blocks);
+        // face_index 0 = axis 0, dir -1: neighbor at x=-1 is out of bounds -> treated as air -> exposed.
+        assert_eq!(masks[0][voxel_index(0, 0, 0)], stone);
+    }
 }

worker/src/render/gpu.rs

@@ -0,0 +1,356 @@
+use std::sync::mpsc;
+
+use wgpu::util::DeviceExt;
+
+use super::{base_colors, encode_png, ColumnPixel, FaceMasks, RenderBackend};
+
+const TILE_SHADER: &str = r#"
+@group(0) @binding(0) var<storage, read> base_colors: array<u32, 256>;
+@group(0) @binding(1) var<storage, read_write> out_pixels: array<u32, 65536>;
+
+// One invocation per output pixel (256x256 = the CHUNK_SIZE*UPSCALE tile) — nearest-neighbor
+// samples the 16x16 base color grid, mirroring cpu_upscale's loop exactly so both backends
+// produce byte-identical tiles for the same input.
+@compute @workgroup_size(16, 16, 1)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    let x = gid.x;
+    let y = gid.y;
+    if (x >= 256u || y >= 256u) {
+        return;
+    }
+    let bx = x / 16u;
+    let by = y / 16u;
+    out_pixels[y * 256u + x] = base_colors[by * 16u + bx];
+}
+"#;
+
+const FACE_MASK_SHADER: &str = r#"
+@group(0) @binding(0) var<storage, read> blocks: array<u32, 4096>;
+@group(0) @binding(1) var<storage, read_write> out_masks: array<u32, 24576>;
+
+fn voxel_index(x: i32, y: i32, z: i32) -> u32 {
+    return u32((y * 16 + z) * 16 + x);
+}
+
+fn block_at(x: i32, y: i32, z: i32) -> u32 {
+    if (x < 0 || x >= 16 || y < 0 || y >= 16 || z < 0 || z >= 16) {
+        return 0u;
+    }
+    return blocks[voxel_index(x, y, z)];
+}
+
+// One invocation per voxel (16x16x16 = one section) — computes all 6 face directions for that
+// voxel in a single pass rather than 6 separate dispatches, since the neighbor lookups are cheap
+// and this keeps dispatch overhead to one call per section. face_index = axis*2 + (dir==+1), axis
+// 0=x/1=y/2=z, matching RenderBackend::compute_face_masks's doc comment. workgroup_size 4x4x4=64
+// (not e.g. 8x8x8=512) to stay within wgpu's portable default 256-invocations-per-workgroup limit
+// — this only targets one 16x16x16 section per dispatch anyway, so the smaller workgroup costs
+// nothing but a few extra (very cheap) group launches.
+@compute @workgroup_size(4, 4, 4)
+fn main(@builtin(global_invocation_id) gid: vec3<u32>) {
+    if (gid.x >= 16u || gid.y >= 16u || gid.z >= 16u) {
+        return;
+    }
+    let x = i32(gid.x);
+    let y = i32(gid.y);
+    let z = i32(gid.z);
+    let block = block_at(x, y, z);
+    let idx = voxel_index(x, y, z);
+
+    var offsets_x = array<i32, 6>(-1, 1, 0, 0, 0, 0);
+    var offsets_y = array<i32, 6>(0, 0, -1, 1, 0, 0);
+    var offsets_z = array<i32, 6>(0, 0, 0, 0, -1, 1);
+
+    for (var f = 0u; f < 6u; f = f + 1u) {
+        var value = 0u;
+        if (block != 0u) {
+            let nx = x + offsets_x[f];
+            let ny = y + offsets_y[f];
+            let nz = z + offsets_z[f];
+            if (block_at(nx, ny, nz) == 0u) {
+                value = block;
+            }
+        }
+        out_masks[f * 4096u + idx] = value;
+    }
+}
+"#;
+
+/// GPU rendering backend: offloads the two embarrassingly-parallel per-pixel/per-voxel steps
+/// (tile upscale/shading, per-voxel face-visibility extraction) to wgpu compute shaders, per the
+/// plan's "partial GPU rendering" design — greedy-mesh merge/compaction (`mesh.rs`) stays on CPU
+/// regardless of which `RenderBackend` produced the face masks, since that step is sequential.
+///
+/// Each `rasterize_tile`/`compute_face_masks` call does its own buffer upload, dispatch, and
+/// blocking readback — correct and simple, but it means per-call dispatch overhead isn't
+/// amortized across a whole render batch the way `main.rs`'s rayon pool amortizes CPU work
+/// across many chunks at once. Batching multiple chunks into one dispatch is a natural follow-up
+/// if profiling shows dispatch overhead dominates at real workload sizes; not attempted here to
+/// keep this Phase 8 change reviewable and to keep `RenderBackend`'s existing per-tile/per-section
+/// trait shape (shared with the CPU backend) intact.
+pub struct GpuRenderBackend {
+    device: wgpu::Device,
+    queue: wgpu::Queue,
+    tile_pipeline: wgpu::ComputePipeline,
+    tile_bind_layout: wgpu::BindGroupLayout,
+    face_pipeline: wgpu::ComputePipeline,
+    face_bind_layout: wgpu::BindGroupLayout,
+    adapter_name: String,
+}
+
+impl GpuRenderBackend {
+    /// Requests a high-performance adapter and device; returns `None` (never panics/errors) if
+    /// no compatible GPU is found, so callers (`main.rs`) can fall back to CPU with a logged
+    /// warning rather than crashing — see the plan's Phase 8 auto-fallback requirement.
+    pub async fn try_new() -> Option<Self> {
+        let instance = wgpu::Instance::new(wgpu::InstanceDescriptor {
+            backends: wgpu::Backends::all(),
+            ..Default::default()
+        });
+        let adapter = instance
+            .request_adapter(&wgpu::RequestAdapterOptions {
+                power_preference: wgpu::PowerPreference::HighPerformance,
+                compatible_surface: None,
+                force_fallback_adapter: false,
+            })
+            .await?;
+        let adapter_name = adapter.get_info().name;
+        let (device, queue) = adapter
+            .request_device(
+                &wgpu::DeviceDescriptor { label: Some("mcmapper-worker-gpu"), ..Default::default() },
+                None,
+            )
+            .await
+            .ok()?;
+
+        let (tile_pipeline, tile_bind_layout) =
+            build_pipeline(&device, "tile-shading", TILE_SHADER);
+        let (face_pipeline, face_bind_layout) =
+            build_pipeline(&device, "face-mask-extraction", FACE_MASK_SHADER);
+
+        Some(Self { device, queue, tile_pipeline, tile_bind_layout, face_pipeline, face_bind_layout, adapter_name })
+    }
+
+    /// The selected GPU's name (e.g. "NVIDIA GeForce RTX 3090"), for `main.rs`'s startup log line.
+    pub fn adapter_name(&self) -> &str {
+        &self.adapter_name
+    }
+
+    /// Runs a compute shader over one input storage buffer producing one output storage buffer,
+    /// blocking until the result is read back. `input`/`output_len` are in `u32` elements.
+    fn dispatch_u32(
+        &self,
+        pipeline: &wgpu::ComputePipeline,
+        bind_layout: &wgpu::BindGroupLayout,
+        input: &[u32],
+        output_len: usize,
+        workgroups: (u32, u32, u32),
+    ) -> Vec<u32> {
+        let input_buf = self.device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
+            label: Some("mcmapper-gpu-input"),
+            contents: bytemuck::cast_slice(input),
+            usage: wgpu::BufferUsages::STORAGE,
+        });
+        let output_size = (output_len * std::mem::size_of::<u32>()) as u64;
+        let output_buf = self.device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("mcmapper-gpu-output"),
+            size: output_size,
+            usage: wgpu::BufferUsages::STORAGE | wgpu::BufferUsages::COPY_SRC,
+            mapped_at_creation: false,
+        });
+        let staging_buf = self.device.create_buffer(&wgpu::BufferDescriptor {
+            label: Some("mcmapper-gpu-staging"),
+            size: output_size,
+            usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST,
+            mapped_at_creation: false,
+        });
+
+        let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor {
+            label: Some("mcmapper-gpu-bind-group"),
+            layout: bind_layout,
+            entries: &[
+                wgpu::BindGroupEntry { binding: 0, resource: input_buf.as_entire_binding() },
+                wgpu::BindGroupEntry { binding: 1, resource: output_buf.as_entire_binding() },
+            ],
+        });
+
+        let mut encoder = self.device.create_command_encoder(&wgpu::CommandEncoderDescriptor {
+            label: Some("mcmapper-gpu-encoder"),
+        });
+        {
+            let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor {
+                label: Some("mcmapper-gpu-pass"),
+                timestamp_writes: None,
+            });
+            pass.set_pipeline(pipeline);
+            pass.set_bind_group(0, &bind_group, &[]);
+            pass.dispatch_workgroups(workgroups.0, workgroups.1, workgroups.2);
+        }
+        encoder.copy_buffer_to_buffer(&output_buf, 0, &staging_buf, 0, output_size);
+        self.queue.submit(Some(encoder.finish()));
+
+        let slice = staging_buf.slice(..);
+        let (tx, rx) = mpsc::channel();
+        slice.map_async(wgpu::MapMode::Read, move |result| {
+            let _ = tx.send(result);
+        });
+        // wgpu's futures only resolve once the device is polled — there's no background executor
+        // driving that here, so poll synchronously until the map callback above has fired. This
+        // runs on whatever thread called us (a rayon worker thread in `main.rs`'s batch stage,
+        // never the tokio reactor), so blocking here doesn't stall other async work.
+        self.device.poll(wgpu::Maintain::Wait);
+        rx.recv().expect("map_async callback must fire after Maintain::Wait").expect("buffer mapping failed");
+
+        let data = slice.get_mapped_range();
+        let result: Vec<u32> = bytemuck::cast_slice(&data).to_vec();
+        drop(data);
+        staging_buf.unmap();
+        result
+    }
+}
+
+impl RenderBackend for GpuRenderBackend {
+    fn name(&self) -> &'static str {
+        "gpu"
+    }
+
+    fn rasterize_tile(&self, columns: &[ColumnPixel]) -> anyhow::Result<Vec<u8>> {
+        let base = base_colors(columns);
+        let packed_input: Vec<u32> = base.iter().map(|&[r, g, b]| pack_rgba(r, g, b)).collect();
+        let packed_output = self.dispatch_u32(
+            &self.tile_pipeline,
+            &self.tile_bind_layout,
+            &packed_input,
+            65536,
+            (16, 16, 1),
+        );
+
+        let mut rgb = Vec::with_capacity(65536 * 3);
+        for pixel in packed_output {
+            let [r, g, b, _a] = pixel.to_le_bytes();
+            rgb.extend_from_slice(&[r, g, b]);
+        }
+        encode_png(&rgb, 256, 256)
+    }
+
+    fn compute_face_masks(&self, blocks: &[u16; 4096]) -> FaceMasks {
+        let packed_input: Vec<u32> = blocks.iter().map(|&b| b as u32).collect();
+        let packed_output =
+            self.dispatch_u32(&self.face_pipeline, &self.face_bind_layout, &packed_input, 24576, (4, 4, 4));
+
+        let mut out: FaceMasks = [[0u16; 4096]; 6];
+        for face_index in 0..6 {
+            for voxel in 0..4096 {
+                out[face_index][voxel] = packed_output[face_index * 4096 + voxel] as u16;
+            }
+        }
+        out
+    }
+}
+
+fn pack_rgba(r: u8, g: u8, b: u8) -> u32 {
+    u32::from_le_bytes([r, g, b, 255])
+}
+
+fn build_pipeline(
+    device: &wgpu::Device,
+    label: &str,
+    source: &str,
+) -> (wgpu::ComputePipeline, wgpu::BindGroupLayout) {
+    let module = device.create_shader_module(wgpu::ShaderModuleDescriptor {
+        label: Some(label),
+        source: wgpu::ShaderSource::Wgsl(source.into()),
+    });
+    let bind_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor {
+        label: Some(label),
+        entries: &[
+            wgpu::BindGroupLayoutEntry {
+                binding: 0,
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::Buffer {
+                    ty: wgpu::BufferBindingType::Storage { read_only: true },
+                    has_dynamic_offset: false,
+                    min_binding_size: None,
+                },
+                count: None,
+            },
+            wgpu::BindGroupLayoutEntry {
+                binding: 1,
+                visibility: wgpu::ShaderStages::COMPUTE,
+                ty: wgpu::BindingType::Buffer {
+                    ty: wgpu::BufferBindingType::Storage { read_only: false },
+                    has_dynamic_offset: false,
+                    min_binding_size: None,
+                },
+                count: None,
+            },
+        ],
+    });
+    let pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
+        label: Some(label),
+        bind_group_layouts: &[&bind_layout],
+        push_constant_ranges: &[],
+    });
+    let pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor {
+        label: Some(label),
+        layout: Some(&pipeline_layout),
+        module: &module,
+        entry_point: Some("main"),
+        compilation_options: Default::default(),
+        cache: None,
+    });
+    (pipeline, bind_layout)
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+    use crate::render::CpuRenderBackend;
+
+    /// Real hardware isn't guaranteed in every environment this runs in — skip rather than fail
+    /// when no compatible adapter is found, matching `main.rs`'s own runtime fallback behavior.
+    async fn gpu_or_skip() -> Option<GpuRenderBackend> {
+        let backend = GpuRenderBackend::try_new().await;
+        if backend.is_none() {
+            eprintln!("[test] no compatible GPU adapter found, skipping GPU test");
+        }
+        backend
+    }
+
+    #[tokio::test]
+    async fn rasterize_tile_matches_cpu_output_byte_for_byte() {
+        let Some(gpu) = gpu_or_skip().await else { return };
+        let cpu = CpuRenderBackend::new();
+        let columns = vec![
+            ColumnPixel { local_x: 0, local_z: 0, block_id: 1, block_meta: 0 },
+            ColumnPixel { local_x: 5, local_z: 9, block_id: 2, block_meta: 0 },
+            ColumnPixel { local_x: 15, local_z: 15, block_id: 12, block_meta: 0 },
+        ];
+        let gpu_bytes = gpu.rasterize_tile(&columns).unwrap();
+        let cpu_bytes = cpu.rasterize_tile(&columns).unwrap();
+        assert_eq!(gpu_bytes, cpu_bytes);
+    }
+
+    #[tokio::test]
+    async fn compute_face_masks_matches_cpu_output_exactly() {
+        let Some(gpu) = gpu_or_skip().await else { return };
+        let cpu = CpuRenderBackend::new();
+        let mut blocks = [0u16; 4096];
+        blocks[0] = (1 << 4) | 0;
+        blocks[1] = (1 << 4) | 0; // adjacent along x, hides the face between them
+        blocks[4095] = (2 << 4) | 0; // far corner, isolated
+
+        let gpu_masks = gpu.compute_face_masks(&blocks);
+        let cpu_masks = cpu.compute_face_masks(&blocks);
+        assert_eq!(gpu_masks, cpu_masks);
+    }
+
+    #[tokio::test]
+    async fn name_reports_gpu() {
+        let Some(gpu) = gpu_or_skip().await else { return };
+        assert_eq!(gpu.name(), "gpu");
+        // adapter_name is only ever read for the startup log line in main.rs — assert it's at
+        // least populated so that log line isn't silently empty.
+        assert!(!gpu.adapter_name.is_empty());
+    }
+}

worker/src/render/hybrid.rs

@@ -0,0 +1,61 @@
+use super::{ColumnPixel, CpuRenderBackend, FaceMasks, GpuRenderBackend, RenderBackend};
+
+/// A middle ground between `cpu` and `gpu`: offloads tile shading (the cheaper, more uniform of
+/// the two GPU-candidate steps — one dispatch, always exactly 65536 output pixels) to the GPU,
+/// but keeps per-voxel face-visibility extraction on CPU. Meshing is comparatively rare (only
+/// dirty *sections* need remeshing, whereas every dirty chunk needs its tile re-rasterized) and
+/// each section's face-mask dispatch has more fixed overhead relative to its 4096-voxel workload
+/// than the tile shader's 65536-pixel one — `hybrid` is for deployments where a GPU is available
+/// but not obviously worth it for the smaller, more frequent meshing workload specifically.
+pub struct HybridRenderBackend {
+    gpu: GpuRenderBackend,
+    cpu: CpuRenderBackend,
+}
+
+impl HybridRenderBackend {
+    pub fn new(gpu: GpuRenderBackend) -> Self {
+        Self { gpu, cpu: CpuRenderBackend::new() }
+    }
+}
+
+impl RenderBackend for HybridRenderBackend {
+    fn name(&self) -> &'static str {
+        "hybrid"
+    }
+
+    fn rasterize_tile(&self, columns: &[ColumnPixel]) -> anyhow::Result<Vec<u8>> {
+        self.gpu.rasterize_tile(columns)
+    }
+
+    fn compute_face_masks(&self, blocks: &[u16; 4096]) -> FaceMasks {
+        self.cpu.compute_face_masks(blocks)
+    }
+}
+
+#[cfg(test)]
+mod tests {
+    use super::*;
+
+    #[tokio::test]
+    async fn hybrid_uses_gpu_for_tiles_and_cpu_for_face_masks() {
+        let Some(gpu) = GpuRenderBackend::try_new().await else {
+            eprintln!("[test] no compatible GPU adapter found, skipping");
+            return;
+        };
+        let cpu = CpuRenderBackend::new();
+        let hybrid = HybridRenderBackend::new(gpu);
+
+        assert_eq!(hybrid.name(), "hybrid");
+
+        let columns = vec![ColumnPixel { local_x: 3, local_z: 3, block_id: 2, block_meta: 0 }];
+        assert_eq!(hybrid.rasterize_tile(&columns).unwrap(), {
+            // Re-fetch a fresh GPU backend to compare against, since `hybrid` consumed the first.
+            let Some(gpu2) = GpuRenderBackend::try_new().await else { unreachable!() };
+            gpu2.rasterize_tile(&columns).unwrap()
+        });
+
+        let mut blocks = [0u16; 4096];
+        blocks[0] = (1 << 4) | 0;
+        assert_eq!(hybrid.compute_face_masks(&blocks), cpu.compute_face_masks(&blocks));
+    }
+}

worker/src/render/mod.rs

@@ -1,6 +1,13 @@
 mod cpu;
+mod gpu;
+mod hybrid;
 
 pub use cpu::CpuRenderBackend;
+pub use gpu::GpuRenderBackend;
+pub use hybrid::HybridRenderBackend;
+
+use image::{ImageEncoder, RgbImage};
+use std::io::Cursor;
 
 /// One rendered column within a chunk, in chunk-local coordinates (0..16).
 pub struct ColumnPixel {
@@ -10,14 +17,85 @@ pub struct ColumnPixel {
     pub block_meta: u8,
 }
 
-/// Swappable rendering strategy (CPU via rayon / GPU via wgpu / hybrid) selected at startup
-/// via the `RENDER_BACKEND` env var. Only the CPU path exists so far; GPU/hybrid land in
-/// Phase 8, with automatic fallback to CPU if `gpu`/`hybrid` is requested but no compatible
-/// GPU is found.
+/// Native tile resolution (one pixel per block within a chunk) before upscaling for display.
+pub const CHUNK_SIZE: u32 = 16;
+/// Upscale factor applied so tiles are a reasonable size for a Leaflet `tileSize: 256` layer —
+/// Phase 1 has one native zoom level (see api's tile route), so this is purely cosmetic, not a
+/// multi-resolution pyramid (that's Phase 2's job).
+pub const UPSCALE: u32 = 16;
+
+/// Per-voxel face-visibility extraction for one 16x16x16 section: for each of the 6
+/// axis+direction combinations (index = `axis*2 + (dir==1 as usize)`, axis 0=x/1=y/2=z, matching
+/// `mesh.rs`'s own iteration order), a full section-sized grid (same x/y/z-major flat layout as
+/// the `blocks` input) where each entry is the voxel's own blockId if a face should be drawn in
+/// that direction (the voxel is solid and its neighbor along that axis/direction is air or out of
+/// bounds) or 0 otherwise. This is the embarrassingly-parallel per-voxel step Phase 8 targets for
+/// GPU offload; the caller (`mesh::mesh_section`) still does greedy-mesh merge/compaction on CPU
+/// regardless of which backend produced these masks, since that step is sequential/branchy — see
+/// the plan's "partial GPU rendering" note.
+pub type FaceMasks = [[u16; 4096]; 6];
+
+/// Swappable rendering strategy (CPU via rayon / GPU via wgpu compute / hybrid) selected at
+/// startup via the `RENDER_BACKEND` env var, with automatic fallback to CPU (logged) if `gpu`/
+/// `hybrid` is requested but no compatible GPU adapter is found — see `main.rs`.
 pub trait RenderBackend: Send + Sync {
     fn name(&self) -> &'static str;
 
     /// Rasterize one chunk's worth of columns (up to 256, sparse if the chunk isn't fully
     /// synced yet) into a single-resolution top-down PNG tile, returned as encoded bytes.
     fn rasterize_tile(&self, columns: &[ColumnPixel]) -> anyhow::Result<Vec<u8>>;
+
+    fn compute_face_masks(&self, blocks: &[u16; 4096]) -> FaceMasks;
+}
+
+/// Resolves a batch of columns into the native-resolution (16x16) base color grid, flat in
+/// row-major (z-major, matching `ColumnPixel`'s local_x/local_z) order. Shared by every backend
+/// — palette lookup is cheap (256 entries at most) and keeping it in one place avoids maintaining
+/// two copies of `palette::color_for`'s logic (one in Rust, one duplicated into WGSL).
+pub(crate) fn base_colors(columns: &[ColumnPixel]) -> [[u8; 3]; 256] {
+    let mut base = [crate::palette::color_for(0, 0); 256]; // air everywhere until overwritten
+    for col in columns {
+        if col.local_x as u32 >= CHUNK_SIZE || col.local_z as u32 >= CHUNK_SIZE {
+            continue;
+        }
+        let color = crate::palette::color_for(col.block_id, col.block_meta);
+        base[(col.local_z as usize) * 16 + col.local_x as usize] = color;
+    }
+    base
+}
+
+/// Nearest-neighbor-upscales a 16x16 base color grid into a `CHUNK_SIZE*UPSCALE` square, on the
+/// CPU. The GPU backend performs the equivalent expansion via a compute dispatch instead (see
+/// `gpu.rs`) and feeds its readback through `encode_png` below — both paths are expected to
+/// produce byte-identical output for the same input, only *how* the loop runs differs.
+pub(crate) fn cpu_upscale(base: &[[u8; 3]; 256]) -> RgbImage {
+    let size = CHUNK_SIZE * UPSCALE;
+    let mut upscaled = RgbImage::new(size, size);
+    for y in 0..size {
+        for x in 0..size {
+            let bx = (x / UPSCALE) as usize;
+            let by = (y / UPSCALE) as usize;
+            let [r, g, b] = base[by * 16 + bx];
+            upscaled.put_pixel(x, y, image::Rgb([r, g, b]));
+        }
+    }
+    upscaled
+}
+
+pub(crate) fn upscale_and_encode(base: &[[u8; 3]; 256]) -> anyhow::Result<Vec<u8>> {
+    let upscaled = cpu_upscale(base);
+    encode_png(upscaled.as_raw(), upscaled.width(), upscaled.height())
+}
+
+/// PNG-encodes a flat RGB8 pixel buffer. Shared by every backend so tile output is
+/// byte-identical regardless of which one produced the raw pixels.
+pub(crate) fn encode_png(rgb: &[u8], width: u32, height: u32) -> anyhow::Result<Vec<u8>> {
+    let mut bytes = Vec::new();
+    image::codecs::png::PngEncoder::new(&mut Cursor::new(&mut bytes)).write_image(
+        rgb,
+        width,
+        height,
+        image::ExtendedColorType::Rgb8,
+    )?;
+    Ok(bytes)
 }