diff --git a/README.md b/README.md index 0903fbf..99a4d02 100644 --- a/README.md +++ b/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: ``` diff --git a/docker-compose.yml b/docker-compose.yml index caa7f86..d01362e 100644 --- a/docker-compose.yml +++ b/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 diff --git a/worker/.dockerignore b/worker/.dockerignore new file mode 100644 index 0000000..14ee500 --- /dev/null +++ b/worker/.dockerignore @@ -0,0 +1,2 @@ +target/ +.env diff --git a/worker/.env.example b/worker/.env.example index 80561c6..758ae81 100644 --- a/worker/.env.example +++ b/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 diff --git a/worker/Cargo.lock b/worker/Cargo.lock index 0773541..fcc9e27 100644 --- a/worker/Cargo.lock +++ b/worker/Cargo.lock @@ -14,6 +14,15 @@ version = "0.2.21" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "683d7910e743518b0e34f1186f92494becacb047c7b6bf616c96772180fef923" +[[package]] +name = "android_system_properties" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ae221649c9976a6f6c56ae1facf410f3ddb33cc661c4b7b61020a912d4237fbc" +dependencies = [ + "libc", +] + [[package]] name = "anyhow" version = "1.0.104" @@ -29,6 +38,21 @@ dependencies = [ "rustversion", ] +[[package]] +name = "arrayvec" +version = "0.7.8" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d3fb67a6e08acf24fdeccbac2cb6ac4305825bd1f117462e0e6f2f193345ad56" + +[[package]] +name = "ash" +version = "0.38.0+1.3.281" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0bb44936d800fea8f016d7f2311c6a4f97aebd5dc86f09906139ec848cf3a46f" +dependencies = [ + "libloading", +] + [[package]] name = "async-trait" version = "0.1.92" @@ -573,6 +597,27 @@ version = "1.8.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "2af50177e190e07a26ab74f8b1efbfe2ef87da2116221318cb1c2e82baf7de06" +[[package]] +name = "bit-set" +version = "0.8.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08807e080ed7f9d5433fa9b275196cfc35414f66a0c79d864dc51a0d825231a3" +dependencies = [ + "bit-vec", +] + +[[package]] +name = "bit-vec" +version = "0.8.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "5e764a1d40d510daf35e07be9eb06e75770908c27d411ee6c92109c9840eaaf7" + +[[package]] +name = "bitflags" +version = "1.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bef38d45163c2f1dde094a7dfd33ccf595c92905c8f8f4fdc18d06fb1037718a" + [[package]] name = "bitflags" version = "2.13.1" @@ -582,6 +627,12 @@ dependencies = [ "serde_core", ] +[[package]] +name = "block" +version = "0.1.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "0d8c1fef690941d3e7788d328517591fecc684c084084702d6ff1641e993699a" + [[package]] name = "block-buffer" version = "0.10.4" @@ -611,6 +662,20 @@ name = "bytemuck" version = "1.25.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "95832e849adfb21180ccb6826a99da14e5d266ae5c2e668e1602cf234f153797" +dependencies = [ + "bytemuck_derive", +] + +[[package]] +name = "bytemuck_derive" +version = "1.12.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fc0e56a716f1e132ff6bf4bdac1c944a3fcdc1cae65f70a4a2a1ac3b401d2d1f" +dependencies = [ + "proc-macro2", + "quote", + "syn 3.0.3", +] [[package]] name = "byteorder" @@ -658,6 +723,12 @@ version = "1.0.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "9330f8b2ff13f34540b44e946ef35111825727b38d33286ef986142615121801" +[[package]] +name = "cfg_aliases" +version = "0.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fd16c4719339c4530435d38e511904438d07cce7950afa3718a84ac36c10e89e" + [[package]] name = "cmake" version = "0.1.58" @@ -673,6 +744,16 @@ version = "0.5.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0c9ea0ac24bc397ab3c98583a3c9ba74fa56b09a4449bbe172b9b1ddb016027a" +[[package]] +name = "codespan-reporting" +version = "0.11.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3538270d33cc669650c4b093848450d380def10c331d38c768e34cac80576e6e" +dependencies = [ + "termcolor", + "unicode-width", +] + [[package]] name = "combine" version = "4.6.7" @@ -699,6 +780,16 @@ version = "0.10.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "a6ef517f0926dd24a1582492c791b6a4818a4d94e789a334894aa15b0d12f55c" +[[package]] +name = "core-foundation" +version = "0.9.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "91e195e091a93c46f7102ec7818a2aa394e1e1771c3ab4825963fa03e45afb8f" +dependencies = [ + "core-foundation-sys", + "libc", +] + [[package]] name = "core-foundation" version = "0.10.1" @@ -715,6 +806,17 @@ version = "0.8.7" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "773648b94d0e5d620f64f280777445740e61fe701025087ec8b57f45c791888b" +[[package]] +name = "core-graphics-types" +version = "0.1.3" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "45390e6114f68f718cc7a830514a96f903cccd70d02a8f6d9f643ac4ba45afaf" +dependencies = [ + "bitflags 1.3.2", + "core-foundation 0.9.4", + "libc", +] + [[package]] name = "cpufeatures" version = "0.2.17" @@ -893,6 +995,15 @@ dependencies = [ "syn 3.0.3", ] +[[package]] +name = "document-features" +version = "0.2.12" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d4b8a88685455ed29a21542a33abd9cb6510b6b129abadabdcef0f4c55bc8f61" +dependencies = [ + "litrs", +] + [[package]] name = "dotenvy" version = "0.15.7" @@ -1055,6 +1166,33 @@ version = "0.2.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "77ce24cb58228fbb8aa041425bb1050850ac19177686ea6e0f41a70416f56fdb" +[[package]] +name = "foreign-types" +version = "0.5.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d737d9aa519fb7b749cbc3b962edcf310a8dd1f4b67c91c4f83975dbdd17d965" +dependencies = [ + "foreign-types-macros", + "foreign-types-shared", +] + +[[package]] +name = "foreign-types-macros" +version = "0.2.4" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ea5190182e6915eb873ddbc16e23b711b6eb1f9c00a0d0a3a91b5f6228475225" +dependencies = [ + "proc-macro2", + "quote", + "syn 3.0.3", +] + +[[package]] +name = "foreign-types-shared" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "aa9a19cbb55df58761df49b23516a86d432839add4af60fc256da840f66ed35b" + [[package]] name = "form_urlencoded" version = "1.2.2" @@ -1174,6 +1312,89 @@ dependencies = [ "r-efi", ] +[[package]] +name = "gl_generator" +version = "0.14.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1a95dfc23a2b4a9a2f5ab41d194f8bfda3cabec42af4e39f08c339eb2a0c124d" +dependencies = [ + "khronos_api", + "log", + "xml-rs", +] + +[[package]] +name = "glow" +version = "0.14.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d51fa363f025f5c111e03f13eda21162faeacb6911fe8caa0c0349f9cf0c4483" +dependencies = [ + "js-sys", + "slotmap", + "wasm-bindgen", + "web-sys", +] + +[[package]] +name = "glutin_wgl_sys" +version = "0.6.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2c4ee00b289aba7a9e5306d57c2d05499b2e5dc427f84ac708bd2c090212cf3e" +dependencies = [ + "gl_generator", +] + +[[package]] +name = "gpu-alloc" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "45cf04b2726f02df5508c6de726acdc90cdf97ac771a9a0ffd8ba10a6e696bf9" +dependencies = [ + "bitflags 2.13.1", + "gpu-alloc-types", +] + +[[package]] +name = "gpu-alloc-types" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b2bbed164dd10ed526c2e4fe3e721ca4a71c61730e5aafac6844b417b3227058" +dependencies = [ + "bitflags 2.13.1", +] + +[[package]] +name = "gpu-allocator" +version = "0.27.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c151a2a5ef800297b4e79efa4f4bec035c5f51d5ae587287c9b952bdf734cacd" +dependencies = [ + "log", + "presser", + "thiserror 1.0.69", + "windows", +] + +[[package]] +name = "gpu-descriptor" +version = "0.3.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b89c83349105e3732062a895becfc71a8f921bb71ecbbdd8ff99263e3b53a0ca" +dependencies = [ + "bitflags 2.13.1", + "gpu-descriptor-types", + "hashbrown 0.15.5", +] + +[[package]] +name = "gpu-descriptor-types" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "fdf242682df893b86f33a73828fb09ca4b2d3bb6cc95249707fc684d27484b91" +dependencies = [ + "bitflags 2.13.1", +] + [[package]] name = "group" version = "0.13.0" @@ -1272,6 +1493,12 @@ version = "0.4.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "7f24254aa9a54b5c858eaee2f5bccdb46aaf0e486a595ed5fd8f86ba55232a70" +[[package]] +name = "hexf-parse" +version = "0.2.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dfa686283ad6dd069f105e5ab091b04c62850d3e4cf5d67debad1933f55023df" + [[package]] name = "hkdf" version = "0.12.4" @@ -1630,6 +1857,34 @@ version = "1.0.18" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "8f42a60cbdf9a97f5d2305f08a87dc4e09308d1276d28c869c684d7777685682" +[[package]] +name = "jni-sys" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "41a652e1f9b6e0275df1f15b32661cf0d4b78d4d87ddec5e0c3c20f097433258" +dependencies = [ + "jni-sys 0.4.1", +] + +[[package]] +name = "jni-sys" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c6377a88cb3910bee9b0fa88d4f42e1d2da8e79915598f65fb0c7ee14c878af2" +dependencies = [ + "jni-sys-macros", +] + +[[package]] +name = "jni-sys-macros" +version = "0.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "38c0b942f458fe50cdac086d2f946512305e5631e720728f2a61aabcd47a6264" +dependencies = [ + "quote", + "syn 2.0.119", +] + [[package]] name = "jobserver" version = "0.1.35" @@ -1651,6 +1906,23 @@ dependencies = [ "wasm-bindgen", ] +[[package]] +name = "khronos-egl" +version = "6.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6aae1df220ece3c0ada96b8153459b67eebe9ae9212258bb0134ae60416fdf76" +dependencies = [ + "libc", + "libloading", + "pkg-config", +] + +[[package]] +name = "khronos_api" +version = "3.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e2db585e1d738fc771bf08a151420d3ed193d9d895a36df7f6f8a9456b911ddc" + [[package]] name = "lazy_static" version = "1.5.0" @@ -1666,6 +1938,16 @@ version = "0.2.189" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "3eaf3ede3fee6db1a4c2ee091bf8a8b4dccdc6d17f656fb07896ee72867612f2" +[[package]] +name = "libloading" +version = "0.8.9" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d7c4b02199fee7c5d21a5ae7d8cfa79a6ef5bb2fc834d6e9058e89c825efdc55" +dependencies = [ + "cfg-if", + "windows-link", +] + [[package]] name = "libm" version = "0.2.16" @@ -1678,7 +1960,7 @@ version = "0.1.19" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "2026a5056764a10b2bf5d56488cba40da507f5493a6a429340e2004d9ed085fa" dependencies = [ - "bitflags", + "bitflags 2.13.1", "libc", "plain", "redox_syscall 0.9.1", @@ -1700,6 +1982,12 @@ version = "0.8.2" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "92daf443525c4cce67b150400bc2316076100ce0b3686209eb8cf3c31612e6f0" +[[package]] +name = "litrs" +version = "1.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "11d3d7f243d5c5a8b9bb5d6dd2b1602c0cb0b9db1621bafc7ed66e35ff9fe092" + [[package]] name = "lock_api" version = "0.4.14" @@ -1724,6 +2012,15 @@ dependencies = [ "hashbrown 0.16.1", ] +[[package]] +name = "malloc_buf" +version = "0.0.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "62bb907fe88d54d8d9ce32a3cceab4218ed2f6b7d35617cafe9adf84e43919cb" +dependencies = [ + "libc", +] + [[package]] name = "mcmapper-worker" version = "0.1.0" @@ -1733,12 +2030,14 @@ dependencies = [ "aws-credential-types", "aws-sdk-s3", "base64", + "bytemuck", "image", "rayon", "redis", "sqlx", "tokio", "uuid", + "wgpu", ] [[package]] @@ -1767,6 +2066,21 @@ version = "2.8.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98" +[[package]] +name = "metal" +version = "0.29.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7ecfd3296f8c56b7c1f6fbac3c71cefa9d78ce009850c45000015f206dc7fa21" +dependencies = [ + "bitflags 2.13.1", + "block", + "core-graphics-types", + "foreign-types", + "log", + "objc", + "paste", +] + [[package]] name = "miniz_oxide" version = "0.8.9" @@ -1798,6 +2112,36 @@ dependencies = [ "pxfm", ] +[[package]] +name = "naga" +version = "23.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "364f94bc34f61332abebe8cad6f6cd82a5b65cff22c828d05d0968911462ca4f" +dependencies = [ + "arrayvec", + "bit-set", + "bitflags 2.13.1", + "cfg_aliases", + "codespan-reporting", + "hexf-parse", + "indexmap", + "log", + "rustc-hash", + "spirv", + "termcolor", + "thiserror 1.0.69", + "unicode-xid", +] + +[[package]] +name = "ndk-sys" +version = "0.5.0+25.2.9519653" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "8c196769dd60fd4f363e11d948139556a344e79d451aeb2fa2fd040738ef7691" +dependencies = [ + "jni-sys 0.3.1", +] + [[package]] name = "num-bigint" version = "0.4.8" @@ -1859,6 +2203,15 @@ dependencies = [ "libm", ] +[[package]] +name = "objc" +version = "0.2.7" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "915b1b472bc21c53464d6c8461c9d3af805ba1ef837e1cac254428f4a77177b1" +dependencies = [ + "malloc_buf", +] + [[package]] name = "once_cell" version = "1.21.4" @@ -1918,6 +2271,12 @@ dependencies = [ "windows-link", ] +[[package]] +name = "paste" +version = "1.0.15" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a" + [[package]] name = "pem-rfc7468" version = "0.7.0" @@ -1984,7 +2343,7 @@ version = "0.18.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "60769b8b31b2a9f263dae2776c37b1b28ae246943cf719eb6946a1db05128a61" dependencies = [ - "bitflags", + "bitflags 2.13.1", "crc32fast", "fdeflate", "flate2", @@ -2015,6 +2374,12 @@ dependencies = [ "zerocopy", ] +[[package]] +name = "presser" +version = "0.3.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e8cf8e6a8aa66ce33f63993ffc4ea4271eb5b0530a9002db8455ea6050c77bfa" + [[package]] name = "primeorder" version = "0.13.6" @@ -2033,6 +2398,12 @@ dependencies = [ "unicode-ident", ] +[[package]] +name = "profiling" +version = "1.0.18" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "3d595e54a326bc53c1c197b32d295e14b169e3cfeaa8dc82b529f947fba6bcf5" + [[package]] name = "pxfm" version = "0.1.30" @@ -2084,6 +2455,18 @@ dependencies = [ "getrandom 0.2.17", ] +[[package]] +name = "range-alloc" +version = "0.1.5" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ca45419789ae5a7899559e9512e58ca889e41f04f1f2445e9f4b290ceccd1d08" + +[[package]] +name = "raw-window-handle" +version = "0.6.2" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "20675572f6f24e9e76ef639bc5552774ed45f1c30e2951e1e99c59888861c539" + [[package]] name = "rayon" version = "1.12.0" @@ -2134,7 +2517,7 @@ version = "0.5.18" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ed2bf2547551a7053d6fdfafda3f938979645c44812fbfcda098faae3f1a362d" dependencies = [ - "bitflags", + "bitflags 2.13.1", ] [[package]] @@ -2143,7 +2526,7 @@ version = "0.9.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "07507be7b4a5f9f26eeb41eeaebb1f5a7ff29dfb29739facc21d35bf8b11c21e" dependencies = [ - "bitflags", + "bitflags 2.13.1", ] [[package]] @@ -2152,6 +2535,12 @@ version = "0.1.9" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "cab834c73d247e67f4fae452806d17d3c7501756d98c8808d7c9c7aa7d18f973" +[[package]] +name = "renderdoc-sys" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "19b30a45b0cd0bcca8037f3d0dc3421eaf95327a17cad11964fb8179b4fc4832" + [[package]] name = "rfc6979" version = "0.4.0" @@ -2196,6 +2585,12 @@ dependencies = [ "zeroize", ] +[[package]] +name = "rustc-hash" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "08d43f7aa6b08d49f382cde6a7982047c3426db949b1424bc4b7ec9ae12c6ce2" + [[package]] name = "rustc_version" version = "0.4.1" @@ -2332,8 +2727,8 @@ version = "3.7.0" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "b7f4bc775c73d9a02cde8bf7b2ec4c9d12743edf609006c7facc23998404cd1d" dependencies = [ - "bitflags", - "core-foundation", + "bitflags 2.13.1", + "core-foundation 0.10.1", "core-foundation-sys", "libc", "security-framework-sys", @@ -2498,6 +2893,15 @@ version = "0.4.12" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "0c790de23124f9ab44544d7ac05d60440adc586479ce501c1d6d7da3cd8c9cf5" +[[package]] +name = "slotmap" +version = "1.1.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "bdd58c3c93c3d278ca835519292445cb4b0d4dc59ccfdf7ceadaab3f8aeb4038" +dependencies = [ + "version_check", +] + [[package]] name = "smallvec" version = "1.15.2" @@ -2542,6 +2946,15 @@ version = "0.10.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "023a211cb3138dbc438680b32560ad89f699977624c9f8dbb95a47d5b4c07dd3" +[[package]] +name = "spirv" +version = "0.3.0+sdk-1.3.268.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "eda41003dc44290527a59b13432d4a0379379fa074b70174882adfbdfd917844" +dependencies = [ + "bitflags 2.13.1", +] + [[package]] name = "spki" version = "0.7.3" @@ -2593,7 +3006,7 @@ dependencies = [ "serde_json", "sha2 0.10.9", "smallvec", - "thiserror", + "thiserror 2.0.20", "tokio", "tokio-stream", "tracing", @@ -2648,7 +3061,7 @@ checksum = "aa003f0038df784eb8fecbbac13affe3da23b45194bd57dba231c8f48199c526" dependencies = [ "atoi", "base64", - "bitflags", + "bitflags 2.13.1", "byteorder", "bytes", "crc", @@ -2677,7 +3090,7 @@ dependencies = [ "smallvec", "sqlx-core", "stringprep", - "thiserror", + "thiserror 2.0.20", "tracing", "uuid", "whoami", @@ -2691,7 +3104,7 @@ checksum = "db58fcd5a53cf07c184b154801ff91347e4c30d17a3562a635ff028ad5deda46" dependencies = [ "atoi", "base64", - "bitflags", + "bitflags 2.13.1", "byteorder", "crc", "dotenvy", @@ -2715,7 +3128,7 @@ dependencies = [ "smallvec", "sqlx-core", "stringprep", - "thiserror", + "thiserror 2.0.20", "tracing", "uuid", "whoami", @@ -2740,7 +3153,7 @@ dependencies = [ "serde", "serde_urlencoded", "sqlx-core", - "thiserror", + "thiserror 2.0.20", "tracing", "url", "uuid", @@ -2752,6 +3165,12 @@ version = "1.2.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "6ce2be8dc25455e1f91df71bfa12ad37d7af1092ae736f3a6cd0e37bc7810596" +[[package]] +name = "static_assertions" +version = "1.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "a2eb9349b6444b326872e140eb1cf5e7c522154d69e7a0ffb0fb81c06b37543f" + [[package]] name = "stringprep" version = "0.1.5" @@ -2802,13 +3221,42 @@ dependencies = [ "syn 2.0.119", ] +[[package]] +name = "termcolor" +version = "1.4.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "06794f8f6c5c898b3275aebefa6b8a1cb24cd2c6c79397ab15774837a0bc5755" +dependencies = [ + "winapi-util", +] + +[[package]] +name = "thiserror" +version = "1.0.69" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "b6aaf5339b578ea85b50e080feb250a3e8ae8cfcdff9a461c9ec2904bc923f52" +dependencies = [ + "thiserror-impl 1.0.69", +] + [[package]] name = "thiserror" version = "2.0.20" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "ec86235f5fcc2a73650310756d2ac5b138a5780bbbdfae3eeccec992c435ba4f" dependencies = [ - "thiserror-impl", + "thiserror-impl 2.0.20", +] + +[[package]] +name = "thiserror-impl" +version = "1.0.69" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4fee6c4efc90059e10f81e6d42c60a18f76588c3d74cb83a0b242a2b6c7504c1" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.119", ] [[package]] @@ -3043,6 +3491,18 @@ version = "0.1.4" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "7df058c713841ad818f1dc5d3fd88063241cc61f49f5fbea4b951e8cf5a8d71d" +[[package]] +name = "unicode-width" +version = "0.1.14" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "7dd6e30e90baa6f72411720665d41d89b9a3d039dc45b8faea1ddd07f617f6af" + +[[package]] +name = "unicode-xid" +version = "0.2.6" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "ebc1c04c71510c7f702b52b7c350734c9ff1295c464a03335b00bb84fc54f853" + [[package]] name = "untrusted" version = "0.9.0" @@ -3136,6 +3596,16 @@ dependencies = [ "wasm-bindgen-shared", ] +[[package]] +name = "wasm-bindgen-futures" +version = "0.4.77" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6b7777d5cc23d0e91404e53ce2d5e8ec7acae3026b16233dba62cd3246457950" +dependencies = [ + "js-sys", + "wasm-bindgen", +] + [[package]] name = "wasm-bindgen-macro" version = "0.2.127" @@ -3168,6 +3638,16 @@ dependencies = [ "unicode-ident", ] +[[package]] +name = "web-sys" +version = "0.3.104" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c435338968042f4f59a557f690a253676d47ce13ceb55d70100e7facf6620a30" +dependencies = [ + "js-sys", + "wasm-bindgen", +] + [[package]] name = "webpki-roots" version = "0.26.11" @@ -3186,6 +3666,112 @@ dependencies = [ "rustls-pki-types", ] +[[package]] +name = "wgpu" +version = "23.0.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "80f70000db37c469ea9d67defdc13024ddf9a5f1b89cb2941b812ad7cde1735a" +dependencies = [ + "arrayvec", + "cfg_aliases", + "document-features", + "js-sys", + "log", + "naga", + "parking_lot", + "profiling", + "raw-window-handle", + "smallvec", + "static_assertions", + "wasm-bindgen", + "wasm-bindgen-futures", + "web-sys", + "wgpu-core", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-core" +version = "23.0.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "d63c3c478de8e7e01786479919c8769f62a22eec16788d8c2ac77ce2c132778a" +dependencies = [ + "arrayvec", + "bit-vec", + "bitflags 2.13.1", + "cfg_aliases", + "document-features", + "indexmap", + "log", + "naga", + "once_cell", + "parking_lot", + "profiling", + "raw-window-handle", + "rustc-hash", + "smallvec", + "thiserror 1.0.69", + "wgpu-hal", + "wgpu-types", +] + +[[package]] +name = "wgpu-hal" +version = "23.0.1" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "89364b8a0b211adc7b16aeaf1bd5ad4a919c1154b44c9ce27838213ba05fd821" +dependencies = [ + "android_system_properties", + "arrayvec", + "ash", + "bit-set", + "bitflags 2.13.1", + "block", + "bytemuck", + "cfg_aliases", + "core-graphics-types", + "glow", + "glutin_wgl_sys", + "gpu-alloc", + "gpu-allocator", + "gpu-descriptor", + "js-sys", + "khronos-egl", + "libc", + "libloading", + "log", + "metal", + "naga", + "ndk-sys", + "objc", + "once_cell", + "parking_lot", + "profiling", + "range-alloc", + "raw-window-handle", + "renderdoc-sys", + "rustc-hash", + "smallvec", + "thiserror 1.0.69", + "wasm-bindgen", + "web-sys", + "wgpu-types", + "windows", + "windows-core", +] + +[[package]] +name = "wgpu-types" +version = "23.0.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "610f6ff27778148c31093f3b03abc4840f9636d58d597ca2f5977433acfe0068" +dependencies = [ + "bitflags 2.13.1", + "js-sys", + "web-sys", +] + [[package]] name = "whoami" version = "1.6.1" @@ -3196,12 +3782,85 @@ dependencies = [ "wasite", ] +[[package]] +name = "winapi-util" +version = "0.1.11" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "c2a7b1c03c876122aa43f3020e6c3c3ee5c05081c9a00739faf7503aeba10d22" +dependencies = [ + "windows-sys 0.61.2", +] + +[[package]] +name = "windows" +version = "0.58.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "dd04d41d93c4992d421894c18c8b43496aa748dd4c081bac0dc93eb0489272b6" +dependencies = [ + "windows-core", + "windows-targets 0.52.6", +] + +[[package]] +name = "windows-core" +version = "0.58.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "6ba6d44ec8c2591c134257ce647b7ea6b20335bf6379a27dac5f1641fcf59f99" +dependencies = [ + "windows-implement", + "windows-interface", + "windows-result", + "windows-strings", + "windows-targets 0.52.6", +] + +[[package]] +name = "windows-implement" +version = "0.58.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "2bbd5b46c938e506ecbce286b6628a02171d56153ba733b6c741fc627ec9579b" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.119", +] + +[[package]] +name = "windows-interface" +version = "0.58.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "053c4c462dc91d3b1504c6fe5a726dd15e216ba718e84a0e46a88fbe5ded3515" +dependencies = [ + "proc-macro2", + "quote", + "syn 2.0.119", +] + [[package]] name = "windows-link" version = "0.2.1" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "f0805222e57f7521d6a62e36fa9163bc891acd422f971defe97d64e70d0a4fe5" +[[package]] +name = "windows-result" +version = "0.2.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "1d1043d8214f791817bab27572aaa8af63732e11bf84aa21a45a78d6c317ae0e" +dependencies = [ + "windows-targets 0.52.6", +] + +[[package]] +name = "windows-strings" +version = "0.1.0" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "4cd9b125c486025df0eabcb585e62173c6c9eddcec5d117d3b6e8c30e2ee4d10" +dependencies = [ + "windows-result", + "windows-targets 0.52.6", +] + [[package]] name = "windows-sys" version = "0.48.0" @@ -3356,6 +4015,12 @@ version = "0.6.3" source = "registry+https://github.com/rust-lang/crates.io-index" checksum = "1ffae5123b2d3fc086436f8834ae3ab053a283cfac8fe0a0b8eaae044768a4c4" +[[package]] +name = "xml-rs" +version = "0.8.29" +source = "registry+https://github.com/rust-lang/crates.io-index" +checksum = "e450f9b2ed1dff33c94c12589a87338689467b9c4f5d8a5710bd09a847d2c8a7" + [[package]] name = "xmlparser" version = "0.13.6" diff --git a/worker/Cargo.toml b/worker/Cargo.toml index 942263f..a3500b2 100644 --- a/worker/Cargo.toml +++ b/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 diff --git a/worker/Dockerfile b/worker/Dockerfile index eea1a55..d3e5068 100644 --- a/worker/Dockerfile +++ b/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 diff --git a/worker/examples/benchmark.rs b/worker/examples/benchmark.rs new file mode 100644 index 0000000..70e9ebd --- /dev/null +++ b/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 { + 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(§ions[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, §ions, batch); + + match GpuRenderBackend::try_new().await { + Some(gpu) => { + println!("GPU adapter: {}", gpu.adapter_name()); + run_backend("gpu", &gpu, &columns, §ions, 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, §ions, batch); + } + None => println!("hybrid (skipped — no compatible GPU adapter found)"), + } +} diff --git a/worker/src/lib.rs b/worker/src/lib.rs new file mode 100644 index 0000000..72da688 --- /dev/null +++ b/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; diff --git a/worker/src/main.rs b/worker/src/main.rs index 3091fda..6157bcc 100644 --- a/worker/src/main.rs +++ b/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 = Box::new(CpuRenderBackend::new()); + let backend: Box = 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; } diff --git a/worker/src/mesh.rs b/worker/src/mesh.rs index d33f35a..237c82a 100644 --- a/worker/src/mesh.rs +++ b/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()); diff --git a/worker/src/render/cpu.rs b/worker/src/render/cpu.rs index b1f01d2..5606a76 100644 --- a/worker/src/render/cpu.rs +++ b/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> { - 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); + } } diff --git a/worker/src/render/gpu.rs b/worker/src/render/gpu.rs new file mode 100644 index 0000000..56c41a6 --- /dev/null +++ b/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 base_colors: array; +@group(0) @binding(1) var out_pixels: array; + +// 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) { + 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 blocks: array; +@group(0) @binding(1) var out_masks: array; + +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) { + 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(-1, 1, 0, 0, 0, 0); + var offsets_y = array(0, 0, -1, 1, 0, 0); + var offsets_z = array(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 { + 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 { + 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::()) 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 = 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> { + let base = base_colors(columns); + let packed_input: Vec = 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 = 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 { + 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()); + } +} diff --git a/worker/src/render/hybrid.rs b/worker/src/render/hybrid.rs new file mode 100644 index 0000000..c44bc20 --- /dev/null +++ b/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> { + 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)); + } +} diff --git a/worker/src/render/mod.rs b/worker/src/render/mod.rs index 5d9a06a..9a61114 100644 --- a/worker/src/render/mod.rs +++ b/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>; + + 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> { + 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> { + 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) }