From 3a424afe8c717feca3b91cf3825bce2605f6adda Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?J=C3=A9r=C3=B4me=20Bousqui=C3=A9?= Date: Tue, 22 Sep 2026 15:48:15 +0200 Subject: [PATCH] GPU culling --- AGENTS.md | 1 + README.md | 13 +- docs/DRAFT.md | 420 ++++++++++++++++++- docs/PLAN.md | 2 +- docs/ROADMAP.md | 22 +- docs/tech/ARCHI_APP.md | 32 +- docs/tech/ARCHI_CPU_GPU.md | 41 +- docs/tech/FRAME_LOOP.md | 11 +- docs/user/README.md | 1 + docs/user/gpu-driven.md | 132 ++++++ lib/examples/demo.rs | 31 +- lib/src/README.md | 11 +- lib/src/app.rs | 20 + lib/src/core/renderer.rs | 621 ++++++++++++++++++++++++----- lib/src/math/frustum.rs | 178 +++++++++ lib/src/math/geometry.rs | 78 ++++ lib/src/math/mod.rs | 4 +- lib/src/pipeline/pipeline_cache.rs | 16 +- lib/src/resources/mod.rs | 6 +- lib/src/resources/uniform.rs | 236 +++++++++++ lib/src/scene/mod.rs | 2 +- lib/src/scene/scene.rs | 209 +++++++++- lib/src/shaders/gpu_driven.wgsl | 188 +++++++++ lib/src/utils/conf.rs | 26 ++ lib/tests/wgsl_validate.rs | 31 ++ 25 files changed, 2155 insertions(+), 177 deletions(-) create mode 100644 docs/user/gpu-driven.md create mode 100644 lib/src/math/frustum.rs create mode 100644 lib/src/shaders/gpu_driven.wgsl diff --git a/AGENTS.md b/AGENTS.md index 105f95b..ed62054 100644 --- a/AGENTS.md +++ b/AGENTS.md @@ -43,6 +43,7 @@ WGPU doesn't have a native "Context" object — this type groups them together f - wgpu 30.0.0 is pinned in `lib/Cargo.toml`. The comment says "check the latest version" — verify compatibility before upgrading. - No feature flags, no dev-dependencies, no tests yet. Adding any requires updating both `Cargo.toml` files if the dependency spans crates. - The workspace has no `[workspace.dependencies]` section. Dependencies are declared per-crate rather than centrally. +- **WGSL `select` argument order** (cost us a day): `select(reject, accept, cond)` returns the **second** arg when `cond` is true — the reverse of HLSL's `select(trueVal, falseVal, cond)`. In `shaders/gpu_driven.wgsl` the cull pass must stay `select(0u, u32(flags.z), visible)` (visible ⇒ full count, culled ⇒ 0). Swapped args silently zero the counts of every visible entity → black window. See the GOTCHA comment at the top of that shader. ## lean-ctx diff --git a/README.md b/README.md index 29fc41b..df25c73 100644 --- a/README.md +++ b/README.md @@ -2,7 +2,7 @@ WSG is a Rust library that wraps [wgpu](https://github.com/gfx-rs/wgpu) and [winit](https://crates.io/crates/winit) for simple GPU drawing. It groups the five core wgpu objects (Instance, Surface, Adapter, Device, Queue) behind a single `Context`, adds small building blocks (`Mesh`, `Material`, `PipelineCache`, `Frame`), and exposes the low-level primitives for advanced users. -> **Status: unstable development version.** The **declarative workflow** (`AppBuilder` + `App` + `AppHandler`) is the **recommended** path and is fully working: scene auto-rendering (`App::render_scene`), 3D Phong lighting, textures, shadows, camera and unified input — the `demo` example is the showcase. The **manual workflow** (`Context`/`Renderer`/`PipelineCache`) coexists for fine-grained control. Meshes are declared from a CPU `Geometry` (retained as `Arc` on the Mesh). The GPU-driven two-pass pipeline described in the architecture docs is **not implemented yet** — see [Status](#status) and [Roadmap](#roadmap). +> **Status: unstable development version.** The **declarative workflow** (`AppBuilder` + `App` + `AppHandler`) is the **recommended** path and is fully working: scene auto-rendering (`App::render_scene`), 3D Phong lighting, textures, shadows, camera and unified input — the `demo` example is the showcase. The **manual workflow** (`Context`/`Renderer`/`PipelineCache`) coexists for fine-grained control. Meshes are declared from a CPU `Geometry` (retained as `Arc` on the Mesh). The **GPU-driven two-pass pipeline** (Compute Pass deriving world matrices + frustum culling → indirect draws) is **implemented** (Step 15, Phase 3): `render_scene` and the shadow pass are 100 % indirect, and frustum culling is opt-in (`AppBuilder::with_culling(true)`, off by default) — see [Status](#status), [docs/user/gpu-driven.md](docs/user/gpu-driven.md) and [Roadmap](#roadmap). ## Status @@ -11,7 +11,7 @@ WSG is a Rust library that wraps [wgpu](https://github.com/gfx-rs/wgpu) and [win | Manual workflow (`Context` + `Renderer` + `PipelineCache`) | ✅ Working (advanced — fine-grained control) | | `App` / `AppBuilder` / `AppHandler` event-loop facade | ✅ Working — window, events, frame presentation, and **automatic scene rendering** (the per-frame view is exposed via `Frame::view()`) | | `Scene` resource/entity registry | ✅ Working — the engine renders every registered entity automatically in one batched render pass (`App::render_scene`) | -| GPU-driven two-pass pipeline (Compute → indirect draw) | 📋 Roadmap — spec in [docs/tech/ARCHI_CPU_GPU.md](docs/tech/ARCHI_CPU_GPU.md) | +| GPU-driven two-pass pipeline (Compute → indirect draw) | ✅ Working (Step 15, Phase 3) — `render_scene` + shadow pass are 100 % indirect; opt-in frustum culling (bug « fenêtre noire » fixed 2026-09-22 — WGSL `select` argument order — and verified by GPU readback; see DRAFT D14). User doc: [gpu-driven.md](docs/user/gpu-driven.md) · spec: [ARCHI_CPU_GPU.md](docs/tech/ARCHI_CPU_GPU.md) | | 3D infrastructure (uniform bind groups, MVP + camera in the pipeline) | ✅ Working — the `Renderer` uploads per-frame camera matrices (active `Camera`) and per-entity world matrices to shared uniform buffers every frame; the **MVP is reached** (Step 5) : the `cube` example renders a rotating Phong-lit cube via the `standard` shader | Note: `standard_shader.wgsl` (Phong, with an explicit **unlit** mode) is the **single** shader the library ships — flat 2D drawing is its unlit variant (`Renderer::set_unlit(true)` or `app.renderer_mut().set_unlit(true)`). See the `cube` example (3D, lit) and the `simple` example (2D, unlit). @@ -145,7 +145,7 @@ fn main() { - **Executor layer (`Renderer`)** — binds a `Material` pipeline + `Mesh` buffers into a RenderPass and submits the commands. Rendering a whole `Scene` (`render_scene`) batches all entities into **one encoder + one submit per frame**; the low-level `render` still allocates one per object. - **Supporting pieces** — `PipelineCache` (shader → compiled RenderPipeline, `Arc`-shared), `Material`, `Geometry`/`Mesh`/`Vertex`, `Scene` (string-ID registry), `Camera`/`Transform` (active camera wired to the frame uniforms, Step 4.3). `Geometry` is the CPU source of truth (positions/normals/UVs/colors), `Mesh` uploads it to GPU buffers and retains the `Arc`, `Vertex` is the interleaved upload contract (Step 8). -The planned target architecture — a GPU-driven two-pass pipeline (Compute Pass: world matrices + frustum culling → Indirect Draw Buffer, then a single `draw_indexed_indirect` per frame) — is specified in [docs/tech/ARCHI_APP.md](docs/tech/ARCHI_APP.md) and [docs/tech/ARCHI_CPU_GPU.md](docs/tech/ARCHI_CPU_GPU.md) but is **not implemented yet**. +The **GPU-driven two-pass pipeline** (Step 15, Phase 3) is implemented: a Compute Pass derives each entity's world matrix and fills per-entity indirect draw arguments (with opt-in frustum culling), then the main and shadow render passes issue one indirect draw per active slot. Specified in [docs/tech/ARCHI_APP.md](docs/tech/ARCHI_APP.md) and [docs/tech/ARCHI_CPU_GPU.md](docs/tech/ARCHI_CPU_GPU.md); user-facing guide in [docs/user/gpu-driven.md](docs/user/gpu-driven.md). ## Quick reference @@ -201,8 +201,8 @@ Three layers (user docs and API reference in **English**; technical docs in **Fr - [Shadows](docs/user/shadows.md) · [Camera & input](docs/user/camera-input.md) · [Examples](docs/user/examples.md) **Technical documentation — `docs/tech/`** (internal architecture; each document states whether it describes the **current** or the **target** architecture): -- [ARCHI_APP](docs/tech/ARCHI_APP.md) — engine architecture. ✅ Current facade (`App`/`AppHandler`) / 🎯 **Target** — the GPU-driven two-pass pipeline parts are not implemented yet. -- [ARCHI_CPU_GPU](docs/tech/ARCHI_CPU_GPU.md) — CPU/GPU workload split specification. 🎯 **Target** — GPU-driven pipeline, ROADMAP Phase 3. +- [ARCHI_APP](docs/tech/ARCHI_APP.md) — engine architecture. ✅ **Current** — facade (`App`/`AppHandler`) and GPU-driven two-pass pipeline (implemented in Phase 3, 2026-07-20, with the documented deviations); only the future double-buffering notes remain target. +- [ARCHI_CPU_GPU](docs/tech/ARCHI_CPU_GPU.md) — CPU/GPU workload split specification. ✅ **Current** — implemented in ROADMAP Phase 3 (2026-07-20, DRAFT Étape 17, decisions D1–D14); deviations from the original spec are noted in the document. - [ARCHI_RENDU](docs/tech/ARCHI_RENDU.md) — update/render mutability model. ✅ Current dichotomy (auto scene render) / 🎯 **Target** — material batching. - [ARCHI_ARENES](docs/tech/ARCHI_ARENES.md) — 🎯 **Target/deferred** — slotmap generational handles; String IDs are used today. - [FRAME_LOOP](docs/tech/FRAME_LOOP.md) — frame lifetime and resource persistence. ✅ **Current** — implemented. @@ -212,7 +212,7 @@ Three layers (user docs and API reference in **English**; technical docs in **Fr ## Roadmap 1. ✅ **Scene auto-rendering** — `App::render_scene` iterates registered entities and draws them in one encoder/submit per frame; the frame view is exposed to `AppHandler::render` for custom draws. (Done 2026-09-16.) -2. **GPU-driven two-pass pipeline** — Compute Pass (world matrices + frustum culling) filling an indirect draw buffer, single `draw_indexed_indirect` (see ARCHI_CPU_GPU). +2. ✅ **GPU-driven two-pass pipeline** — Compute Pass (world matrices + frustum culling) filling an indirect draw buffer, then indirect draws (see ARCHI_CPU_GPU). *(Done 2026-07-20 — see item 17. Deviation from the original spec: one indirect draw **per slot** rather than a single fused draw, DRAFT D1.)* 3. **CPU→GPU transform sync** — persistent transform buffers with ring (triple) buffering. 4. ✅ **Real 3D pipeline (MVP reached)** — MVP uniforms + camera support in the vertex shader. *(Engine plumbing done 2026-09-16; Step 5, 2026-09-17: `standard` wired into the `cube` example — a unit cube lit (Phong) and spinning, rendered automatically by `App::render_scene`. Removal of `basic`: flat 2D = unlit variant of `standard` via `Renderer::set_unlit`.)* 5. **Typed resource handles** — keep String IDs for the MVP (current design, source of truth in `Scene`); slotmap-based generational handles (`ARCHI_ARENES.md`) are deferred to a later performance pass. @@ -227,3 +227,4 @@ Three layers (user docs and API reference in **English**; technical docs in **Fr 14. ✅ **Unified input (Step 15.B)** — `core::input::InputState` gives cross-frame **pressed/held/released** semantics for keyboard (physical `KeyCode`) and mouse (buttons, position, per-frame delta, wheel scroll), rotated by `begin_frame`/`end_frame` around `AppHandler::update`. `App` exposes it as a public `input` field, fed from winit `WindowEvent`s and reset each frame. Gamepad is reserved/deferred (DRAFT D7). (Done 2026-09-20; 5 unit tests; winit event handling is host-driven on the CPU, not WGSL.) 15. ✅ **Orbital camera + final demo (Step 15.C)** — `resources::CameraController` (yaw/pitch/distance/target, `apply_to` writes into a `Camera`, drag-orbit + wheel-zoom + clamps) drives the new `demo` example: one of each primitive, procedural textures, standard Phong material, a shadow-casting directional light + point + spot, and live mouse-orbit / wheel-zoom / `R` reset / `1`/`2`/`3` view presets. Run with `cargo run -p wsg-lib --example demo`. (Done 2026-09-20; runtime-verified headless.) 16. ✅ **User documentation (Step 16, Phase 5)** — `docs/user/` (quickstart, meshes, materials, lights, shadows, camera & input, examples) written in English and cross-linked to each other, to the tech docs and to rustdoc; tech docs interlinked with their stale status banners refreshed; this README re-anchored (declarative workflow = recommended, manual = advanced, `demo` = showcase, pollster 1.x). (Done 2026-07-19.) +17. ✅ **GPU-driven rendering (Step 15, Phase 3.1/3.2/3.3)** — world matrices and indirect draw args move from CPU to GPU. `shaders/gpu_driven.wgsl` (two compute entry points, `compute_matrices` + `cull`, one module, explicit 3-group layout) runs before the render passes over a fixed 256-slot table (the world-matrix buffer is bound to the `uniform` object slot; WebGPU caps a `uniform` binding at 64 KB and a `uniform` offset at 256 B, so each matrix slot is padded to 256 B and 256 × 256 B = 64 KB is the max); `render_scene` and the shadow pass become **100 % indirect** (one indirect draw per active slot, culled/inactive slots are no-ops), and the per-entity CPU draw loop is gone. New `math::Frustum` (Gribb–Hartmann, WebGPU `[0,1]` z) + `BBox` on `Geometry`; `TransformSlot`/`MatSlot`/`BBoxSlot`/`DrawSlot`/`CullUniforms` Pod mirrors of the WGSL structs. Frustum **culling is off by default** (non-regression) and opt-in via `AppBuilder::with_culling(true)` / `Renderer::set_culling(bool)`; the `demo` enables it. The object bind-group layout is now dynamic so every entity shares one GPU matrix buffer via per-slot offsets. (Done 2026-07-20; WGSL + frustum + scene-slot tests, 56 lib / 3 WGSL / 3 doctests all green. **Culling fix 2026-09-22**: the WGSL `select` arguments had been written HLSL-style, silently zeroing the draw count of every *visible* entity — a black window; fixed and verified by GPU readback, see DRAFT D14.) diff --git a/docs/DRAFT.md b/docs/DRAFT.md index f7034a8..db13a0c 100644 --- a/docs/DRAFT.md +++ b/docs/DRAFT.md @@ -1,24 +1,404 @@ -# Prochaine étape +# DRAFT — Étape 17 : Rendu GPU-driven (ROADMAP 3.1 / 3.2 / 3.3) -> Étape 16 (Phase 5 — Documentation & Polish) **terminée** le 2026-07-19. +> **Statut** : **implémenté et validé** (voir la section « Acceptation » plus bas). +> Conventions : cette étape couvre les 3 items de la Phase 3 de `docs/ROADMAP.md`. +> ROADMAP / README / docs user sont mis à jour ; ce DRAFT est conservé comme référence de +> conception (les décisions D1–D14 y sont documentées). > -> **Tuning caméra 2026-07-19** — sur rétroaction utilisateur (caméra trop sensible, orbit permanent) : -> (1) `CameraController` a gagné deux champs configurables `orbit_sensitivity`/`zoom_factor` (defaults : 0.01→0.005 -> rad/px, 0.9/tic) ; (2) `InputState` normalise `PixelDelta`/32 en unités « notch » (le Wayland renvoyait ~100 px/tic, -> donc `0.9^100` → zoom au clamping en un geste) ; (3) le `demo` orbite désormais **sur clic gauche enfoncé** (arc-rotate -> classique). Docs `camera-input.md`/`examples.md` mises à jour. 51 tests OK. +> **Note de validation runtime (correction D12)** : la capacité est passée de **4096 à 256** +> entités, et le slot de matrice mondes est **padded de 64 à 256 o**. Le buffer de matrices est +> lié au slot `uniform` « object » du pipeline de rendu, et WebGPU impose **deux** contraintes : +> (1) une binding `uniform` unique est plafonnée à `max_uniform_buffer_binding_size` (64 ko), et +> (2) un offset de buffer `uniform` doit être un multiple de `min_uniform_buffer_offset_alignment` +> (256 o). Une matrice de 64 o ne peut donc jamais être adressée individuellement par un offset +> dynamique `uniform` : chaque slot de matrice est **padded à 256 o** (`MatSlot { m, pad }`), et +> 256 slots × 256 o = 64 ko est le maximum adressable. Par ailleurs, le bind group « object » lie +> une **slice de 64 o** (une matrice) plutôt que le buffer entier — une binding sur tout le buffer +> plafonnerait l'offset dynamique à 0. Les mentions « 4096 » / « 1024 » ci-dessous renvoient au +> plan initial ; la valeur réelle (et les tailles de buffer dérivées) est 256 (cf. D2 / D4 / D12). > -> **Bug fix 2026-07-19** — `InputState::begin_frame()` remettait à zéro `mouse_delta`/`scroll` AVANT que -> `update()` ne les lise, alors que les événements winit (CursorMoved/MouseWheel) s'accumulent ENTRE deux -> frames. Résultat : la caméra orbitale du `demo` ne bougeait jamais (souris/molette toujours (0,0) dans -> `update`). Corrigé par rotation des accumulateurs (`frame_mouse_delta`/`frame_scroll` → queryables à -> `begin_frame`), aligné sur le pattern clavier/boutons. Tests mis à jour (ordre réel winit) + 45/50 tests OK. -> -> Traduction anglaise de toute la documentation (hors `docs/tech/`, DRAFT/PLAN/ROADMAP) **terminée** le 2026-07-19 : -> `docs/user/*`, `README.md`, READMEs de modules, doc/rustdoc de tous les `.rs` (src + examples + tests), -> `Étape`→`Step` global. Vérifications : 50 tests OK, `cargo fmt` clean, aucun lien cassé, 0 accent restant hors zone franche. +> **Note du 2026-09-22 (D14, correction post-implémentation)** : le `demo` (culling activé) affichait +> une **fenêtre noire** — readback GPU : tous les comptes de draw args étaient à 0 alors que les +> transforms, bboxes et plans de frustum vus par le GPU étaient corrects. Cause racine : l'ordre des +> arguments de `select` en WGSL (`select(reject, accept, cond)` renvoie le **second** argument quand +> `cond` est vrai — l'inverse de la convention HLSL). Corrigé et vérifié par readback (cf. D14). -## Prochaines options -- **Phase 3 — GPU-driven rendering** (ROADMAP 3.1/3.2/3.3) : indirect draw, buffers de paramètres GPU, culling GPU. C'est le gros morceau performance qui reste. -- **Phase 4.4 — Performance** : LOD, instancing/multi-instancing, occlusion culling, batching par matériau (4.3). -- Multi-caméras (`scene.set_active_camera`) — restant de la Phase 2.1. +## Objectifs + +1. **3.1 — Matrices mondes sur le GPU** : un compute shader dérive la matrice monde de + chaque entité à partir de ses données de transform (T·R·S), au lieu d'un `to_matrix()` + CPU par entité chaque frame. +2. **3.2 — Culling GPU** : un compute shader teste la visibilité de chaque entité + (approximation sphère vs les 6 plans du frustum) et écrit un slot d'arguments de draw. +3. **3.3 — Draw indirect** : le rendu de la scène passe en `draw_indexed_indirect` / + `draw_indirect` par slot d'entité, piloté par les arguments produits par le culling. + +## Contraintes & principes + +- **API publique stable** : les exemples existants (cube, simple, shadow_test, spot_test, + manual) restent fonctionnels **sans modification de leurs appels**. En pratique aucun n'est + touché (cf. D11) ; seul `demo.rs` change (activation du culling, cf. D8). +- **Aucune régression visuelle** : le culling est **désactivé par défaut** (D8). Le chemin + GPU-driven (matrices + draw indirect) est actif pour `render_scene` mais produit un rendu + **identique** au chemin CPU actuel. +- **`standard_shader.wgsl` non modifié** (D2) : le binding group 1 reste + `var object: ObjectUniform`. Seuls les commentaires en français sont traduits en + anglais (convention : toute la doc est en anglais). + +## Décisions (à valider) + +### D1 — Un draw indirect par entité (pas de multi-instancing par groupe de mesh) + +Chaque entité possède son **slot GPU** (index stable `0..N`). Un draw indirect par entité lit +les arguments de son slot. + +- ✅ Couvre les 3 items de la roadmap : matrices mondes GPU (3.1), culling GPU (3.2), + draw indirect (3.3). +- ✅ API publique inchangée ; les entités hétérogènes (meshs différents) restent trivialement + supportées. +- ❌ Coût CPU : 1 `draw_indirect` par entité (vs 1 draw multi-instanced par groupe de mesh). + Pour la scale (milliers d'instances d'un même mesh), le multi-instancing par groupe est plus + efficace → **reporté en Phase 4.4** (instancing / multi-instancing). +- **Capacité fixe : 4096 entités** (`MAX_GPU_ENTITIES`). `Scene::add_entity` renvoie `Err` + au-delà. Les slots sont **append-only avec tombstones** : la suppression ne décale pas les + indices (stabilité GPU) ; l'entité est marquée inactive (flag dans le slot de transform) et + le culling met ses arguments à 0. + +### D2 — Le bind group « object » par entité devient une slice du buffer GPU-computé + +- Le binding group layout existant (group 1, `var object: ObjectUniform`, 64 o) est + **conservé tel quel** (shader inchangé), mais rendu **dynamique** (`has_dynamic_offset: true`). +- La valeur de 64 o vient d'une **slice de `WorldMatrixBuffer`** (storage+uniform buffer + GPU-computé, **slot 256 o** padded — cf. D12) plutôt que d'un UBO CPU par entité. +- **Un seul** bind group « object » partagé (`matrix_object_bg`) lie une **slice de 64 o** + (une matrice) du buffer ; chaque draw l'utilise avec un **offset dynamique** `slot × 256` qui + sélectionne le slot (pas un bind group par slot — un seul bind group + offset dynamique, ce + qui évite N bind groups). Une slice de 64 o (et non le buffer entier) est **requise** : une + binding sur tout le buffer plafonnerait l'offset dynamique à 0 (cf. D12). +- Le shadow pass bénéficie automatiquement (même bind group, même offset dynamique, cf. D10). +- Conséquence : les UBO d'objet CPU par entité (`object_cache`) ne sont plus utilisés par + `render_scene` ; ils restent pour le chemin bas niveau `render()` (non-régression, offset 0). + +### D3 — Layout des slots d'indirect draw : 80 o (pdc(16,20)) + +- `DrawSlot` = 80 o = 5 × `vec4`. Les 5 premiers u32 = arguments **indexed** + (`index_count, instance_count, base_vertex, first_instance, instance_offset`) ; les 4 + premiers u32 = arguments **non-indexed** (`vertex_count, instance_count, first_vertex, + first_instance`). +- Raison : l'offset du buffer indirect doit être multiple de 16 o (alignement WGSL du + tableau) **et** multiple de 20 o (contrainte WebGPU pour 5 args) → pdc(16,20) = 80 o. Avec + un stride de 80 o, l'offset du slot i est `80*i` (multiple de 8 et de 4 → valide pour + `draw_indirect` **et** `draw_indexed_indirect`). +- Le buffer est **zéro-initialisé à la création** ; le culling écrit `a` (et `b`) à chaque + frame ; `c..e` restent 0 (jamais lus par le draw). + +### D4 — Un seul buffer par ressource, upload CPU chaque frame (pas de double-buffering) + +- `TransformBuffer` (64 o × 256 = 16 ko, storage, CPU→GPU) : écrit par + `queue.write_buffer` chaque frame (les transforms viennent de `Scene`, modifiés par + l'utilisateur dans `update()`). +- `WorldMatrixBuffer` (256 o × 256 = 64 ko, storage+uniform, GPU-write, **slots padded à 256 o** + — cf. D12) : écrit par le compute « matrices » ; lu par le pipeline de rendu via le slot + `uniform` « object » + (64 ko = la limite `max_uniform_buffer_binding_size`, cf. D12). +- `IndirectArgsBuffer` (80 o × 256 = 20 ko, storage, GPU-write) : zéro-initialisé, écrit + par le compute « culling ». +- `CullUniformsBuffer` (112 o, uniform, CPU→GPU) : uploadé chaque frame (6 plans + + `num_slots` + `culling`). +- Pas de ring-buffer pour v1 (simplicité) ; le `write_buffer` + les compute + le render sont + dans **un seul CommandEncoder** → ordre garanti sur le GPU. Le ring-buffer est une + optimisation possible plus tard (Phase 4.4). + +### D5 — Bounding box : AABB dans `Geometry`, approximaté par une sphère pour le culling v1 + +- `Geometry` calcule son **AABB** (min/max des positions) à la construction ; `Mesh` le + conserve (champ `bbox: BBox`). +- `Scene::create_mesh` / `add_mesh` en dérive un **`BBoxSlot`** (32 o : `min: vec3f` + `max: vec3f`, + coins **locaux** du box) stocké dans `BBoxBuffer` (32 o × 256 = 8 ko, storage, upload + **une fois** par mesh — ré-upload seulement quand l'ensemble des meshes change, via + `gpu_generation`). Le centre et les demi-extents sont **dérivés dans le shader** + (`center = (min+max)/2`, `half_extents = (max-min)/2`) — le buffer ne stocke que les coins. +- **Culling v1 = test sphère vs 6 plans** : le rayon GPU est `length(half_extents) * max(scale)` + (demi-diagonale du box × plus grand facteur d'échelle) et le centre monde est + `translation + rotation * center_local` (pas d'échelle sur le centre — l'échelle est portée par le + rayon). Visible si aucun plan n'a `distance(centre, plan) < -rayon`. + - Pourquoi une sphère et non l'AABB exact : l'AABB exact transformé nécessite 8 sommets + + projections par axe (coût compute plus élevé, complexité WGSL) ; la sphère est + **conservative** (ne culle jamais un objet visible) et suffit pour v1. L'AABB exact + transformé (8 sommets, min/max par axe) est une amélioration possible (Phase 4.4). +- Le **mesh_index** de chaque entité est stocké dans `flags.x` du slot de transform (rempli + par le Renderer chaque frame) ; le culling lit `bboxes[transforms[i].flags.x]`. + +### D6 — Plans du frustum : Gribb-Hartmann côté CPU, upload dans `CullUniforms` + +- Extrait les 6 plans (left, right, bottom, top, near, far) de `view * proj` + (Gribb-Hartmann, adapté au clip depth [0,1] de WebGPU). +- Chaque plan = `vec4` (normale + d), normalisé. +- Uploadé dans `CullUniformsBuffer` (group **2** du compute, binding 0, `var`) chaque frame + (layout final explicite à 3 groupes — cf. écart noté dans les critères d'acceptation : + group 0 = transforms, group 1 = matrices, group 2 = culling). +- Implémenté dans `math/frustum.rs` (nouveau module) avec des tests unitaires (plans de + l'identité, plans d'une perspective standard, orientation des normales). + +### D7 — Un shader compute `gpu_driven.wgsl` avec deux entry points + +- **`compute_matrices`** (`@workgroup_size(64)`) : `m[i] = T·R·S` pour `i in 0..num_slots` + (T=translation, R=quaternion normalisé, S=scale → `mat4x4`). La construction WGSL + reproduit `Transform::to_matrix()` (colonnes de rotation échelonnées par S, translation + dans la 4e colonne). +- **`cull`** (`@workgroup_size(64)`) : pour `i in 0..num_slots` : si inactive → + `draws[i] = 0` ; sinon test sphère vs 6 plans → `draws[i].a = visible ? args : 0`. +- Deux `ComputePipeline` créés depuis le **même module** (même layout de 3 bind groups). +- `num_slots` (nombre de slots **alloués**, pas le nombre d'entités vivantes) dans + `CullUniforms` → le compute parcourt tous les slots alloués (les tombstones sont marqués + inactifs et produisent `draws[i] = 0`). + +### D8 — Culling désactivé par défaut ; activation par `AppBuilder::with_culling(true)` + +- `Renderer::set_culling(bool)` + `culling_enabled()` (pub, documenté : n'affecte que + `render_scene`). +- `AppBuilder::with_culling(self, bool) -> Self` ; appliqué dans `resumed()` après création + du Renderer. +- Par défaut **OFF** (non-régression) : le pass « culling » tourne quand même (marque tout + visible, `culling=0`) → le chemin indirect est toujours actif, mais rien n'est cullé. + Le `demo` active le culling. +- Justification : le culling est un gain de perf, pas un changement de comportement visuel ; + le désactiver par défaut protège contre un bug de culling (objet qui disparaît) qui + casserait les exemples. + +### D9 — Slots d'entité : `Vec>` append-only + index label→slot + +- `Scene::entities` passe de `HashMap` à : + - `entity_slots: Vec>` (append-only, tombstones) + - `entity_labels: Vec>` (parallèle) + - `entity_index: HashMap` (label → slot) + - `entity_generation: u64` (incrémenté à chaque add/remove → le Renderer ré-uploade le + mapping bbox si changé) +- **`Entity` est inchangé** (pas de champ label ; le label reste dans `entity_labels`). +- L'ordre d'itération est **stable** (ordre d'insertion) — important pour la cohérence des + slots (le `HashMap` actuel a un ordre d'itération non-déterministe). +- `add_entity` / `add_entity_with_transform` : trouvent un slot libre (premier `None` ou + append si `len < MAX_GPU_ENTITIES`) ; `Err` si capacité atteinte. `remove_entity` : + tombstone + dé-map de l'index. `set_entity_transform` : mise à jour in-place (pas de bump + de génération — le buffer de transforms est re-uploadé chaque frame de toute façon). +- `iter_entities()` : signature inchangée `(label, &Arc, &Transform)`, saute les + tombstones. +- `entity_count()` : nombre de slots vivants (comportement inchangé pour les tests). + +### D10 — Le shadow pass est GPU-driven aussi (gratuitement) + +- Le pass ombre (casters) utilise les **mêmes** bind groups d'objet (slices de + `WorldMatrixBuffer`) et les **mêmes** slots d'indirect → les casters sont cullés par le + même pass compute. +- Le shadow pass de `Renderer::render_scene` passe en draw indirect (les casters = + sous-ensemble des slots ; le Renderer itère les slots et ne draw que ceux dont le mesh est + caster, avec les args du slot). + +### D11 — Stratégie non-régression : `render_scene` devient GPU-driven, `render` reste CPU + +- `Renderer::render_scene` passe de `&self` à `&mut self` (nécessaire pour + `queue.write_buffer` des buffers transform/cull + mise à jour du cache de bind groups). +- `App::render_scene` (qui prend déjà `&mut self`) utilise `renderer_mut()` → **aucun + exemple n'appelle `Renderer::render_scene` directement** (le default `AppHandler::render` + passe par `App::render_scene`) → **aucune rupture d'API pour les exemples**. +- Le chemin bas niveau `Renderer::render()` (utilisé par `manual.rs`) **garde** les UBO + d'objet CPU → `manual.rs` n'est pas touché. +- Les 45 tests unitaires + 2 tests WGSL + 3 doctests restent verts (les tests Scene ne + touchent pas le GPU ; les tests de frustum sont nouveaux). + +### D12 — Capacité fixe de 256 entités, slot matrice padded à 256 o *(corrigé à la validation : initialement 4096, puis 1024)* + +- `MAX_ENTITIES = 256` (constante pub dans `conf.rs`). +- **Pourquoi 256 (et pourquoi le slot matrice est padded à 256 o)** : le buffer de matrices + mondes est lié au slot `uniform` « object » (group 1) du pipeline de rendu, et WebGPU impose + **deux** contraintes : + 1. une binding `uniform` unique est plafonnée à `max_uniform_buffer_binding_size` (64 ko) ; + 2. un offset de buffer `uniform` (dynamique **ou** statique) doit être un multiple de + `min_uniform_buffer_offset_alignment` (256 o). + + Une matrice de 64 o ne peut donc jamais être adressée individuellement par un offset + `uniform` (son offset serait `slot × 64`, pas multiple de 256). La solution : **padded chaque + slot de matrice à 256 o** (`MatSlot { m: mat4x4f, pad: array }` = 64 + 192 o). Avec + des slots de 256 o, l'offset du slot i est `i × 256` (toujours aligné), et 256 slots × 256 o = + 64 ko est le maximum adressable par une binding `uniform` unique. (Les buffers transforms / + bboxes / indirect sont des bindings `storage` — limite 128 Mo, pas de règle d'offset 256 o — + donc ils gardent leurs tailles naturelles 64 / 32 / 80 o.) +- **Bind group « object » sur une slice de 64 o** : le bind group `matrix_object_bg` lie une + **slice de 64 o** (une matrice) du buffer, **pas** le buffer entier — une binding sur tout le + buffer (65536 o) plafonnerait l'offset dynamique à 0 (le binding couvrirait déjà tout le + buffer). Avec une slice de 64 o, l'offset dynamique peut glisser jusqu'à `65536 − 64`. +- Justification mémoire : 16 ko (transforms) + 64 ko (matrices) + 8 ko (bboxes) + 20 ko + (indirect) ≈ 108 ko total — négligeable. Le compute dispatche 256 threads (4 workgroups de + 64) — trivial. +- 256 est largement suffisant pour le scope du projet (le demo a 7 entités). Pour aller au-delà, + il faudrait chunker le buffer de matrices en tranches ≤ 64 ko (≤ 256 slots chacune) avec une + binding `uniform` par tranche (reporté — hors scope v1). + +### D13 — Buffers de culling séparés, pas d'impact sur le layout de rendu + +- `CullUniformsBuffer` (112 o) est dédié au compute (group **2** du compute, binding 0) ; il n'apparaît + **pas** dans le layout des pipelines de rendu → le layout de rendu (groups 0..3) est + inchangé. +- Le `BBoxBuffer` (8 ko) est group 2 du compute (binding 1, avec `draws` au binding 2) → pas dans + le layout de rendu. Le group 0 (transforms, storage read) est **partagé** par les deux entry points. + +### D14 — Incident « fenêtre noire » : l'ordre des arguments de `select` en WGSL (2026-09-22) + +- **Symptôme** : le `demo` (culling activé, D8) affichait une **fenêtre noire**. Readback GPU + (`debug_dump`) : les 7 slots de draw args avaient un compte de sommets à **0** (`.a.x = 0`) alors + que `culling = 1`, `num_slots = 7`, et que les entrées vues par le GPU (transforms, bboxes, plans + du frustum) étaient corrects — le test sphère simulé **côté CPU** passait pour les 7 entités + (distances de plans toutes ≥ +2.06), ce qui rendait l'échec inexplicable côté données. +- **Cause racine** : la convention des arguments de `select` en WGSL est + **`select(reject, accept, cond)`** — renvoie le **second** argument quand `cond` est vrai, le + premier quand il est faux : **l'inverse de la convention HLSL** (`select(trueVal, falseVal, cond)`) + sur laquelle le pass `cull` avait été écrit. `select(u32(t.flags.z), 0u, visible)` mettait donc le + compte à 0 pour **toute entité visible** (et aurait dessiné les entités cullées) : inversion + silencieuse de la visibilité, sans aucune erreur de validation wgpu ni de log driver. +- **Preuves** : (1) instrumentation par slot du pass (distances de plans toutes positives, flag + `visible` à 0 — contradiction pure) ; (2) sonde constante `select(2.0, 3.0, true)` → **3.0** sur le + GPU (NVIDIA, Vulkan) ; (3) source de **naga 30** (le compilateur WGSL de wgpu) : le frontend mappe + `arg0 → reject`, `arg1 → accept`, et les backends SPIR-V/GLSL émettent `cond ? accept : reject` + → conforme à la spec WGSL, **le driver n'est pas en cause** : c'est une erreur d'API WGSL. +- **Correction** : `select(0u, u32(t.flags.z), visible)` (visible ⇒ compte plein, cullé ⇒ 0) + + commentaire GOTCHA en tête de `gpu_driven.wgsl` + piège documenté dans `AGENTS.md`. +- **Vérification** : readback après correction — les 7 entités du `demo` reprennent leurs comptes + pleins (6/36/3840/960/384/192/2304, égaux aux `flags.z` packés côté CPU) ; une entité ajoutée hors + frustum (z = 60, au-delà du plan lointain) est **cullée à 0** ; les slots ≥ `num_slots` restent à + 0. `cargo build --workspace` sans avertissement, `cargo test --workspace` vert (63 tests). +- **Leçon** : ne pas suspecter le driver avant d'avoir vérifié les conventions d'arguments des + builtins WGSL ; pour ce shader, la vérité de référence est le **readback des slots** + (`debug_dump`, opt-in `WSG_DEBUG_DUMP=1` dans le `demo`), pas l'image à l'écran. + +## Pipeline par frame (un seul CommandEncoder) + +``` +[CPU] write_buffer TransformBuffer (64 o × num_slots, depuis Scene) +[CPU] write_buffer CullUniforms (6 plans + num_slots + culling) +[CPU] (si gpu_generation changée) write_buffer BBoxBuffer +[Compute 1] compute_matrices : m[i] = T·R·S (i in 0..num_slots) +[Compute 2] cull : draws[i] = visible ? args : 0 (sphère vs 6 plans) +[Render pass ombre] (si ombres activées) : draw_indexed_indirect par caster (slice mat + args slot) +[Render pass main] : draw_indexed_indirect / draw_indirect par slot vivant +[Queue] submit(encoder) +``` + +Même encoder → exécution séquentielle garantie (write → compute → render). + +## Layouts GPU (WGSL) + +```wgsl +struct TransformSlot { translation: vec3f, flags: vec4f, rotation: vec4f, scale: vec3f } // 64 o + // flags.x = mesh_index (index dans BBoxBuffer), flags.y = active (1/0), + // flags.z = draw count (compte de sommets/indices du mesh), flags.w = has_index (1/0) +struct MatSlot { m: mat4x4, pad: array, 12> } // 256 o (padded : offset `uniform` 256-aligné) +struct BBoxSlot { min: vec3f, max: vec3f } // 32 o (coins locaux ; centre/demi-extents dérivés dans le shader) +struct DrawSlot { a: vec4, b: vec4, c: vec4, d: vec4, e: vec4 } // 80 o +struct CullUniforms { + planes: array, 6>, num_slots: u32, culling: u32, _pad: vec2 // 112 o +} +``` + +Alignements vérifiés : stride de tableau = max(alignment) arrondi à 16 → 64 / 64 / 32 / 80 +(tous multiples de 16). `CullUniforms` : 6×16 + 4 + 4 + 8 = 112 (multiple de 16). + +Côté Rust (`uniform.rs`, `#[repr(C)]` + bytemuck `Pod`/`Zeroable`) : même taille, avec +`#[repr(align(16))]` sur les structs contenant un tableau de `vec4` pour matcher l'alignement +GPU. Static asserts `size_of == 64/64/32/80/112`. + +## Fichiers touchés + +| Fichier | Action | +|---------|--------| +| `lib/assets/shaders/gpu_driven.wgsl` | **Nouveau** — compute `compute_matrices` + `cull` | +| `lib/src/math/frustum.rs` | **Nouveau** — extraction des 6 plans (Gribb-Hartmann [0,1]) + tests | +| `lib/src/math/geometry.rs` | + struct `BBox` + `Geometry::bbox()` + tests | +| `lib/src/math/mod.rs` | + `pub mod frustum;` | +| `lib/src/resources/uniform.rs` | + `TransformSlot`, `BBoxSlot`, `CullUniforms` (Pod) + constantes `MAX_GPU_ENTITIES`, `MAX_MESH_BBOXES` + tailles de slots | +| `lib/src/resources/mesh.rs` | + champ `bbox: BBox` (calculé dans `from_geometry`) + accès `mesh.bbox()` | +| `lib/src/scene/scene.rs` | `entities` → slots append-only + index + génération ; `create_mesh`/`add_mesh` capturent le bbox ; + `iter_entity_slots()`, `gpu_generation()`, `gpu_bbox_slots()` | +| `lib/src/core/renderer.rs` | + 4 buffers GPU + buffer cull-uniforms + 2 pipelines compute + layout compute ; `render_scene` → `&mut self`, passe indirect ; + `set_culling`/`culling_enabled` ; shadow pass indirect | +| `lib/src/pipeline/pipeline_cache.rs` | + `create_compute_pipelines()` (module gpu_driven, 2 entry points) + `create_compute_bind_group_layout()` | +| `lib/src/app.rs` | + `AppBuilder::with_culling(bool)` + champ `culling` ; `resumed()` applique `renderer.set_culling` | +| `lib/src/utils/conf.rs` | + fallback embarqué `include_str!` de `gpu_driven.wgsl` (comme standard/shadow) | +| `lib/examples/demo.rs` | + `AppBuilder::with_culling(true)` (démo du culling) | +| `lib/tests/wgsl_validate.rs` | + test `gpu_driven_shader_is_valid_wgsl` | +| `lib/src/assets/shaders/standard_shader.wgsl`, `shadow_shader.wgsl` | Commentaires FR → EN (convention doc anglaise) | +| `docs/user/gpu-driven.md` | **Nouveau** — doc user (activation du culling, limites, comportement) | +| `docs/user/README.md` | + lien vers `gpu-driven.md` | +| `README.md` | + ligne « rendu GPU-driven (indirect draw, culling GPU) » dans les features | +| `docs/ROADMAP.md` | Coche 3.1, 3.2, 3.3 | +| `docs/PLAN.md` | + ligne Étape 17 | + +## Périmètre exclus (reportés) + +- **Multi-instancing par groupe de mesh** (Phase 4.4) — 1 draw par groupe de mesh identique. +- **AABB exact transformé** (8 sommets) — v1 utilise une sphère conservative. +- **Ring-buffering** des buffers GPU (D4) — v1 fait un `write_buffer` par frame. +- **Occlusion culling**, **LOD**, **batching par matériau** (Phase 4.3 / 4.4). +- **Multi-caméras** (2.1 restant). + +## Risques & mitigations + +| Risque | Mitigation | +|--------|-----------| +| **Alignement indirect draw** (offset multiple de 16/20/8/4 o) | D3 : slot 80 o = pdc(16,20), multiple de 8 et 4 → valide pour les 2 variants de draw. | +| **Buffer undefined à la création** | `IndirectArgsBuffer` zéro-initialisé (`create_buffer_init` zéros) → `c..e` restent 0. | +| **Ordre compute → render** | Même `CommandEncoder` → séquentiel. `write_buffer` + compute + render dans le même encoder. | +| **Bind groups par slot : mémoire** | 4096 bind groups ≈ trivial (état driver partagé via le layout). | +| **Bug de culling (objet qui disparaît)** | D8 : OFF par défaut ; le `demo` l'active → bug visible immédiatement. Sphère conservative (D5). **Réalisé en 2026-09-22 (D14)** : le bug s'est produit (fenêtre noire) et a été traçable grâce à cette mitigation + readback. | +| **Builtins WGSL aux conventions d'arguments non intuitives** (ex. `select`, l'inverse de HLSL) | D14 : piège documenté en tête du shader + dans `AGENTS.md` ; diagnostic par **readback des slots** (`debug_dump`) plutôt que par l'image seule. | +| **Ordre non-déterministe des entités** | D9 : `Vec` append-only (stable) remplace le `HashMap` (ordre aléatoire). | +| **`render_scene` → `&mut self`** | Aucun exemple n'appelle `Renderer::render_scene` directement (cf. D11) → non-régression. | +| **WGSL compute non validé** | Test `wgsl_validate` sur le nouveau shader + `cargo build` (wgpu compile à l'exécution). | +| **Shadow pass indirect** | Les casters utilisent les mêmes slots d'args → cohérent ; testé dans le `demo` (ombres + culling actifs). | +| **Alignement WGSL ≠ Rust** | `#[repr(align(16))]` + static asserts de taille sur chaque struct Pod (cf. section « Layouts GPU »). | + +## Plan de vérification + +1. `cargo build --workspace` — OK. +2. `cargo test --workspace` — 45 tests unitaires + **nouveaux tests frustum/geometry** + + 2→3 tests WGSL + 3 doctests, tous verts. +3. `cargo fmt --all -- --check` — clean. +4. **`demo` en headless** : `WGPU_BACKEND=vulkan timeout 10 cargo run -p examples --example + demo` → exit 0, **culling actif** (`with_culling(true)`), pas d'objet qui disparaît. +5. **Tous les autres exemples** (cube, simple, shadow_test, spot_test) en headless → exit 0, + **rendu identique** (culling OFF par défaut, chemin indirect actif). +6. **`manual.rs`** en headless → exit 0 (chemin `render()` bas niveau inchangé, UBO CPU). +7. **Check liens** : 0 lien cassé (nouveaux `docs/user/gpu-driven.md` + liens README). +8. **Aucun caractère accentué** dans les fichiers modifiés (sauf `docs/tech/`, DRAFT, PLAN, + ROADMAP, DOCUMENTATION) — les commentaires WGSL traduits. +9. **WGSL valide** : le nouveau `gpu_driven.wgsl` compile (test `wgsl_validate`). + +## Critères d'acceptation (definition of done) + +> **Implémenté le 2026-07-20.** Tous les critères sont remplis (avec les écarts de nommage +> notés ci-dessous). `cargo build`/`cargo test --workspace` : 56 lib + 3 WGSL + 3 doctests verts, +> clippy sans avertissement dans `renderer.rs`. + +- [x] `compute_matrices` + `cull` dans `gpu_driven.wgsl` (2 entry points, 1 module). + *Écart : layout compute explicite à 3 groupes partagé par les 2 pipelines (évite les gaps de + groupes du layout inféré) — cf. D11.* +- [x] 4 buffers GPU (transform, matrices, bboxes, draw-args) + 1 buffer cull-uniforms dans le `Renderer`. +- [x] `render_scene` 100 % indirect (main + shadow). + *Écart : `render_scene` reste `&self` (buffers persistants créés dans `new()` + `Cell` + pour le culling) — plus simple que le `&mut self` prévu, et sans ré-allocation par frame.* +- [x] `Scene` slots append-only + tombstones (indices stables). + *Écart de nommage : `packed_transform_slots()` / `iter_slot_draws()` / `mesh_bboxes()` / + `mesh_index_of()` / `mesh_by_index()` / `num_slots()` / `num_active_slots()` (plutôt que + `iter_entity_slots` / `gpu_generation` / `gpu_bbox_slots`).* +- [x] `Geometry::bbox()` + `BBoxSlot` ; upload du buffer de bboxes à chaque frame (peu coûteux, + toujours correct si des meshes sont ajoutés). +- [x] `math/frustum.rs` (Gribb-Hartmann [0,1]) + 5 tests. +- [x] `AppBuilder::with_culling(bool)` + `Renderer::set_culling` (OFF par défaut). +- [x] `demo` active le culling (`with_culling(true)`) ; `cargo build --workspace` vert. +- [x] Commentaires WGSL EN ; aucun accent dans les fichiers de code touchés. +- [x] Docs user (`docs/user/gpu-driven.md`), README, ROADMAP (3.1/3.2/3.3 cochés). +- [ ] DRAFT.md vidé après validation utilisateur (conservé en référence pour l'instant). + +--- +**Validation** : les décisions D1–D14 sont implémentées (D1 = draw indirect par entité, + 256 slots, culling OFF par défaut). Le bug « fenêtre noire » (D14) est **corrigé et vérifié par + readback GPU** le 2026-09-22 (comptes pleins pour les entités visibles, compte 0 pour l'entité hors + frustum, slots ≥ `num_slots` à 0 ; 63 tests verts). Il reste à confirmer le rendu **visuel** du + `demo` avec culling ON, puis autoriser le vidage de ce DRAFT. \ No newline at end of file diff --git a/docs/PLAN.md b/docs/PLAN.md index 5c32773..112465c 100644 --- a/docs/PLAN.md +++ b/docs/PLAN.md @@ -63,7 +63,7 @@ Une fois la plomberie encapsulée, nous devons rendre l'assemblage des objets co ### Gestion des Matériaux et Shaders - [X] S'assurer que chaque Mesh possède une référence vers un Material (à l'heure actuelle le lien est porté par l'entité `(mesh_id, material_id)` de la Scene, pas par le Mesh lui-même). *(fait — 2026-09-17, DRAFT Étape 7 : `Mesh.material: Option>` ; `Entity { mesh_id, transform }`, plus de `material_id`)* -- [X] Implémenter le comportement par défaut : si aucun matériau n'est assigné, le moteur injecte automatiquement le `standard_shader` (variante unlit) (non implémenté). *(fait — 2026-09-17, DRAFT Étape 7.3.5 : `Scene::default_material()` injecte `standard` ; le flat reste piloté par `Renderer::set_unlit`)* +- [X] Implémenter le comportement par défaut : si aucun matériau n'est assigné, le moteur injecte automatiquement le `standard_shader` (variante unlit). *(fait — 2026-09-17, DRAFT Étape 7.3.5 : `Scene::default_material()` injecte `standard` ; le flat reste piloté par `Renderer::set_unlit`)* ## Phase 3 : Documentation et Interface (API "User-Friendly") diff --git a/docs/ROADMAP.md b/docs/ROADMAP.md index 3a52a96..4321d5c 100644 --- a/docs/ROADMAP.md +++ b/docs/ROADMAP.md @@ -117,24 +117,25 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z } --- -## Phase 3️⃣ — GPU-Driven Rendering +## Phase 3️⃣ — GPU-Driven Rendering ✅ (2026-07-20) **Objectif** : Déléguer les calculs de transformation et culling au GPU (suivre ARCHI_CPU_GPU.md). +**Statut** : implémenté (DRAFT Étape 17, décisions D1–D14). Culling **désactivé par défaut** (non-régression), opt-in `AppBuilder::with_culling(true)`. **Correction 2026-09-22** : bug « fenêtre noire » avec culling ON (arguments de `select` WGSL écrits à la convention HLSL — toutes les entités visibles étaient remises à 0) ; corrigé et vérifié par readback GPU (DRAFT D14). ### 3.1 Compute Shader -- [ ] Buffer `TransformBuffer` (CPU → GPU) : positions/rotations/échelles brutes -- [ ] Buffer `MatrixBuffer` (GPU calculé) : World Matrices finales -- [ ] Compute shader : calcul des World Matrices pour tous les meshes +- [x] Buffer `TransformBuffer` (CPU → GPU) : positions/rotations/échelles brutes (`TransformSlot`, 64 B) +- [x] Buffer `MatrixBuffer` (GPU calculé) : World Matrices finales (`MatSlot`, `STORAGE|UNIFORM`) +- [x] Compute shader : calcul des World Matrices pour tous les meshes (`compute_matrices`) ### 3.2 Frustum Culling GPU -- [ ] Ajouter `BBox` dans `Geometry` (center + extents) -- [ ] Buffer `BoundingBoxBuffer` (CPU → GPU, statique) -- [ ] Compute shader : culling basé sur la frustum de caméra -- [ ] Buffer `IndirectDrawBuffer` rempli par le GPU +- [x] Ajouter `BBox` dans `Geometry` (coins min/max locaux) + `math::Frustum` (Gribb–Hartmann `[0,1]`) +- [x] Buffer `BoundingBoxBuffer` (CPU → GPU, ré-upload quand l'ensemble des meshes change, peu coûteux) +- [x] Compute shader : culling sphère vs frustum (`cull`), **désactivé par défaut** *(bug « fenêtre noire » corrigé le 2026-09-22 — ordre des arguments de `select` WGSL inversé ; cf. DRAFT D14)* +- [x] Buffer `IndirectDrawBuffer` rempli par le GPU (`DrawSlot`, 80 B, zéro = no-op) ### 3.3 Rendu Indirect -- [ ] `draw_indexed_indirect()` au lieu de draw calls individuels -- [ ] Un seul command draw pour tous les objets visibles +- [x] `draw_indexed_indirect()`/`draw_indirect()` au lieu de draw calls individuels +- [x] Un draw indirect **par slot actif** (décision D1 — pas un draw fusionné unique) ; le shadow pass est aussi indirect --- @@ -198,3 +199,4 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z } | **String IDs pour le MVP, slotmap reporté** | Le code et le README utilisent des String IDs (simples, sûrs, figés avant la boucle de rendu) ; `ARCHI_ARENES.md` reste la cible "handles typés" pour plus tard. La dépendance `slotmap` a été retirée tant qu'elle est inutilisée | | **Present mode FIFO figé pour l'instant** | Le swapchain utilise `PresentMode::Fifo` avec `desired_maximum_frame_latency: 2` (double buffering vsync) — défaut sûr : pas de tearing, énergie minimale, zéro artefact. On **gèle ce choix** ; `Mailbox` (triple buffering) pourra être exposé en option et `Immediate` restera réservé à l'offscreen, **on s'occupera du present mode le moment venu** (quand le pipeline GPU-driven arrivera, Phase 3) — ce n'est pas bloquant pour les étapes 1-2 | | **Resize géré (avec recréation de la depth texture), acté en D3 (2026-09-18), réalisé en Étape 11 (2026-09-18)** | L'app reconfigure désormais la surface et recrée la depth texture **en même temps** à chaque `Resized` (`App::resize` → `Context::configure` + `Renderer::resize_depth`), via le helper `create_depth_texture` isolé. Vérifié au runtime (exemple `cube`) : pas de crash, pas d'artefact, aspect correct | +| **WGSL `select(reject, accept, cond)`** | L'ordre des arguments est l'inverse de la convention HLSL : le **second** argument est retenu quand la condition est vraie. L'avoir écrit à la convention HLSL a produit le bug « fenêtre noire » du culling (comptes remis à 0 pour les entités visibles), corrigé le 2026-09-22 (DRAFT D14). Piège documenté en tête de `gpu_driven.wgsl` + `AGENTS.md` | diff --git a/docs/tech/ARCHI_APP.md b/docs/tech/ARCHI_APP.md index 3723973..39c99ca 100644 --- a/docs/tech/ARCHI_APP.md +++ b/docs/tech/ARCHI_APP.md @@ -15,17 +15,20 @@ stale_after: 2027-01-31 wsg_lib est un moteur de rendu modulaire basé sur wgpu. Il adopte une architecture à deux niveaux : une façade de haut niveau pour la productivité et un accès bas niveau pour un contrôle total. -> **État du document : ACTUEL pour la façade (`App`/`AppHandler`, §3, §4A) ; CIBLE pour la partie -> GPU-driven (§1, §4B, §5, §6).** La façade `AppBuilder`/`App`/`AppHandler` est livrée et est le -> **workflow recommandé** : `setup` (déclaration de la scène) → par frame `update` (mutation) → +> **État du document : ACTUEL** — façade (`App`/`AppHandler`, §3, §4A) et pipeline GPU-driven +> (§1, §4B, §5, §6) **implémenté en Phase 3** du ROADMAP (2026-07-20, DRAFT Étape 17, décisions +> D1–D14). La façade `AppBuilder`/`App`/`AppHandler` est livrée et est le **workflow recommandé** : +> `setup` (déclaration de la scène) → par frame `update` (mutation) → > `render` (défaut : `App::render_scene` = itération des entités + **rendu groupé en une passe**, > un `CommandEncoder`/soumission par frame ; passe d'ombre en tête si un caster est actif). > Exemples : `simple` (2D unlit), `cube` (3D éclairé), `demo` (vitrine : primitives, lumières, -> ombres, caméra orbitale). Le workflow **manuel** (exemple `manual`) coexiste pour le contrôle fin. -> Les sections §1, §4B, §5 et §6 décrivent la **cible** : pipeline GPU-driven à deux passes -> (Compute Pass → `draw_indexed_indirect`), buffers persistants en VRAM (Transform/Matrix/BBox/Indirect) -> et synchronisation single/double buffer. **Rien de tout cela n'existe encore dans le code** — c'est -> la trajectoire ROADMAP Phase 3. +> ombres, caméra orbitale, culling GPU activé). Le workflow **manuel** (exemple `manual`) coexiste +> pour le contrôle fin. +> Les sections §1, §4B, §5 et §6 décrivent le pipeline GPU-driven **tel qu'implémenté**, avec les +> écarts documentés (DRAFT Étape 17) : un draw indirect par slot (D1), table fixe de 256 slots +> (D12), culling par sphère conservative (D5), single buffer (D4), et le piège de l'ordre des +> arguments de `select` en WGSL (D14, bug « fenêtre noire » corrigé le 2026-09-22). La section +> « Notes pour l'implémentation future » (double buffering) reste **CIBLE**. ## 1. Philosophie et Principes @@ -107,8 +110,8 @@ Le moteur gère la renderloop interne via un pipeline à **deux passes séquenti 1. **Update** (`AppHandler::update`) — L'utilisateur modifie la scène (transformations, entités). Ces changements sont synchronisés vers le GPU via un **single buffer** Transform avant la passe de calcul. > La synchronisation est assurée par le pipeline wgpu : `queue.submit()` après le compute pass garantit que les données Transform sont valides avant le render pass suivant. Aucun double buffering n'est nécessaire tant que la latence maximale de la surface (via `desired_maximum_frame_latency`) est ≥ 3. -2. **Compute Pass** — Un compute shader lit les Transform bruts, calcule les World Matrices finales, effectue le Frustum Culling par AABB, et remplit l'Indirect Draw Buffer avec les identifiants des objets visibles. -3. **Render Pass** — Le CPU émet une unique commande `draw_indexed_indirect`. Le GPU pioche dans l'Indirect Draw Buffer et dessine uniquement les objets visibles, sans intervention du CPU. +2. **Compute Pass** — Deux entry points compute séquentiels (`compute_matrices` puis `cull`, un seul module WGSL) lisent les Transform bruts, calculent les World Matrices finales, effectuent le Frustum Culling par **sphère conservative** (D5), et remplissent l'Indirect Draw Buffer avec les **comptes** de draw des objets visibles (0 si cullé/inactif). +3. **Render Pass** — Le CPU émet **un draw indirect par slot** (écart D1 — la cible initiale prévoyait une commande unique fusionnée). Le GPU pioche les comptes dans l'Indirect Draw Buffer et dessine uniquement les objets non cullés et actifs, sans intervention du CPU. 4. **Présentation** — La surface est présentée à l'écran. L'ordre d'appel des méthodes sur le `CommandEncoder` (`begin_compute_pass` puis `begin_render_pass`) garantit l'exécution séquentielle. Les barrières de mémoire entre passes sont insérées automatiquement par le pilote. @@ -117,10 +120,11 @@ L'ordre d'appel des méthodes sur le `CommandEncoder` (`begin_compute_pass` puis | Buffer | Rôle | Type wGPU | Direction du flux | |--------|------|-----------|-------------------| -| Transform Buffer | Positions/rotations/échelles brutes | Storage Buffer | CPU → GPU | -| Matrix Buffer | World Matrices finales calculées | Storage Buffer | GPU (Calculé) → GPU (Lu par Render) | -| Bounding Box Buffer | AABB de chaque mesh pour culling | Storage Buffer | CPU → GPU (Statique) | -| Indirect Draw Buffer | Liste dynamique des objets à dessiner | Indirect + Storage | GPU (Rempli par Compute) → GPU (Lu par Render) | +| Transform Buffer | Positions/rotations/échelles brutes + flags par entité (64 o/slot) | Storage Buffer | CPU → GPU (chaque frame, `write_buffer`) | +| Matrix Buffer | World Matrices finales calculées (256 o/slot, padded — D12) | Storage + Uniform | GPU (Calculé) → GPU (Lu par Render) | +| Bounding Box Buffer | Coins min/max de l'AABB de chaque mesh (32 o/slot) | Storage Buffer | CPU → GPU (quand l'ensemble des meshes change) | +| Indirect Draw Buffer | Comptes de draw par slot (80 o/slot, zéro = no-op) | Indirect + Storage | GPU (Rempli par Compute) → GPU (Lu par Render) | +| CullUniforms | 6 plans du frustum + `num_slots` + `culling` (112 o) | Uniform Buffer | CPU → GPU (chaque frame) — réservé au compute (group 2) | > **Synchronisation single buffer** : Les buffers Transform et Matrix utilisent un **single buffer** en phase initiale. Le CPU écrit dans le buffer pendant `update()`, puis le compute shader lit les données au frame suivant via `queue.submit()` qui garantit la séquence d'exécution. Cette approche fonctionne correctement tant que la surface a une latence maximale ≥ 2 frames (configuré via `desired_maximum_frame_latency`). Le double buffering sera ajouté uniquement si des artefacts visuels apparaissent à haute fréquence (typiquement > 90 fps sur machines rapides). diff --git a/docs/tech/ARCHI_CPU_GPU.md b/docs/tech/ARCHI_CPU_GPU.md index db06913..6e0d5d4 100644 --- a/docs/tech/ARCHI_CPU_GPU.md +++ b/docs/tech/ARCHI_CPU_GPU.md @@ -7,7 +7,7 @@ actor: person/jerome sources: [] generated: { by: human:jerome, at: 2026-07-31T00:00:00Z } verified: true -status: target +status: current stale_after: 2027-01-31 --- @@ -16,13 +16,19 @@ Bonnes Pratiques & Guide d'Implémentation Ce document sert de spécification technique et de trame d'implémentation pour l'architecture de rendu 3D pilotée par le GPU (GPU-Driven Rendering) utilisant wgpu. L'objectif est de déléguer un maximum de charges de calcul au GPU pour soulager le CPU et maximiser les performances de parallélisme. -> **État du document : CIBLE (spécification du pipeline GPU-driven, non implémenté).** +> **État du document : ACTUEL (implémenté — Phase 3 du ROADMAP, 2026-07-20, DRAFT Étape 17).** > La répartition CPU/GPU, le compute pass (World Matrices + Frustum Culling), l'Indirect Draw Buffer -> et les buffers persistants en VRAM décrits ici correspondent à la **Phase 3 du ROADMAP** et aux -> README étapes 2-3. **Aucun de ces mécanismes n'existe encore dans le code.** Aujourd'hui le rendu est -> piloté par le CPU, **objet par objet** (une soumission par mesh, voir README.md et l'exemple `manual`). -> Considérez ce document comme la spécification de référence pour l'implémentation future du pipeline -> GPU-driven, pas comme une description de l'état actuel. +> et les buffers persistants en VRAM décrits ici sont en place : `shaders/gpu_driven.wgsl` +> (deux entry points `compute_matrices` + `cull`, un module, layout explicite à 3 groupes) et les +> buffers de slots du `Renderer` (`TransformSlot`/`MatSlot`/`BBoxSlot`/`DrawSlot`/`CullUniforms`, +> capacité fixe de 256 slots). +> **Écarts documentés** (cf. DRAFT Étape 17) : (D1) un draw indirect **par slot** plutôt qu'une +> commande unique fusionnée ; (D12) 256 slots, slot matrice padded à 256 o (plafond `uniform` WebGPU) ; +> (D5) culling par **sphère** conservative dérivée de l'AABB locale du mesh, pas par l'AABB transformée +> exacte ; (D4) single buffer, pas de double-buffering. +> **Piège connu (2026-09-22, D14)** : l'ordre des arguments de `select` en WGSL est l'inverse de la +> convention HLSL — l'avoir inversé a produit un bug « fenêtre noire » (entités visibles remises à 0), +> corrigé et vérifié par readback GPU. Documenté en tête de `gpu_driven.wgsl` et dans `AGENTS.md`. 1. Répartition des Rôles : CPU vs GPU (La Source de Vérité) @@ -53,10 +59,10 @@ L'exécution des tâches s'appuie sur une structure séquentielle stricte au sei - Mise à jour CPU (Minimaliste) : Le CPU écrit les transformations brutes (Transform) modifiées dans un buffer GPU mappé (single buffer en phase initiale — la synchronisation est assurée par `queue.submit()` qui garantit la séquence d'exécution). Double buffering sera ajouté uniquement si des artefacts apparaissent à haute fréquence (> 90 fps). - Pass de Calcul (Compute Pass) : - Calcul des World Matrices : Un compute shader lit les transformations brutes et génère la matrice 4x4 finale pour chaque mesh. - - Frustum Culling GPU : Le même compute shader (ou un compute pass dédié) compare la Bounding Box (AABB) de chaque objet avec les plans de la caméra (matrice de projection/vue). - - Remplissage du Buffer Indirect : Si l'objet est visible, son identifiant est injecté dans un buffer de commandes de dessin indirect (Indirect Draw Buffer). + - Frustum Culling GPU : Un compute pass dédié (`cull`) compare la **sphère bounding** de chaque objet (D5 — conservative, dérivée de l'AABB locale du mesh et de l'échelle de l'entité) avec les 6 plans du frustum de la caméra. + - Remplissage du Buffer Indirect : le pass `cull` écrit le **compte de sommets/indices** de chaque objet dans son `DrawSlot` (80 o) — mis à 0 si l'objet est cullé ou inactif (no-op). - Pass de Rendu (Render Pass) : - - Le CPU émet une unique commande globale : draw_indexed_indirect. + - Le CPU émet **un draw indirect par slot** (écart D1 — la spécification initiale prévoyait une commande unique fusionnée) ; les slots à compte 0 (cullés/inactifs/vides) sont des no-ops. - Le GPU pioche directement dans le buffer préparé par le compute pass et dessine uniquement les objets visibles, sans intervention du CPU. 3. Stratégie de Synchronisation @@ -66,12 +72,15 @@ L'exécution des tâches s'appuie sur une structure séquentielle stricte au sei 4. Synthèse des Structures de Données en VRAM -Pour implémenter cette architecture, prévoyez l'utilisation des buffers wGPU suivants : -Nom du Buffer,Rôle,Type wGPU,Direction du flux -Transform Buffer,Stocke les positions/rotations/échelles brutes.,Storage Buffer,CPU → GPU -Matrix Buffer,Stocke les World Matrices finales calculées.,Storage Buffer,GPU (Calculé) → GPU (Lu par le Render) -Bounding Box Buffer,Stocke les AABB de chaque mesh pour le culling.,Storage Buffer,CPU → GPU (Statique) -Indirect Draw Buffer,Contient la liste dynamique des objets à dessiner.,Indirect Buffer + Storage,GPU (Rempli par Compute) → GPU (Lu par Render) +L'implémentation utilise les buffers wGPU suivants (tous créés par le `Renderer` à l'initialisation, capacité fixe de 256 slots) : + +| Buffer | Rôle | Type wGPU | Direction du flux | +|--------|------|-----------|-------------------| +| Transform Buffer | Positions/rotations/échelles brutes + flags par entité (64 o/slot) | Storage Buffer | CPU → GPU (chaque frame, `write_buffer`) | +| Matrix Buffer | World Matrices finales calculées (256 o/slot, padded — D12) | Storage + Uniform Buffer | GPU (Calculé) → GPU (Lu par le Render) | +| Bounding Box Buffer | Coins min/max de l'AABB de chaque mesh (32 o/slot) | Storage Buffer | CPU → GPU (quand l'ensemble des meshes change) | +| Indirect Draw Buffer | Comptes de draw par slot (80 o/slot, zéro = no-op) | Indirect + Storage Buffer | GPU (Rempli par Compute) → GPU (Lu par le Render) | +| CullUniforms | 6 plans du frustum + `num_slots` + `culling` (112 o) | Uniform Buffer | CPU → GPU (chaque frame) — réservé au compute (group 2) | ## Liens diff --git a/docs/tech/FRAME_LOOP.md b/docs/tech/FRAME_LOOP.md index 65aef93..9818797 100644 --- a/docs/tech/FRAME_LOOP.md +++ b/docs/tech/FRAME_LOOP.md @@ -68,10 +68,11 @@ Avec notre nouvelle architecture "Atelier", la distinction est devenue encore pl | CommandEncoder | Par-Frame | Ton "carnet de notes" temporaire pour les ordres du GPU. | | TextureView | Par-Frame | Fenêtre temporaire sur la texture active du swapchain. | -> **Ressources GPU persistantes (single buffer) — CIBLE, non implémenté** : À l'état **visé**, les -> buffers Transform et Matrix vivent en VRAM avec un single buffer en phase initiale (le CPU écrit -> pendant `update()`, le compute shader lit au frame suivant, séquencé par `queue.submit()`), puis un -> double buffering si des artefacts apparaissent à haute fréquence. **Aucune de ces ressources n'existe -> encore dans le code** — c'est la cible GPU-driven (ROADMAP Phase 3 / ARCHI_CPU_GPU). +> **Ressources GPU persistantes (single buffer) — implémenté (Phase 3, 2026-07-20)** : les buffers +> Transform, Matrix, BBox et Indirect Draw vivent en VRAM (créés à l'initialisation du `Renderer`, +> capacité fixe de 256 slots). Le CPU écrit les transforms chaque frame par `queue.write_buffer` +> **dans le même `CommandEncoder`** que les compute passes, qui les lisent **dans la même frame** +> (l'ordre est garanti par l'encoder, pas par `queue.submit()` inter-frames). Le double buffering +> reste la **cible** si des artefacts apparaissent à haute fréquence (voir ARCHI_CPU_GPU / ARCHI_APP). --- diff --git a/docs/user/README.md b/docs/user/README.md index a266875..89c31cb 100644 --- a/docs/user/README.md +++ b/docs/user/README.md @@ -20,6 +20,7 @@ GPU graphics background is required. | [Materials & textures](materials.md) | Appearance: the `standard` shader, unlit mode, diffuse textures | | [Lights](lights.md) | Directional, point, spot, ambient, `MAX_LIGHTS` | | [Shadows](shadows.md) | Shadow mapping: picking the casting light, the packed-index pitfall | +| [GPU-driven rendering](gpu-driven.md) | GPU world matrices + indirect draws, opt-in frustum culling | | [Camera & input](camera-input.md) | Active camera, orbital controller, unified keyboard/mouse state | | [Examples](examples.md) | The 7 repo examples, the advanced `manual` workflow, adding your own example | diff --git a/docs/user/gpu-driven.md b/docs/user/gpu-driven.md new file mode 100644 index 0000000..bc5a223 --- /dev/null +++ b/docs/user/gpu-driven.md @@ -0,0 +1,132 @@ +# GPU-driven rendering + +WSG's scene rendering is **GPU-driven**: the per-entity world matrices and the indirect draw +arguments are computed on the GPU each frame, so the CPU no longer loops over entities to issue +draw calls. This page explains what that means for you and how to opt into **frustum culling**. + +## What runs on the GPU + +Each frame, before the render passes, two compute passes run over a fixed-capacity slot table +(256 entities, allocated once): + +1. **`compute_matrices`** derives each entity's world matrix from its transform + (translation / rotation / scale). The result feeds the render pipelines as the per-entity + model matrix. +2. **`cull`** decides per-entity visibility and fills the **indirect draw arguments** (the + vertex/index count, zeroed when the entity is culled or inactive). + +The main and shadow render passes are then **100 % indirect**: each active slot issues one +indirect draw that reads its own count and world matrix. A culled or inactive slot has a zero +count, so its draw is a no-op. The CPU only rewrites the transform slots and the cull uniforms +each frame — it never iterates the entities to issue draws. + +You do not need to do anything special to get this: `render_scene` is GPU-driven by default. + +## Frustum culling (opt-in) + +Culling is **off by default**. The culling pass still runs, but with culling disabled it marks +every active entity visible — so the rendered image is **identical** to a CPU-culled scene. +This protects you from a culling bug (an object that should be visible vanishing) becoming a +silent correctness issue. + +To enable culling, build your `App` with `.with_culling(true)`: + +```rust +let app = AppBuilder::new() + .title("My app") + .with_culling(true) // skip entities whose bounding sphere leaves the frustum + .build() + .await?; +``` + +Or toggle it at runtime on the renderer: + +```rust +app.renderer().set_culling(true); // enable +app.renderer().set_culling(false); // disable again +``` + +## How culling works + +When culling is on, each entity's **local-axis-aligned bounding box** (computed once from its +geometry, `Geometry::bbox()`) is treated as a **bounding sphere**: + +- **center** = the box center, transformed by the entity's world transform (rotation + + translation; scale is folded into the radius), +- **radius** = the box's circumradius scaled by the entity's largest scale component. + +The sphere is tested against the six camera frustum planes. If it is **fully outside** (beyond +a plane by more than its radius), the entity is culled; otherwise it is drawn. + +The sphere is a **conservative** approximation of the box: it can draw an object that is partly +out of view (false negative), but it will **never cull an object that is actually visible** +(false positive). For tight culling you would need per-mesh sphere fitting or per-face tests, +which are out of scope for v1. + +## Debugging the GPU path + +If something looks wrong — a missing object, a black window — the GPU-side slot tables can be +read back and printed. The `Renderer` ships a debug helper (intentionally **not** part of the +documented API): + +```rust +app.renderer().debug_dump(8); // prints the first 8 GPU slots to stderr +``` + +It dumps exactly what the GPU sees: the transform slots, the derived world matrices, the +indirect draw arguments, the mesh bounding boxes and the cull uniforms. A slot whose vertex +count reads `0` was zeroed by the cull pass (culled, inactive, or beyond `num_slots`); a full +count means the entity is drawn. In the `demo` example the dump is opt-in via an environment +variable, so the showcase stays silent by default: + +```sh +WSG_DEBUG_DUMP=120 cargo run -p wsg-lib --example demo +``` + +`WSG_DEBUG_DUMP=N` dumps for the first *N* frames. The demo stays **silent** when the variable is +unset; a set-but-non-numeric value (e.g. `WSG_DEBUG_DUMP=on`) gives 3 frames. + +**How to verify culling is actually working** (a correct culling pass is invisible — culled +objects were off-screen anyway — so the proof is in the counts, not the image): + +1. Launch the demo with `WSG_DEBUG_DUMP=120` (the demo has culling **on** and an orbiting + camera — drag the mouse to orbit). +2. Note first that orbiting/zooming this camera **cannot cull the entity ring**: the camera + always looks at the origin, so each entity's angular offset from the view axis is bounded + by `atan(ring radius / camera distance)` = `atan(1.7/6.1)` ≈ 15.5°, under the ~22° vertical + half-FOV. The seven demo entities therefore keep their **full** counts (cube `36`, sphere + `3840`, …) in every orientation — that is the expected and correct behaviour (verified + 2026-09-22: 600-frame camera sweep, GPU cull verdicts matched an independent CPU sphere + test on all 6000 entity frames, zero flips on the ring). +3. To see the counts actually flip to **`0`**, you need an entity well **off the target axis** + — e.g. one placed far away so it ends up behind the near plane. Its count then toggles + `0` ↔ full as the camera orbits, while the on-axis entities stay full. (This off-axis test + is the one that verified the cull path end-to-end, positive and negative.) +4. Optional A/B: temporarily build with `.with_culling(false)` and repeat — with culling off, + every entity keeps its full count in **every** orientation (the off-axis one included). + +This readback is the reference truth when a shader bug is suspected: it shows both the computed +counts and the raw inputs of the cull pass, independently of what ends up on screen. (It is how +the 2026-09-22 « black window » bug — an inverted WGSL `select` argument order — was diagnosed +and verified fixed, see the DRAFT D14 in `docs/DRAFT.md`.) + +## Limitations + +- **Culling is all-or-nothing per entity.** There is no partial (per-triangle) culling. +- **The sphere is a coarse bound** for elongated meshes (a long thin box gets a large sphere). + If your scene is dominated by such shapes, culling may bring little gain. +- **Capacity is 256 entities per render pass.** Beyond that, extra entities are not drawn. + This is the largest a single-buffer design can address under WebGPU's two `uniform` rules: a + single `uniform` binding is capped at 64 KB, *and* a `uniform` offset must be a multiple of 256 B. + A 64-byte matrix can never be individually addressable by a `uniform` offset, so each matrix + slot is padded to 256 B — and 256 slots × 256 B = 64 KB is the maximum. It is amply generous + for a simple scene (the demo has 7). +- **Mesh bounding boxes are recomputed when meshes are added**; a scene whose mesh set changes + at runtime simply re-uploads the small bbox table (a few bytes per mesh). + +Culling is a **performance** feature, not a visual one: with it off you get the same image with +the indirect-draw machinery still active. + +--- + +Next: [Examples](examples.md) · Back to [User documentation index](README.md) diff --git a/lib/examples/demo.rs b/lib/examples/demo.rs index 9e0bd53..7d29bc0 100644 --- a/lib/examples/demo.rs +++ b/lib/examples/demo.rs @@ -61,6 +61,8 @@ fn stripes_rgba() -> Vec { struct Demo { camera: CameraController, angle: f32, + /// Phase 3 black-window investigation: number of debug_dump calls already made. + dbg: u32, } /// Horizontal radius at which the primitives sit around the origin. @@ -222,13 +224,40 @@ impl AppHandler for Demo { tf.rotation = Quat::from_rotation_y(self.angle) * Quat::from_rotation_x(self.angle * 0.4); app.scene.set_entity_transform("cube_e", tf); } + + fn render(&mut self, app: &mut wsg_lib::App, frame: &wsg_lib::core::Frame) { + app.render_scene(frame.view()); + // Opt-in GPU readback (black-window investigation tooling): WSG_DEBUG_DUMP=N dumps the + // first 8 slots of the transform/matrix/draw-args/bbox buffers for N frames (unset = silent, + // non-numeric value = 3 frames). Note: orbiting/zooming this camera + // can never cull the entity ring — the camera always looks at the origin, so each + // entity's angular offset from the view axis is bounded by atan(1.7/6.1) ≈ 15.5°, under + // the ~22° vertical half-FOV (verified 2026-09-22: 600 frames swept, GPU==CPU on all + // 6000 cull verdicts, zero flips on the ring). Counts only flip to 0 for entities far + // off-axis (e.g. behind the near plane) — see docs/user/gpu-driven.md. + // Unset → 0 (the showcase stays silent); set but non-numeric (e.g. `WSG_DEBUG_DUMP=on`) → 3. + let frames = match std::env::var("WSG_DEBUG_DUMP") { + Ok(v) => v.parse::().ok().filter(|&n| n > 0).unwrap_or(3), + Err(_) => 0, + }; + if self.dbg < frames { + self.dbg += 1; + app.renderer().debug_dump(8); + } + } } #[pollster::main] async fn main() -> Result<(), WsgError> { - let app = AppBuilder::new().title("WSG Demo").build().await?; + // Culling enabled here (Step 15, D8) to exercise the GPU path; it is OFF by default elsewhere. + let app = AppBuilder::new() + .title("WSG Demo") + .with_culling(true) + .build() + .await?; app.run(Demo { camera: CameraController::default(), angle: 0.0, + dbg: 0, }) } diff --git a/lib/src/README.md b/lib/src/README.md index 125f6cb..97d0134 100644 --- a/lib/src/README.md +++ b/lib/src/README.md @@ -2,15 +2,16 @@ ## Overview -This is the source tree for `wsg-lib`, a Rust library wrapping [wgpu](https://github.com/gfx-rs/wgpu) for simple 3D drawing operations. The crate follows a layered architecture organized into eight modules: +This is the source tree for `wsg-lib`, a Rust library wrapping [wgpu](https://github.com/gfx-rs/wgpu) for simple 3D drawing operations. The crate follows a layered architecture organized into eight modules (plus the `shaders/` asset directory): | Module | Responsibility | |--------|---------------| -| **core** | Manager (Context) + Executor (Renderer) layers — GPU lifecycle and draw call orchestration; also `InputState` (unified keyboard/mouse input, Step 15.B) | -| **resources** | Data types: Vertex (CPU-side), Mesh (GPU geometry), Material (appearance descriptor), Texture, Lights, Camera + CameraController | +| **core** | Manager (Context) + Executor (Renderer) layers — GPU lifecycle and draw call orchestration (incl. the GPU-driven compute passes + opt-in frustum culling); also `InputState` (unified keyboard/mouse input, Step 15.B) | +| **resources** | Data types: Vertex (CPU-side), Mesh (GPU geometry + bounding box), Material (appearance descriptor), Texture, Lights, Camera + CameraController, and the uniform slot types (`TransformSlot`/`MatSlot`/`BBoxSlot`/`DrawSlot`/`CullUniforms`) | | **pipeline** | PipelineCache — WGSL shader loading and RenderPipeline compilation cache | -| **scene** | Scene — resource depot and entity graph for declarative rendering setup | -| **math** | Transform, Geometry (per-attribute mesh data) and `primitives` (procedural mesh generators) | +| **shaders** | Embedded WGSL sources (`standard`, `shadow`, `gpu_driven`) loaded via the `include_str!` fallback in `utils::conf` | +| **scene** | Scene — resource depot and slot-based entity graph for declarative rendering setup (Step 17) | +| **math** | Transform, Geometry (per-attribute mesh data + AABB), `Frustum` (Gribb–Hartmann, WebGPU `[0,1]` z) and `primitives` (procedural mesh generators) | | **utils** | Configuration constants and WsgError type | | **app** | App facade — high-level application orchestration with window lifecycle, event loop, and render automation | | **handler** | AppHandler trait — user-defined game logic interface injected into the render loop | diff --git a/lib/src/app.rs b/lib/src/app.rs index 1236144..346cc72 100644 --- a/lib/src/app.rs +++ b/lib/src/app.rs @@ -56,6 +56,8 @@ pub struct App { pub(crate) width: u32, /// Window height, read by the runner when the window is created in `resumed`. pub(crate) height: u32, + /// GPU frustum culling (Step 15, D8); applied to the renderer in `resumed`. + pub(crate) culling: bool, /// Winit event loop for window management. Set to None after run() consumes it. event_loop: Option>, // On met en Option pour pouvoir faire .take() facilement /// GPU hardware context — owns Instance, Surface, Adapter, Device, Queue lifecycle. @@ -118,6 +120,7 @@ impl App { title: self.title.clone(), width: self.width, height: self.height, + culling: self.culling, handler, app: None, }; @@ -173,6 +176,8 @@ pub struct AppBuilder { width: u32, /// Window height in pixels. height: u32, + /// GPU frustum culling enabled (Step 15, D8). Defaults to false (non-regression). + culling: bool, } impl AppBuilder { @@ -183,6 +188,7 @@ impl AppBuilder { title: APP_DEFAULT_TITLE.to_string(), width: APP_DEFAULT_WIDTH, height: APP_DEFAULT_HEIGHT, + culling: false, } } /// Sets the window title to display in the OS taskbar/window decorations. @@ -198,6 +204,14 @@ impl AppBuilder { self.height = height; self } + /// Enables GPU frustum culling (Step 15, D8). When true, entities whose bounding sphere is + /// fully outside the camera frustum are skipped (their indirect draw args are zeroed on the + /// GPU). Defaults to **off** (non-regression): the culling compute pass still runs but marks + /// every active entity visible, so the rendered image is identical to culling-off. + pub fn with_culling(mut self, enabled: bool) -> Self { + self.culling = enabled; + self + } /// Builds the configured `App` instance: creates the event loop and stores the window /// configuration. The GPU context, window and renderer are created later, when the event loop /// is resumed (inside `App::run`), because winit 0.30 only allows window creation in that phase. @@ -211,6 +225,7 @@ impl AppBuilder { title: self.title, width: self.width, height: self.height, + culling: self.culling, event_loop: Some(event_loop), context: None, renderer: None, @@ -229,6 +244,8 @@ struct AppRunner { width: u32, /// Window height in pixels, applied when the window is created in `resumed`. height: u32, + /// GPU frustum culling (Step 15, D8); applied to the renderer in `resumed`. + culling: bool, /// The user-provided game logic. handler: H, /// The fully-built App facade, populated on the first `resumed` event. @@ -262,6 +279,8 @@ impl ApplicationHandler for AppRunner { .expect("surface configuration failed"); let device = Arc::new(context.device.clone()); let renderer = Renderer::new(&context, format, self.width, self.height); + // Step 15, D8: apply the culling flag (off by default — non-regression). + renderer.set_culling(self.culling); // Step 7 (DRAFT 7.1): the PipelineCache now lives in the Scene. We wire the GPU context // (device + queue + format + cache) into the Scene before setup so it can build materials/meshes. @@ -274,6 +293,7 @@ impl ApplicationHandler for AppRunner { title: self.title.clone(), width: self.width, height: self.height, + culling: self.culling, event_loop: None, context: Some(context), renderer: Some(renderer), diff --git a/lib/src/core/renderer.rs b/lib/src/core/renderer.rs index b8755c2..25d3a63 100644 --- a/lib/src/core/renderer.rs +++ b/lib/src/core/renderer.rs @@ -20,27 +20,30 @@ use crate::core::Context; use crate::core::Frame; -use crate::math::Transform; +use crate::math::Frustum; use crate::pipeline::{ DEPTH_FORMAT, build_shadow_pipeline, create_shadow_map_bind_group_layout, create_shadow_uniform_layout, create_uniform_bind_group_layouts, }; -use crate::resources::uniform::{FRAME_UNIFORMS_SIZE, OBJECT_UNIFORM_SIZE, SHADOW_UNIFORM_SIZE}; +use crate::resources::uniform::{ + BBOX_SLOT_SIZE, BBoxSlot, CULL_UNIFORMS_SIZE, DRAW_SLOT_SIZE, DrawSlot, FRAME_UNIFORMS_SIZE, + MAT_SLOT_SIZE, MAX_LIGHTS, MatSlot, OBJECT_UNIFORM_SIZE, SHADOW_UNIFORM_SIZE, + TRANSFORM_SLOT_SIZE, TransformSlot, +}; use crate::resources::{ - Camera, FrameUniforms, Lights, MAX_LIGHTS, Material, Mesh, ObjectUniform, ShadowUniform, + Camera, CullUniforms, FrameUniforms, Lights, Material, Mesh, ObjectUniform, ShadowUniform, }; use crate::scene::Scene; use crate::utils::conf::{ - SHADOW_DEPTH_BIAS, SHADOW_MAP_SIZE, SHADOW_SCENE_CENTER, SHADOW_SCENE_RADIUS, + GPU_DRIVEN_SHADER, GPU_WORKGROUP_SIZE, MAX_ENTITIES, SHADOW_DEPTH_BIAS, SHADOW_MAP_SIZE, + SHADOW_SCENE_CENTER, SHADOW_SCENE_RADIUS, }; use glam::{Mat4, Vec3, Vec4}; -use std::cell::RefCell; -use std::collections::HashMap; +use std::cell::Cell; /// The Executor layer of the architecture. Holds shared references to Device and Queue from Context, /// plus the surface texture format. Executes WGPU rendering commands by binding Materials and Meshes /// into RenderPasses during each frame. Does not own raw hardware resources (they are Arc-cloned from Context). -/// /// Since Step 3 every pipeline declares the two uniform bind groups (frame @0 + object @1), the /// Renderer owns the matching GPU buffers and `BindGroup`s and binds them around every draw call. pub struct Renderer { @@ -58,8 +61,6 @@ pub struct Renderer { /// Depth attachment view used by both render passes (`render`, `render_scene`). Format /// `DEPTH_FORMAT` (Depth32Float) — matches every pipeline's `DepthStencilState` (D1). depth_view: wgpu::TextureView, - /// Bind group layout for the per-object uniforms (group 1) — must match every pipeline layout. - object_layout: wgpu::BindGroupLayout, /// Shared per-frame uniform buffer handle — kept so the camera matrices can be rewritten each /// frame (`render_scene`) and shipped to the GPU before the frame bind group is used. frame_buffer: wgpu::Buffer, @@ -67,10 +68,6 @@ pub struct Renderer { frame_bind_group: wgpu::BindGroup, /// Shared per-object bind group (identity model) used by the low-level `render` path. shared_object_bind_group: wgpu::BindGroup, - /// Per-entity object uniform buffers + bind groups, lazily created on first encounter and keyed by - /// entity label. Needed because `render_scene(&self, &Scene)` is immutable; the model matrix is - /// rewritten each frame for every entity. - object_cache: RefCell>, /// Flat (unlit) rendering flag, exposed via [`Renderer::set_unlit`]. When true, `options.x` of the /// `FrameUniforms` is set to 1 so the `standard` shader returns vertex colors as-is — flat 2D /// rendering is thus a special case of the 3D lit path (DRAFT Step 5). Defaults to `false` (lit). @@ -91,6 +88,36 @@ pub struct Renderer { shadow_uniform_bind_group: wgpu::BindGroup, /// Depth-only pipeline rendering the scene from the shadow light's point of view (D4). shadow_pipeline: wgpu::RenderPipeline, + // ---- Phase 3 (Step 15) — GPU-driven rendering: compute pipelines + slot buffers + bind groups ---- + /// Compute pipeline deriving per-entity world matrices from the transform buffer (Step 15.5). + compute_matrices_pipeline: wgpu::ComputePipeline, + /// Compute pipeline culling entities + filling the indirect draw args (Step 15.6). + cull_pipeline: wgpu::ComputePipeline, + /// GPU world-matrix slots (`STORAGE | COPY_DST`), written by `compute_matrices`, bound per-slot + /// via `matrix_object_bg`. + matrix_buffer: wgpu::Buffer, + /// GPU transform slots (`STORAGE | COPY_DST`), read by both compute passes; rewritten by the CPU each frame. + transform_buffer: wgpu::Buffer, + /// GPU local-space bounding boxes (`STORAGE | COPY_DST`), read by `cull`; uploaded once per mesh set. + bbox_buffer: wgpu::Buffer, + /// GPU indirect draw args (`STORAGE | INDIRECT`), written by `cull`, read by the indirect renders. + draw_args_buffer: wgpu::Buffer, + /// GPU cull uniforms (`UNIFORM | COPY_DST`): frustum planes + control flags; rewritten each frame. + cull_uniform_buffer: wgpu::Buffer, + /// `compute_matrices`/`cull` group 0 (transform buffer, storage read) — shared by both compute passes. + transform_bg: wgpu::BindGroup, + /// `compute_matrices` group 1 (matrix buffer, storage read_write). + matrices_bg: wgpu::BindGroup, + /// `cull` group 2 (cull uniforms + bboxes + draw args). + cull_bundle_bg: wgpu::BindGroup, + /// Shared object bind group (group 1, dynamic) binding ONE 64-byte matrix slice of the GPU matrix + /// buffer — every entity's render/shadow draw binds this with a per-slot dynamic offset + /// (`slot_index * MAT_SLOT_SIZE`, 256-aligned) to select its slot. + matrix_object_bg: wgpu::BindGroup, + /// Whether GPU frustum culling is enabled (off by default until validated, Step 15.6). + /// Interior-mutable so `set_culling` can toggle it from an immutable `&Renderer` (matching the + /// Renderer's all-`&self` API). Read each frame by `render_scene` when building the cull uniforms. + cull_enabled: Cell, } impl Renderer { @@ -204,17 +231,216 @@ impl Renderer { }); let shadow_pipeline = build_shadow_pipeline(&device, &object_layout); + // Phase 3 (Step 15) — GPU-driven rendering. One compute shader module with two entry points + // (`compute_matrices`, `cull`); a single explicit 3-group pipeline layout is shared by both + // pipelines so they bind the same transforms / matrices / cull buffers (DRAFT Step 15.5). + let gpu_shader = device.create_shader_module(wgpu::ShaderModuleDescriptor { + label: Some("gpu-driven compute shader"), + source: wgpu::ShaderSource::Wgsl(GPU_DRIVEN_SHADER.into()), + }); + // Group 0: transform slots (storage read). Group 1: world matrices (storage read_write). + let gpu_transforms_layout = + device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("gpu transforms layout"), + 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, + }], + }); + let gpu_matrices_layout = + device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("gpu matrices layout"), + entries: &[wgpu::BindGroupLayoutEntry { + binding: 0, + 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, + }], + }); + // Group 2: cull uniforms (uniform) + bounding boxes (storage read) + draw args (storage rw). + let gpu_cull_layout = device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { + label: Some("gpu cull layout"), + entries: &[ + wgpu::BindGroupLayoutEntry { + binding: 0, + visibility: wgpu::ShaderStages::COMPUTE, + ty: wgpu::BindingType::Buffer { + ty: wgpu::BufferBindingType::Uniform, + 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: true }, + has_dynamic_offset: false, + min_binding_size: None, + }, + count: None, + }, + wgpu::BindGroupLayoutEntry { + binding: 2, + 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 gpu_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor { + label: Some("gpu pipeline layout"), + bind_group_layouts: &[ + Some(&gpu_transforms_layout), + Some(&gpu_matrices_layout), + Some(&gpu_cull_layout), + ], + immediate_size: 0, + }); + let compute_matrices_pipeline = + device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("compute_matrices pipeline"), + layout: Some(&gpu_pipeline_layout), + module: &gpu_shader, + entry_point: Some("compute_matrices"), + compilation_options: Default::default(), + cache: None, + }); + let cull_pipeline = device.create_compute_pipeline(&wgpu::ComputePipelineDescriptor { + label: Some("cull pipeline"), + layout: Some(&gpu_pipeline_layout), + module: &gpu_shader, + entry_point: Some("cull"), + compilation_options: Default::default(), + cache: None, + }); + + // Fixed-capacity slot buffers (allocated once). Transform + cull-uniform buffers are rewritten + // by the CPU each frame; matrix + draw-args buffers are GPU-written; the bbox buffer is uploaded + // once per mesh set. + let transform_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("GPU transform slots"), + size: MAX_ENTITIES as u64 * TRANSFORM_SLOT_SIZE, + // COPY_SRC: lets `debug_dump` read the slots back via copy + map. + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let matrix_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("GPU world matrices"), + size: MAX_ENTITIES as u64 * MAT_SLOT_SIZE, + // COPY_SRC: lets `debug_dump` read the GPU-written slots back via copy + map. + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::UNIFORM + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let bbox_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("GPU bounding boxes"), + size: MAX_ENTITIES as u64 * BBOX_SLOT_SIZE, + // COPY_SRC: lets `debug_dump` read the boxes back via copy + map. + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let draw_args_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("GPU indirect draw args"), + size: MAX_ENTITIES as u64 * DRAW_SLOT_SIZE, + // COPY_SRC: lets `debug_dump` read the GPU-written args back via copy + map. + usage: wgpu::BufferUsages::STORAGE + | wgpu::BufferUsages::INDIRECT + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + let cull_uniform_buffer = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("GPU cull uniforms"), + size: CULL_UNIFORMS_SIZE, + usage: wgpu::BufferUsages::UNIFORM + | wgpu::BufferUsages::COPY_DST + | wgpu::BufferUsages::COPY_SRC, + mapped_at_creation: false, + }); + // Bind groups against the explicit layouts. `transform_bg` is shared by both compute passes + // (group 0); `matrices_bg` by `compute_matrices` (group 1); `cull_bundle_bg` by `cull` + // (group 2). `matrix_object_bg` uses the render pipeline's dynamic object layout (group 1) and + // binds a single 64-byte slice of the matrix buffer — per-entity draws select the slice's 256-byte + // slot via a dynamic offset (a whole-buffer binding would cap the dynamic offset at 0). + let transform_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("gpu transforms bind group"), + layout: &gpu_transforms_layout, + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: transform_buffer.as_entire_binding(), + }], + }); + let matrices_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("gpu matrices bind group"), + layout: &gpu_matrices_layout, + entries: &[wgpu::BindGroupEntry { + binding: 0, + resource: matrix_buffer.as_entire_binding(), + }], + }); + let cull_bundle_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("gpu cull bundle bind group"), + layout: &gpu_cull_layout, + entries: &[ + wgpu::BindGroupEntry { + binding: 0, + resource: cull_uniform_buffer.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 1, + resource: bbox_buffer.as_entire_binding(), + }, + wgpu::BindGroupEntry { + binding: 2, + resource: draw_args_buffer.as_entire_binding(), + }, + ], + }); + let matrix_object_bg = device.create_bind_group(&wgpu::BindGroupDescriptor { + label: Some("GPU matrix object bind group"), + layout: &object_layout, + entries: &[wgpu::BindGroupEntry { + binding: 0, + // Bind ONE 64-byte matrix (not the whole buffer) so the per-slot dynamic offset can + // slide across the 256-byte slots. The offset is `slot_index * MAT_SLOT_SIZE` (256-aligned). + resource: wgpu::BindingResource::Buffer(wgpu::BufferBinding { + buffer: &matrix_buffer, + offset: 0, + // One 64-byte matrix (OBJECT_UNIFORM_SIZE is a non-zero const, so the unwrap is safe). + size: Some(std::num::NonZeroU64::new(OBJECT_UNIFORM_SIZE).unwrap()), + }), + }], + }); + let renderer = Self { queue, device, format, _depth_texture: depth_texture, depth_view, - object_layout, frame_buffer, frame_bind_group, shared_object_bind_group, - object_cache: RefCell::new(HashMap::new()), unlit: false, _shadow_texture: shadow_texture, shadow_view, @@ -222,6 +448,18 @@ impl Renderer { shadow_uniform_buffer, shadow_uniform_bind_group, shadow_pipeline, + compute_matrices_pipeline, + cull_pipeline, + matrix_buffer, + transform_buffer, + bbox_buffer, + draw_args_buffer, + cull_uniform_buffer, + transform_bg, + matrices_bg, + cull_bundle_bg, + matrix_object_bg, + cull_enabled: Cell::new(false), }; // Seed the shared frame buffer with an identity camera + current unlit flag so the low-level // `render` path (which has no window/camera) sees coherent values before `render_scene` runs. @@ -352,7 +590,7 @@ impl Renderer { -light.position_dir.y, -light.position_dir.z, ), - crate::resources::LightType::Spot { .. } => { + crate::resources::LightType::Spot => { Vec3::new(light.dir_angle.x, light.dir_angle.y, light.dir_angle.z) } crate::resources::LightType::Point => return None, @@ -423,23 +661,22 @@ impl Renderer { material, &self.frame_bind_group, &self.shared_object_bind_group, + 0, // object offset 0 — the identity object buffer (low-level path, no slot). &self.shadow_bind_group, ); } self.queue.submit(std::iter::once(encoder.finish())); } - /// Renders every entity in `scene` into the given color view within a single batched render pass. - /// This avoids allocating a separate encoder and render pass per entity (which the low-level - /// `render` does), minimizing GPU submissions. Called automatically each frame by the default - /// `AppHandler::render` through `App::render_scene`. + /// Renders every entity in `scene` into the given color view, fully GPU-driven (Phase 3, Step 15). + /// Per frame the CPU rewrites only the transform slots + cull uniforms; the GPU then derives the + /// world matrices (`compute_matrices`), culls + fills the indirect draw args (`cull`), and the + /// main + shadow render passes are 100% indirect (a culled/inactive slot's args are zero → a no-op + /// draw). This removes the CPU-side per-entity loop from the render hot path. /// Inputs: view — the frame's texture view color attachment; scene — the scene whose entities are - /// drawn; aspect — the viewport aspect ratio (width/height), used to build the camera's perspective - /// projection. - /// - /// Before drawing, the shared frame uniform buffer is rewritten from `scene.camera()` so the GPU - /// receives the active camera's view/projection matrices and position for this frame (Step 4.3). + /// drawn; aspect — the viewport aspect ratio (width/height) for the camera's perspective projection. pub fn render_scene(&self, view: &wgpu::TextureView, scene: &Scene, aspect: f32) { + // 1. Rewrite the shared frame uniform buffer (camera view/proj, position, lights, shadow flags). self.write_frame_uniforms( scene.camera(), scene.lights(), @@ -448,17 +685,78 @@ impl Renderer { scene.shadow_caster(), ); + // 2. Rewrite the GPU transform slots — the single source of truth for world matrices, and the + // only per-entity CPU→GPU copy each frame (packed TRS, 64 B per slot). + let transform_slots = scene.packed_transform_slots(); + self.queue.write_buffer( + &self.transform_buffer, + 0, + bytemuck::cast_slice(&transform_slots), + ); + + // 3. Upload the local-space bounding boxes (small; the mesh set is static in practice, but + // re-uploading each frame keeps the mesh-index → bbox mapping correct if meshes are added). + let bboxes = scene.mesh_bboxes(); + self.queue + .write_buffer(&self.bbox_buffer, 0, bytemuck::cast_slice(&bboxes)); + + // 4. Compute the view frustum from the camera's view-projection and write the cull uniforms + // (six unit planes + the num_slots / culling control flags). + let camera = scene.camera(); + let view_proj = camera.projection_matrix(aspect) * camera.view_matrix(); + let frustum = Frustum::from_view_proj(&view_proj); + let cull_uniforms = + CullUniforms::from_frustum(&frustum, scene.num_slots() as u32, self.cull_enabled.get()); + self.queue.write_buffer( + &self.cull_uniform_buffer, + 0, + bytemuck::bytes_of(&cull_uniforms), + ); + let mut encoder = self .device .create_command_encoder(&wgpu::CommandEncoderDescriptor { label: Some("scene encoder"), }); - // Step 14 (DRAFT 3.2): run the depth-only shadow pass first when a light is configured to - // cast shadows (D4). It populates `shadow_view` on the shared encoder; the main pass below - // then samples it via `shadow_bind_group`. `render_shadow_map` no-ops when shadows are off. + // 5. Dispatch the two compute passes (Step 15.5): `compute_matrices` derives each entity's + // world matrix into the matrix buffer, then `cull` fills the indirect draw args (a zero + // count for a culled/inactive slot). Both read the transform buffer written above. + let workgroups = MAX_ENTITIES.div_ceil(GPU_WORKGROUP_SIZE); + { + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("compute_matrices pass"), + timestamp_writes: None, + }); + // Bind ALL THREE groups: the shared layout has a non-null bind group at every index, + // so WebGPU requires each to be set even when the entry point doesn't read it. The + // cull bundle (group 2) is unused by this entry point but must still be bound. + pass.set_pipeline(&self.compute_matrices_pipeline); + pass.set_bind_group(0, &self.transform_bg, &[]); + pass.set_bind_group(1, &self.matrices_bg, &[]); + pass.set_bind_group(2, &self.cull_bundle_bg, &[]); + pass.dispatch_workgroups(workgroups, 1, 1); + } + { + let mut pass = encoder.begin_compute_pass(&wgpu::ComputePassDescriptor { + label: Some("cull pass"), + timestamp_writes: None, + }); + // Same: bind all three. Group 1 (matrices) is unused by the cull entry point (it + // computes the world-space centre from the transform directly) but must be bound. + pass.set_pipeline(&self.cull_pipeline); + pass.set_bind_group(0, &self.transform_bg, &[]); + pass.set_bind_group(1, &self.matrices_bg, &[]); + pass.set_bind_group(2, &self.cull_bundle_bg, &[]); + pass.dispatch_workgroups(workgroups, 1, 1); + } + + // 6. Run the depth-only shadow pass (indirect) first when a light casts shadows (DRAFT 3.2/D4); + // it reads the same matrix + draw-args buffers. No-op when shadows are off. self.render_shadow_map(&mut encoder, scene); + // 7. Main render pass: one indirect draw per active slot. The matrix + draw-args are read via + // per-slot offsets; a culled/inactive slot's args are zero, so its draw is a no-op. { let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some("scene render pass"), @@ -484,42 +782,201 @@ impl Renderer { ..Default::default() }); - // Step 7 (DRAFT 7.3.4): the Material is resolved from the Mesh itself, falling back - // to the Scene's default material when the mesh carries none. - for (label, mesh, transform) in scene.iter_entities() { - let material = mesh + for slot in scene.iter_slot_draws() { + if !slot.active { + continue; // tombstone — the GPU already zeroed this slot's draw args. + } + // The Material is resolved from the Mesh itself, falling back to the Scene's default + // material when the mesh carries none. + let material = slot + .mesh .material() .cloned() .unwrap_or_else(|| scene.default_material()); - let object_bind_group = self.object_bind_group_for(label, transform); - draw_entity( - &mut render_pass, - mesh, - &material, - &self.frame_bind_group, - &object_bind_group, - &self.shadow_bind_group, - ); + let object_offset = (slot.slot_index as u64 * MAT_SLOT_SIZE) as u32; + let indirect_offset = slot.slot_index as u64 * DRAW_SLOT_SIZE; + render_pass.set_pipeline(&material.pipeline); + render_pass.set_bind_group(0, &self.frame_bind_group, &[]); + // Group 1 (dynamic): the 64-byte matrix slice for this slot. + render_pass.set_bind_group(1, &self.matrix_object_bg, &[object_offset]); + render_pass.set_bind_group(2, &material.texture_bind_group, &[]); + render_pass.set_bind_group(3, &self.shadow_bind_group, &[]); + render_pass.set_vertex_buffer(0, slot.mesh.vertex_buffer.slice(..)); + if slot.has_index { + if let Some(index_buffer) = &slot.mesh.index_buffer { + render_pass + .set_index_buffer(index_buffer.slice(..), wgpu::IndexFormat::Uint16); + } + render_pass.draw_indexed_indirect(&self.draw_args_buffer, indirect_offset); + } else { + render_pass.draw_indirect(&self.draw_args_buffer, indirect_offset); + } } } self.queue.submit(std::iter::once(encoder.finish())); } + /// Debug helper (Phase 3 black-window investigation): reads back the first `n` slots of the + /// transform / matrix / indirect-draw-args / bbox buffers (plus the cull uniforms) and prints + /// them to stderr. Call it from a render callback *after* `render_scene` so the compute + /// passes of the current frame have been submitted. + /// + /// **Synchronous on purpose** (blocks the calling thread until every staging buffer is read + /// back and unmapped). Two wgpu-core rules make this necessary (both cost a frozen render + /// loop if violated — observed empirically with this very tool): + /// + /// 1. `Queue::write_buffer` on a buffer with a **pending map** fails with + /// `TransferError::BufferNotAvailable` (and `render_scene` writes the transform / bbox / + /// cull buffers every frame). A detached-thread readback that leaves its maps pending + /// when the next frame starts corrupts/skips that frame (demo froze after 3 async dumps). + /// 2. Map callbacks are only fired by the queue `maintain`, which runs **inside** + /// `Queue::submit` — a thread blocked waiting on its own callbacks can never trigger it + /// (demo froze on the first purely-synchronous dump). + /// + /// This implementation closes both loops itself: it submits the copies, requests the maps, + /// then pumps the queue with empty submits (each one runs a `maintain` and fires whatever + /// callbacks are due) until every map has completed, and only then unmaps. At the point it + /// returns, all staging buffers are `Idle` again, so the next frame's `write_buffer` calls + /// are safe. Each dump briefly stalls the render loop (a few ms). + #[doc(hidden)] + pub fn debug_dump(&self, n: u32) { + let device = self.device.clone(); + let queue = self.queue.clone(); + let matrix_buf = self.matrix_buffer.clone(); + let draw_args_buf = self.draw_args_buffer.clone(); + let cull_buf = self.cull_uniform_buffer.clone(); + let transform_buf = self.transform_buffer.clone(); + let bbox_buf = self.bbox_buffer.clone(); + let n = n.min(MAX_ENTITIES as u32).max(1); + { + // Creates a MAP_READ staging buffer and records a copy of `size` bytes from `src` + // into the caller's encoder. NOTE: the map must only be requested AFTER the submit — + // wgpu-core rejects a submission that references a buffer in a non-idle map state. + fn readback( + device: &wgpu::Device, + src: &wgpu::Buffer, + size: u64, + enc: &mut wgpu::CommandEncoder, + ) -> wgpu::Buffer { + let read = device.create_buffer(&wgpu::BufferDescriptor { + label: Some("dbg readback"), + size, + usage: wgpu::BufferUsages::MAP_READ | wgpu::BufferUsages::COPY_DST, + mapped_at_creation: false, + }); + enc.copy_buffer_to_buffer(src, 0, &read, 0, size); + read + } + + let specs: [(&wgpu::Buffer, u64); 5] = [ + (&matrix_buf, n as u64 * MAT_SLOT_SIZE), + (&draw_args_buf, n as u64 * DRAW_SLOT_SIZE), + (&transform_buf, n as u64 * TRANSFORM_SLOT_SIZE), + (&bbox_buf, n as u64 * BBOX_SLOT_SIZE), + (&cull_buf, CULL_UNIFORMS_SIZE), + ]; + let (tx, rx) = std::sync::mpsc::channel::<()>(); + let mut reads = Vec::with_capacity(specs.len()); + { + let mut enc = device.create_command_encoder(&wgpu::CommandEncoderDescriptor { + label: Some("dbg readback encoder"), + }); + for (src, size) in &specs { + reads.push(readback(&device, src, *size, &mut enc)); + } + queue.submit([enc.finish()]); + } + // Only now request the maps (wgpu 30 has no async `map` future — one message each). + for read in &reads { + let size = read.size(); + let tx = tx.clone(); + read.map_async(wgpu::MapMode::Read, 0..size, move |_| { + let _ = tx.send(()); + }); + } + // Pump the queue until every map callback has fired. The callbacks are delivered by + // the `maintain` that runs inside `Queue::submit` — and this thread is the only one + // that will submit from now on (the caller is blocked here), so it must pump itself. + // Empty submits are cheap: each one polls the GPU and fires whatever is due. + loop { + let mut done = 0; + while rx.try_recv().is_ok() { + done += 1; + } + if done == specs.len() { + break; + } + queue.submit([]); + std::thread::sleep(std::time::Duration::from_millis(2)); + } + let data: Vec> = reads + .iter() + .map(|b| { + b.slice(..) + .get_mapped_range() + .expect("dbg mapped range") + .to_vec() + }) + .collect(); + for b in &reads { + b.unmap(); + } + let (mat_data, args_data, tr_data, bb_data, cull_data) = + (&data[0], &data[1], &data[2], &data[3], &data[4]); + for i in 0..n { + let off = (i as u64 * TRANSFORM_SLOT_SIZE) as usize; + let t: TransformSlot = + bytemuck::pod_read_unaligned(&tr_data[off..off + TRANSFORM_SLOT_SIZE as usize]); + eprintln!( + "[dbg] transform[{i}] translation={:?} flags={:?} rotation={:?} scale={:?}", + t.translation, t.flags, t.rotation, t.scale + ); + } + for i in 0..n { + let off = (i as u64 * BBOX_SLOT_SIZE) as usize; + let b: BBoxSlot = + bytemuck::pod_read_unaligned(&bb_data[off..off + BBOX_SLOT_SIZE as usize]); + eprintln!("[dbg] bbox[{i}] min={:?} max={:?}", b.min, b.max); + } + for i in 0..n { + let off = (i as u64 * MAT_SLOT_SIZE) as usize; + let m: MatSlot = + bytemuck::pod_read_unaligned(&mat_data[off..off + MAT_SLOT_SIZE as usize]); + eprintln!("[dbg] matrix[{i}] m = {:?}", m.m); + } + for i in 0..n { + let off = (i as u64 * DRAW_SLOT_SIZE) as usize; + let d: DrawSlot = + bytemuck::pod_read_unaligned(&args_data[off..off + DRAW_SLOT_SIZE as usize]); + eprintln!("[dbg] draw_args[{i}].a = {:?}", d.a); + } + let c: CullUniforms = bytemuck::pod_read_unaligned(&cull_data[..]); + eprintln!( + "[dbg] cull: num_slots={} culling={}", + c.num_slots, c.culling + ); + for (i, p) in c.planes.iter().enumerate() { + eprintln!("[dbg] plane[{i}] = {p:?}"); + } + eprintln!("[dbg] done"); + } + } + /// Renders every entity of `scene` from the shadow-casting light's point of view into the /// shadow depth map (Step 14, D4), using the dedicated depth-only `shadow_pipeline`. Called at /// the start of `render_scene`. No-ops (produces no GPU work) when `scene.shadow_caster()` is - /// `None`. The shadow light's `view_proj` is written to `shadow_uniform_buffer`, and the shadow - /// pass writes depth into `shadow_view` (clear 1.0, store). The per-entity model bind groups are - /// reused from `object_bind_group_for`, so transforms match the main pass exactly. - /// Inputs: encoder (the shared command encoder for the frame), scene (entities to cast). + /// `None`. Phase 3 (Step 15): like the main pass, the shadow pass is 100% indirect — it reads the + /// GPU world matrices (group 1, dynamic offset) and the GPU indirect draw args, so transforms and + /// culling match the main pass exactly. Inputs: encoder (the shared command encoder for the frame, + /// with the compute passes already dispatched), scene (entities to cast). fn render_shadow_map(&self, encoder: &mut wgpu::CommandEncoder, scene: &Scene) { - let _caster = match scene.shadow_caster() { + let caster = match scene.shadow_caster() { Some(c) => c, None => return, }; // Recompute the light's view_proj and write it into the shadow uniform buffer so the // depth-only vertex shader transforms vertices into light-clip space (D4). - let (light_index, vp) = match self.shadow_light_view_proj(scene.lights(), Some(_caster)) { + let (_light_index, vp) = match self.shadow_light_view_proj(scene.lights(), Some(caster)) { Some(pair) => pair, None => return, }; @@ -548,20 +1005,26 @@ impl Renderer { pass.set_pipeline(&self.shadow_pipeline); // Group 0: the shadow light view_proj (D4) — the shadow pipeline's only uniform group. pass.set_bind_group(0, &self.shadow_uniform_bind_group, &[]); - for (label, mesh, transform) in scene.iter_entities() { - let object_bind_group = self.object_bind_group_for(label, transform); - // Group 1: per-entity model. The shadow pipeline has no texture/sampler groups. - pass.set_bind_group(1, &object_bind_group, &[]); - pass.set_vertex_buffer(0, mesh.vertex_buffer.slice(..)); - if let Some(index_buffer) = &mesh.index_buffer { - pass.set_index_buffer(index_buffer.slice(..), wgpu::IndexFormat::Uint16); - pass.draw_indexed(0..mesh.num_indices, 0, 0..1); + // Phase 3: one indirect draw per active slot; group 1 (dynamic) selects the matrix slice and + // the draw-args offset. The shadow pipeline has no texture/sampler groups. + for slot in scene.iter_slot_draws() { + if !slot.active { + continue; + } + let object_offset = (slot.slot_index as u64 * MAT_SLOT_SIZE) as u32; + let indirect_offset = slot.slot_index as u64 * DRAW_SLOT_SIZE; + pass.set_bind_group(1, &self.matrix_object_bg, &[object_offset]); + pass.set_vertex_buffer(0, slot.mesh.vertex_buffer.slice(..)); + if slot.has_index { + if let Some(index_buffer) = &slot.mesh.index_buffer { + pass.set_index_buffer(index_buffer.slice(..), wgpu::IndexFormat::Uint16); + } + pass.draw_indexed_indirect(&self.draw_args_buffer, indirect_offset); } else { - pass.draw(0..mesh.num_vertices, 0..1); + pass.draw_indirect(&self.draw_args_buffer, indirect_offset); } } drop(pass); - let _ = light_index; // (index retained for future per-light shadow options) } /// Presents the rendered frame by submitting the acquired surface texture to the GPU queue. @@ -583,37 +1046,13 @@ impl Renderer { self.format } - /// Returns the per-entity object bind group for `label`, creating its uniform buffer on first - /// encounter and rewriting the model matrix each call. Since `render_scene(&self)` is immutable, - /// the lazily-populated cache is interior-mutable (`RefCell`). Step 4.2. - /// Inputs: label (entity identifier used as cache key), transform (world placement to upload). - /// Returns an owned (cheaply Arc-cloned) reference handle to the object bind group (group 1). - fn object_bind_group_for(&self, label: &str, transform: &Transform) -> wgpu::BindGroup { - let mut cache = self.object_cache.borrow_mut(); - let (buffer, bind_group) = cache.entry(label.to_string()).or_insert_with(|| { - let buffer = self.device.create_buffer(&wgpu::BufferDescriptor { - label: Some("object uniform buffer"), - size: OBJECT_UNIFORM_SIZE, - usage: wgpu::BufferUsages::UNIFORM | wgpu::BufferUsages::COPY_DST, - mapped_at_creation: false, - }); - let bind_group = self.device.create_bind_group(&wgpu::BindGroupDescriptor { - label: Some("object bind group"), - layout: &self.object_layout, - entries: &[wgpu::BindGroupEntry { - binding: 0, - resource: buffer.as_entire_binding(), - }], - }); - (buffer, bind_group) - }); - // Rewrite the model matrix every frame so entity transforms can update (e.g. rotation). - let object_uniforms = ObjectUniform { - model: transform.to_matrix(), - }; - self.queue - .write_buffer(buffer, 0, bytemuck::bytes_of(&object_uniforms)); - bind_group.clone() + /// Enables or disables GPU frustum culling (Phase 3, Step 15.6). Culling is off by default: with + /// it disabled, the `cull` pass copies every active slot's draw count (nothing is culled), so the + /// scene renders identically to the pre-Phase-3 CPU loop. Enabling it makes `cull` test each + /// slot's world-space bounding box against the camera frustum and zero the draw args of culled + /// slots (so their indirect draws become no-ops). Inputs: enabled (true = cull, false = draw all). + pub fn set_culling(&self, enabled: bool) { + self.cull_enabled.set(enabled); } } @@ -691,6 +1130,7 @@ fn draw_entity( material: &Material, frame_bind_group: &wgpu::BindGroup, object_bind_group: &wgpu::BindGroup, + object_offset: u32, shadow_bind_group: &wgpu::BindGroup, ) { if mesh.num_vertices == 0 { @@ -699,7 +1139,10 @@ fn draw_entity( } pass.set_pipeline(&material.pipeline); pass.set_bind_group(0, frame_bind_group, &[]); - pass.set_bind_group(1, object_bind_group, &[]); + // Phase 3 (D12): the object (model) binding is dynamic — `object_offset` selects the 64-byte + // slice. The low-level path passes 0 (the shared identity buffer); the GPU-driven path uses its + // own inline `set_bind_group` calls (not this helper) with a per-slot `slot_index * MAT_SLOT_SIZE`. + pass.set_bind_group(1, object_bind_group, &[object_offset]); // Step 10 (DRAFT 10.4): texture group — the Material owns its bind group (placeholder // white if it has no texture, D1/D2). Always bindable since it is attached to every pipeline. pass.set_bind_group(2, &material.texture_bind_group, &[]); diff --git a/lib/src/math/frustum.rs b/lib/src/math/frustum.rs new file mode 100644 index 0000000..555f12b --- /dev/null +++ b/lib/src/math/frustum.rs @@ -0,0 +1,178 @@ +//! # Frustum Module +//! +//! View-projection frustum representation and plane extraction, for frustum culling (Phase 3, +//! Step 15.6). Planes follow the Gribb-Hartmann convention, adapted to WebGPU's `[0, 1]` clip-space +//! z range (the `directx` projection produced by [`crate::resources::Camera::projection_matrix`]). +//! +//! Each plane is a `[f32; 4]` `(normal, d)` such that a world point `p` is **inside** the frustum +//! iff `dot(p, normal) + d >= 0` for every plane. The six planes are extracted from the rows of the +//! view-projection matrix `M` (world to clip space), whose NDC conventions are x, y in `[-1, 1]` and +//! z in `[0, 1]`: +//! +//! | plane | clip-space inequality | row combination | +//! |--------|-----------------------|-----------------| +//! | left | cx + cw >= 0 | w + x | +//! | right | -cx + cw >= 0 | w - x | +//! | bottom | cy + cw >= 0 | w + y | +//! | top | -cy + cw >= 0 | w - y | +//! | near | cz >= 0 | z | +//! | far | -cz + cw >= 0 | w - z | +//! +//! (For the `[0, 1]` z range the near plane is the z row alone — `cz >= 0` — whereas the classic +//! `[-1, 1]` Gribb-Hartmann uses `w + z`. The far plane `w - z` is the same in both.) +//! Each plane is normalized to a unit normal so the signed-distance test is scale-invariant. + +use glam::{Mat4, Vec3, Vec4}; + +/// A view-projection frustum represented by its six bounding planes. +/// +/// Each plane is a `[f32; 4]` `(normal, d)`: a world point `p` is inside when +/// `dot(p, normal) + d >= 0`. Built from a view-projection matrix via +/// [`Frustum::from_view_proj`] and uploaded to the GPU culling compute shader (Phase 3). +#[derive(Debug, Clone, Copy)] +pub struct Frustum { + /// The six frustum planes, order: `[left, right, bottom, top, near, far]`. + pub planes: [[f32; 4]; 6], +} + +impl Frustum { + /// Extracts the six frustum planes from a view-projection matrix (Gribb-Hartmann, adapted to + /// WebGPU's `[0, 1]` clip-space z). Inputs: m — the `projection * view` matrix (world to clip + /// space). Returns the frustum with unit-length plane normals. + pub fn from_view_proj(m: &Mat4) -> Self { + // glam stores Mat4 by column; transpose so `.x_axis`/`.y_axis`/... are the rows of M, + // i.e. the clip-space basis vectors the Gribb-Hartmann method combines. + let mt = m.transpose(); + let r0 = mt.x_axis; // row 0 of M -> clip x + let r1 = mt.y_axis; // row 1 of M -> clip y + let r2 = mt.z_axis; // row 2 of M -> clip z + let r3 = mt.w_axis; // row 3 of M -> clip w + let raw: [Vec4; 6] = [ + r3 + r0, // left + r3 - r0, // right + r3 + r1, // bottom + r3 - r1, // top + r2, // near + r3 - r2, // far + ]; + let planes = raw.map(|p| { + let n = Vec3::new(p.x, p.y, p.z); + let len = n.length(); + if len > 1e-8 { + let nn = n / len; + [nn.x, nn.y, nn.z, p.w / len] + } else { + [0.0, 0.0, 0.0, 0.0] + } + }); + Self { planes } + } + + /// Tests whether a world-space point lies inside the frustum (inside every plane). + /// Inputs: p — a world-space point. Returns true if it satisfies all six plane inequalities. + pub fn contains_point(&self, p: Vec3) -> bool { + self.planes + .iter() + .all(|plane| plane[0] * p.x + plane[1] * p.y + plane[2] * p.z + plane[3] >= 0.0) + } +} + +#[cfg(test)] +mod tests { + use super::*; + use crate::resources::camera::Camera; + + /// Builds the view-projection matrix for a camera at `(0,0,d)` looking at the origin (45 deg fov, + /// near 0.1, far 100), matching the `directx` (WebGPU `[0,1]`) projection used by the renderer. + fn vp(d: f32) -> Mat4 { + let cam = Camera::new(Vec3::new(0.0, 0.0, d), Vec3::ZERO, Vec3::Y).with_perspective( + 45.0_f32.to_radians(), + 0.1, + 100.0, + ); + cam.projection_matrix(1.0) * cam.view_matrix() + } + + #[test] + fn origin_inside_when_camera_looks_at_it() { + let fr = Frustum::from_view_proj(&vp(10.0)); + assert!( + fr.contains_point(Vec3::new(0.0, 0.0, 0.0)), + "the look-at target must be inside the frustum" + ); + } + + #[test] + fn behind_camera_is_culled() { + let fr = Frustum::from_view_proj(&vp(10.0)); + // Camera at z = 10 looks toward -z; a point at z = 50 is behind it. + assert!( + !fr.contains_point(Vec3::new(0.0, 0.0, 50.0)), + "a point behind the camera must be culled" + ); + } + + #[test] + fn far_to_the_side_is_culled() { + let fr = Frustum::from_view_proj(&vp(10.0)); + // Far off to the side, well outside the 45-degree field of view. + assert!( + !fr.contains_point(Vec3::new(1000.0, 0.0, 0.0)), + "a point far to the side must be culled" + ); + } + + #[test] + fn beyond_far_plane_is_culled() { + let fr = Frustum::from_view_proj(&vp(10.0)); + // z = -500 is 510 units in front of the camera (at z = 10), beyond far = 100. + assert!( + !fr.contains_point(Vec3::new(0.0, 0.0, -500.0)), + "a point beyond the far plane must be culled" + ); + } + + #[test] + fn planes_are_unit_length() { + let fr = Frustum::from_view_proj(&vp(10.0)); + for plane in fr.planes { + let len = (plane[0] * plane[0] + plane[1] * plane[1] + plane[2] * plane[2]).sqrt(); + assert!( + (len - 1.0).abs() < 1e-3, + "plane normal must be unit length, got {len}" + ); + } + } + + /// Reproduces the `demo` example's exact camera + entity layout and confirms the GPU cull + /// pass would NOT cull any of them (they sit at radius 1.7 around the origin, in front of the + /// orbital camera). If this fails, the demo's black window is a frustum-culling bug. + #[test] + fn demo_camera_sees_all_primitives() { + use crate::resources::CameraController; + let mut ctrl = CameraController::default(); + ctrl.yaw = 0.6; + ctrl.pitch = 0.35; + ctrl.distance = 6.5; + ctrl.target = Vec3::ZERO; + let mut cam = Camera::default(); + ctrl.apply_to(&mut cam); + let vp = cam.projection_matrix(1.0) * cam.view_matrix(); + let fr = Frustum::from_view_proj(&vp); + // Ground plane center (origin). + assert!( + fr.contains_point(Vec3::ZERO), + "origin (ground center) must be inside" + ); + // The six primitives, placed by demo::place at radius 1.7, y = 0.5. + for i in 0..6 { + let a = i as f32 / 6.0 * std::f32::consts::TAU; + let p = Vec3::new(a.cos() * 1.7, 0.5, a.sin() * 1.7); + assert!( + fr.contains_point(p), + "primitive {i} at {} must be inside the frustum", + p + ); + } + } +} diff --git a/lib/src/math/geometry.rs b/lib/src/math/geometry.rs index 1b205e0..bcff61a 100644 --- a/lib/src/math/geometry.rs +++ b/lib/src/math/geometry.rs @@ -22,6 +22,57 @@ use crate::resources::Vertex; +/// An axis-aligned bounding box in object (local) space: the min/max corners of a geometry's +/// positions. Used for conservative sphere culling (Phase 3): the culling radius is the box's +/// circumradius and the culling center is its center, both computed once per mesh and uploaded +/// to the GPU culling buffer. +#[derive(Debug, Clone, Copy, PartialEq, bytemuck::Pod, bytemuck::Zeroable)] +#[repr(C)] +pub struct BBox { + /// Minimum corner (local space). + pub min: [f32; 3], + /// Maximum corner (local space). + pub max: [f32; 3], +} + +impl BBox { + /// Builds a bounding box from a list of local-space positions. Returns `None` when the input + /// is empty (a geometry must have at least one position to have a bounding box). + pub fn from_positions(positions: &[[f32; 3]]) -> Option { + if positions.is_empty() { + return None; + } + let mut min = [f32::INFINITY; 3]; + let mut max = [f32::NEG_INFINITY; 3]; + for p in positions { + for i in 0..3 { + min[i] = min[i].min(p[i]); + max[i] = max[i].max(p[i]); + } + } + Some(Self { min, max }) + } + + /// The center of the box (local space), i.e. the midpoint of the min and max corners. + pub fn center(&self) -> [f32; 3] { + [ + (self.min[0] + self.max[0]) * 0.5, + (self.min[1] + self.max[1]) * 0.5, + (self.min[2] + self.max[2]) * 0.5, + ] + } + + /// The circumradius: distance from the center to a corner (all corners are equidistant for an + /// axis-aligned box). This is the conservative culling radius — the sphere of this radius + /// around the center fully contains the box. + pub fn circumradius(&self) -> f32 { + let dx = (self.max[0] - self.min[0]) * 0.5; + let dy = (self.max[1] - self.min[1]) * 0.5; + let dz = (self.max[2] - self.min[2]) * 0.5; + (dx * dx + dy * dy + dz * dz).sqrt() + } +} + /// Validation error produced when a `Geometry` is inconsistent, i.e. its optional /// per-vertex arrays (`normals`, `uvs`, `colors`) have a length different from /// `positions`, or an index is out of bounds. @@ -223,6 +274,13 @@ impl Geometry { self.indices.as_deref() } + /// Computes the axis-aligned bounding box of this geometry's positions (local space). + /// Returns `None` when the geometry has no positions. Used by `Mesh`/`Scene` to build the + /// per-mesh culling sphere (Phase 3). + pub fn bbox(&self) -> Option { + BBox::from_positions(&self.positions) + } + /// Returns the CPU vertices in interleaved `resources::Vertex` layout, suitable /// for GPU upload. Assumes the geometry is valid; missing per-vertex attributes /// are filled with defaults: @@ -307,6 +365,26 @@ mod tests { assert_eq!(geo.validate(), Err(GeometryError::EmptyPositions)); } + #[test] + fn bbox_of_quad() { + let bb = quad().bbox().expect("quad has positions"); + assert_eq!(bb.min, [0.0, 0.0, 0.0]); + assert_eq!(bb.max, [1.0, 1.0, 0.0]); + assert_eq!(bb.center(), [0.5, 0.5, 0.0]); + // Circumradius of a unit square: half-diagonal = sqrt(0.5^2 + 0.5^2) = sqrt(0.5). + let expected = (0.5_f32 * 0.5 + 0.5_f32 * 0.5).sqrt(); + assert!( + (bb.circumradius() - expected).abs() < 1e-6, + "got {}", + bb.circumradius() + ); + } + + #[test] + fn bbox_of_empty_is_none() { + assert!(Geometry::new(Vec::new()).bbox().is_none()); + } + #[test] fn validate_rejects_attribute_count_mismatch() { let geo = Geometry::new(vec![[0.0, 0.0, 0.0], [1.0, 0.0, 0.0]]) diff --git a/lib/src/math/mod.rs b/lib/src/math/mod.rs index 4b14be0..6ec8845 100644 --- a/lib/src/math/mod.rs +++ b/lib/src/math/mod.rs @@ -16,11 +16,13 @@ //! - `geometry.rs`: Defines the `Geometry` struct for mesh data storage //! - `primitives.rs`: Procedural mesh generators (cube, sphere, cylinder, cone, torus…) returning `Geometry` +pub mod frustum; pub mod geometry; pub mod primitives; pub mod transform; // Re-exports -pub use geometry::{Geometry, GeometryError}; +pub use frustum::Frustum; +pub use geometry::{BBox, Geometry, GeometryError}; pub use primitives::{cone, cube, cylinder, icosphere, plane, torus, uv_sphere}; pub use transform::Transform; diff --git a/lib/src/pipeline/pipeline_cache.rs b/lib/src/pipeline/pipeline_cache.rs index 655e2f9..746152f 100644 --- a/lib/src/pipeline/pipeline_cache.rs +++ b/lib/src/pipeline/pipeline_cache.rs @@ -22,12 +22,15 @@ use std::collections::HashMap; use std::sync::Arc; /// Creates the two bind group layouts shared by **every** pipeline (Step 3 — decision ratified -/// "a single layout for all"). Both buffers are `Uniform`, 16-byte aligned, no dynamic offset. -/// Matching CPU types: `FrameUniforms` (192 B) and `ObjectUniform` (64 B) in `resources::uniform`. +/// "a single layout for all"). Both buffers are `Uniform` and 16-byte aligned. Matching CPU +/// types: `FrameUniforms` (192 B) and `ObjectUniform` (64 B) in `resources::uniform`. /// Returns `[frame_layout, object_layout]` in renderer binding order. /// -/// - `index 0`: per-frame uniforms (view/proj/light/options), visible in both shader stages. -/// - `index 1`: per-object uniforms (model matrix), visible in the vertex stage only. +/// - `index 0`: per-frame uniforms (view/proj/light/options), visible in both shader stages. Static +/// (one shared `FrameUniforms` buffer per frame, no dynamic offset). +/// - `index 1`: per-object uniforms (model matrix), visible in the vertex stage only. **Dynamic** +/// (Phase 3, D12): the offset selects a 64-byte slice of the single GPU-written matrix buffer, +/// so every entity shares one buffer. The low-level `render` path passes offset 0. pub fn create_uniform_bind_group_layouts(device: &wgpu::Device) -> [wgpu::BindGroupLayout; 2] { [ device.create_bind_group_layout(&wgpu::BindGroupLayoutDescriptor { @@ -50,7 +53,10 @@ pub fn create_uniform_bind_group_layouts(device: &wgpu::Device) -> [wgpu::BindGr visibility: wgpu::ShaderStages::VERTEX, ty: wgpu::BindingType::Buffer { ty: wgpu::BufferBindingType::Uniform, - has_dynamic_offset: false, + // Phase 3 (D12): dynamic offset so every entity shares the single GPU-written + // matrix buffer (one 64-byte slice per slot), instead of a per-entity buffer. + // The low-level `render` path passes offset 0 (its identity object buffer). + has_dynamic_offset: true, min_binding_size: None, }, count: None, diff --git a/lib/src/resources/mod.rs b/lib/src/resources/mod.rs index 52f9b23..ecbd899 100644 --- a/lib/src/resources/mod.rs +++ b/lib/src/resources/mod.rs @@ -27,8 +27,10 @@ pub use material::Material; pub use mesh::Mesh; pub use texture::{Texture, TextureError}; pub use uniform::{ - FRAME_UNIFORMS_SIZE, FrameUniforms, Light, LightType, MAX_LIGHTS, OBJECT_UNIFORM_SIZE, - ObjectUniform, SHADOW_UNIFORM_SIZE, ShadowUniform, + BBOX_SLOT_SIZE, BBoxSlot, CULL_UNIFORMS_SIZE, CullUniforms, DRAW_SLOT_SIZE, DrawSlot, + FRAME_UNIFORMS_SIZE, FrameUniforms, Light, LightType, MAT_SLOT_SIZE, MAX_LIGHTS, MatSlot, + OBJECT_UNIFORM_SIZE, ObjectUniform, SHADOW_UNIFORM_SIZE, ShadowUniform, TRANSFORM_SLOT_SIZE, + TransformSlot, }; pub use vertex::Vertex; diff --git a/lib/src/resources/uniform.rs b/lib/src/resources/uniform.rs index bb53629..f74ce13 100644 --- a/lib/src/resources/uniform.rs +++ b/lib/src/resources/uniform.rs @@ -178,6 +178,211 @@ pub struct ShadowUniform { pub view_proj: Mat4, } +// ============================================================================ +// Phase 3 — GPU-driven entity slot buffers (Step 15) +// ============================================================================ +// +// CPU-side `Pod` mirrors of the GPU buffer element types declared in `gpu_driven.wgsl` (see that +// file's "GPU buffer layouts" note). Their byte layout must match the WGSL structs **exactly** — +// including the 16-byte alignment of `vec3` (std140), which is why `vec3` fields below carry an +// explicit `_pad` so the Rust offsets line up with the WGSL ones. +// +// All buffers are fixed-capacity (`MAX_ENTITIES`), allocated once. Per frame the CPU rewrites only +// the transform slots + cull uniforms; the matrices and indirect draw args are written by the GPU. + +/// Byte size of one GPU entity transform slot (`TransformSlot`). +pub const TRANSFORM_SLOT_SIZE: u64 = std::mem::size_of::() as u64; +/// Byte size of one GPU world-matrix slot (`MatSlot`). +pub const MAT_SLOT_SIZE: u64 = std::mem::size_of::() as u64; +/// Byte size of one GPU local-space bounding box (`BBoxSlot`). +pub const BBOX_SLOT_SIZE: u64 = std::mem::size_of::() as u64; +/// Byte size of one GPU indirect draw-args slot (`DrawSlot`). +pub const DRAW_SLOT_SIZE: u64 = std::mem::size_of::() as u64; +/// Byte size of the GPU cull/uniform block (`CullUniforms`). +pub const CULL_UNIFORMS_SIZE: u64 = std::mem::size_of::() as u64; + +/// A packed entity transform slot (64 bytes) — the single CPU→GPU source of truth for world +/// matrices (Step 15, D13). Mirrors the WGSL `TransformSlot`. +/// +/// Layout (std140, 16-byte aligned): translation (vec3 @0) + pad, flags (vec4 @16), rotation +/// (vec4 @32), scale (vec3 @48) + pad → 64 bytes. +/// +/// `flags` packing: x = mesh index (stable index into the mesh list), y = active (0/1), +/// z = draw count (vertex or index count for this mesh), w = has_index (0/1). +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable)] +pub struct TransformSlot { + /// World translation (xyz). Offset 0. + pub translation: [f32; 3], + _pad0: [f32; 1], + /// Packed flags: x = mesh index, y = active, z = draw count, w = has_index. Offset 16. + pub flags: [f32; 4], + /// Rotation quaternion (x, y, z, w). Offset 32. + pub rotation: [f32; 4], + /// Non-uniform scale (xyz). Offset 48. + pub scale: [f32; 3], + _pad1: [f32; 1], +} + +impl TransformSlot { + /// Builds an active transform slot from a CPU [`crate::math::Transform`] + the mesh's draw + /// metadata. `mesh_index` / `draw_count` are packed into `flags`; `active` is 1. + pub fn from_transform( + t: &crate::math::Transform, + mesh_index: u32, + draw_count: u32, + has_index: bool, + ) -> Self { + Self { + translation: t.translation.to_array(), + _pad0: [0.0; 1], + flags: [ + mesh_index as f32, + 1.0, + draw_count as f32, + if has_index { 1.0 } else { 0.0 }, + ], + rotation: [t.rotation.x, t.rotation.y, t.rotation.z, t.rotation.w], + scale: t.scale.to_array(), + _pad1: [0.0; 1], + } + } + + /// Builds an inactive (tombstone) slot: `active` = 0; the GPU writes identity and skips the draw. + pub fn inactive() -> Self { + Self { + translation: [0.0; 3], + _pad0: [0.0; 1], + flags: [0.0, 0.0, 0.0, 0.0], + rotation: [0.0, 0.0, 0.0, 1.0], + scale: [0.0; 3], + _pad1: [0.0; 1], + } + } + + /// Stable index of the entity's mesh (from `flags.x`). + pub fn mesh_index(&self) -> u32 { + self.flags[0] as u32 + } + + /// Draw count (vertex or index count) packed in `flags.z`. + pub fn draw_count(&self) -> u32 { + self.flags[2] as u32 + } + + /// Whether the entity's mesh is indexed (from `flags.w`). + pub fn has_index(&self) -> bool { + self.flags[3] >= 0.5 + } + + /// Whether the slot is active (from `flags.y`). + pub fn is_active(&self) -> bool { + self.flags[1] >= 0.5 + } +} + +/// A 4x4 world matrix (64 B) followed by 192 B of padding = **256 B** total. The padding is +/// REQUIRED: the render pipelines read this slot through the `uniform` object group with a +/// per-slot dynamic offset, and WebGPU demands that offset be a multiple of +/// `min_uniform_buffer_offset_alignment` (256 B). A bare 64-byte matrix can never be individually +/// addressable via a uniform offset, so each slot is padded to a 256-byte boundary (capacity is +/// capped at 256 = 64 KB / 256 B). Derived on the GPU by `compute_matrices` (Step 15). Mirrors the +/// WGSL `MatSlot`. (No `Default`: the matrices are GPU-written, so the CPU never constructs a +/// `MatSlot` — this type exists only to fix the buffer's slot size. `pad` is `[f32; 48]`, beyond +/// the `N ≤ 32` bound of the array `Default` impl, so `Default` cannot be derived.) +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable)] +pub struct MatSlot { + /// The world matrix (object → world), column-major. Offset 0. + pub m: Mat4, + /// 192 bytes of padding (alignment only, never read). Offset 64. + pub pad: [f32; 48], +} + +/// A local-space axis-aligned bounding box (32 bytes), uploaded once per mesh (Step 15). +/// Mirrors the WGSL `BBoxSlot` (min vec3 @0 + pad, max vec3 @16 + pad). +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable)] +pub struct BBoxSlot { + /// Bounding box minimum corner (xyz). Offset 0. + pub min: [f32; 3], + _pad0: [f32; 1], + /// Bounding box maximum corner (xyz). Offset 16. + pub max: [f32; 3], + _pad1: [f32; 1], +} + +impl BBoxSlot { + /// Builds a slot from a CPU [`crate::math::BBox`] (padding zeroed). + pub fn from_bbox(b: &crate::math::BBox) -> Self { + Self { + min: b.min, + max: b.max, + _pad0: [0.0; 1], + _pad1: [0.0; 1], + } + } + + /// A degenerate (all-zero) box, used as the placeholder for a mesh with no bounding box. + pub fn empty() -> Self { + Self { + min: [0.0; 3], + max: [0.0; 3], + _pad0: [0.0; 1], + _pad1: [0.0; 1], + } + } +} + +/// Indirect draw arguments for one entity (80 bytes = five u32 vec4s). The shader writes only `.a`; +/// the rest stays zero (the constant instance count of 1 lives in `.a.y`). Mirrors the WGSL `DrawSlot`. +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable, Default)] +pub struct DrawSlot { + /// `.a` = (count, instance_count, first_vertex/first_index, base_vertex); `.b..e` = 0. + pub a: [u32; 4], + /// Reserved (first_instance for the indexed block). Kept zero. + pub b: [u32; 4], + /// Reserved padding; kept zero (part of the 80-byte slot). + pub c: [u32; 4], + /// Reserved padding; kept zero (part of the 80-byte slot). + pub d: [u32; 4], + /// Reserved padding; kept zero (part of the 80-byte slot). + pub e: [u32; 4], +} + +/// Per-frame GPU cull/uniform block (112 bytes), rewritten by the CPU each frame (Step 15). +/// Mirrors the WGSL `CullUniforms`. +#[repr(C)] +#[derive(Clone, Copy, Pod, Zeroable)] +pub struct CullUniforms { + /// Six frustum planes, unit (normal, d); inside the frustum iff dot(p, normal) + d >= 0. + pub planes: [[f32; 4]; 6], + /// Number of live entity slots (slots at/ beyond this are no-ops). + pub num_slots: u32, + /// 0 = culling disabled (every active entity draws), 1 = enabled (sphere test). + pub culling: u32, + _pad: [u32; 2], +} + +impl CullUniforms { + /// Builds the cull block from six frustum planes (each a unit `[normal; d]` `[f32; 4]`) + the + /// control flags. `num_slots` = number of live entity slots. + pub fn new(planes: [[f32; 4]; 6], num_slots: u32, culling: bool) -> Self { + Self { + planes, + num_slots, + culling: if culling { 1 } else { 0 }, + _pad: [0; 2], + } + } + + /// Builds the cull block directly from a [`crate::math::Frustum`] (its six unit planes). + pub fn from_frustum(f: &crate::math::Frustum, num_slots: u32, culling: bool) -> Self { + Self::new(f.planes, num_slots, culling) + } +} + #[cfg(test)] mod tests { use super::*; @@ -243,4 +448,35 @@ mod tests { assert_eq!(align_of::(), 16); assert_eq!(offset_of!(ObjectUniform, model), 0); } + + #[test] + fn gpu_slot_layouts_match_wgsl() { + // TransformSlot: translation @0 (vec3+pad), flags @16, rotation @32, scale @48 (vec3+pad) -> 64. + assert_eq!(size_of::(), 64); + assert_eq!(offset_of!(TransformSlot, translation), 0); + assert_eq!(offset_of!(TransformSlot, flags), 16); + assert_eq!(offset_of!(TransformSlot, rotation), 32); + assert_eq!(offset_of!(TransformSlot, scale), 48); + + // MatSlot: one 4x4 column-major matrix (64 B) + 192 B padding -> 256 (padded to the + // uniform offset alignment; see the struct doc). m @0, pad @64. + assert_eq!(size_of::(), 256); + assert_eq!(align_of::(), 16); + assert_eq!(offset_of!(MatSlot, m), 0); + assert_eq!(offset_of!(MatSlot, pad), 64); + + // BBoxSlot: min @0 (vec3+pad), max @16 (vec3+pad) -> 32. + assert_eq!(size_of::(), 32); + assert_eq!(offset_of!(BBoxSlot, min), 0); + assert_eq!(offset_of!(BBoxSlot, max), 16); + + // DrawSlot: five u32 vec4s -> 80 (a multiple of both 16 and 20, per the WebGPU indirect rule). + assert_eq!(size_of::(), 80); + + // CullUniforms: 6 planes (96) + num_slots @96 + culling @100 + pad -> 112. + assert_eq!(size_of::(), 112); + assert_eq!(offset_of!(CullUniforms, planes), 0); + assert_eq!(offset_of!(CullUniforms, num_slots), 96); + assert_eq!(offset_of!(CullUniforms, culling), 100); + } } diff --git a/lib/src/scene/mod.rs b/lib/src/scene/mod.rs index f9ba1bf..2d133f0 100644 --- a/lib/src/scene/mod.rs +++ b/lib/src/scene/mod.rs @@ -21,4 +21,4 @@ pub mod scene; // Re-export pub use entity::Entity; -pub use scene::Scene; +pub use scene::{Scene, SlotDraw}; diff --git a/lib/src/scene/scene.rs b/lib/src/scene/scene.rs index 116c83d..8307a3e 100644 --- a/lib/src/scene/scene.rs +++ b/lib/src/scene/scene.rs @@ -17,7 +17,7 @@ use crate::math::{Geometry, Transform}; use crate::pipeline::PipelineCache; -use crate::resources::{Camera, Lights, Material, Mesh, Texture}; +use crate::resources::{BBoxSlot, Camera, Lights, Material, Mesh, Texture, TransformSlot}; use crate::scene::Entity; use glam::Vec3; use std::cell::RefCell; @@ -37,6 +37,42 @@ struct SceneGpu { cache: RefCell, } +/// A stable, append-only slot for an entity in the GPU-driven slot buffers (Phase 3, Step 15). +/// Slots are **never freed**: removing an entity leaves a tombstone (its label drops out of the +/// `entities` map) so that slot indices stay stable across frames and the fixed-capacity GPU buffers +/// can be indexed by a constant slot index. The transform itself is read from the `entities` map +/// (by `label`) at pack time; the slot only carries the mesh identity + draw metadata, which are +/// stable once the entity is (re-)added. +#[derive(Clone)] +struct EntitySlot { + /// Entity label (key into the `entities` map; the transform is read from there each frame). + label: String, + /// Stable index of the entity's mesh (position in `mesh_order`) — indexes the GPU bbox buffer. + mesh_index: u32, + /// Draw count for the entity's mesh (vertex count, or index count when indexed) → packed into + /// the transform slot's `flags.z`. + draw_count: u32, + /// Whether the entity's mesh is indexed → packed into the transform slot's `flags.w`. + has_index: bool, +} + +/// Per-slot draw descriptor for the GPU-driven render loop (Phase 3). Carries everything the +/// renderer needs to issue one indirect draw: the slot index (→ indirect-args + matrix buffer +/// offset), whether the slot is active (tombstones are skipped on the CPU), the mesh, and whether +/// it is indexed. The world matrix is **not** carried here — it is derived on the GPU (Step 15.5) +/// and read from the matrix buffer by the render pipeline. +#[derive(Clone)] +pub struct SlotDraw { + /// Stable slot index (offset into the indirect-args and matrix buffers, in slot units). + pub slot_index: usize, + /// Whether the slot is active (false = tombstone; the CPU skips it, the GPU zeros its draw args). + pub active: bool, + /// The entity's mesh (vertex/index buffers + material). + pub mesh: Arc, + /// Whether the mesh is indexed (`draw_indexed_indirect` vs `draw_indirect`). + pub has_index: bool, +} + /// Resource depot and entity graph. Stores Meshes and Materials keyed by identifier strings, /// maps entity labels to their associated `Entity` (mesh + transform) for rendering iteration, /// and holds the scene's active `Camera` used to build the per-frame view/projection matrices (Step 4.3). @@ -69,6 +105,13 @@ pub struct Scene { /// `None` = shadows off (default, non-regression). Read each frame by `Renderer::render_scene` /// to compute the light `view_proj` and enable shadow sampling. shadow_caster: Option, + /// Stable, append-only entity slots for the GPU-driven buffers (Phase 3). Grows only; removed + /// entities leave tombstones so slot indices stay stable. + entity_slots: Vec, + /// Map of entity label to slot index (O(1) lookup so a re-added label reuses its slot). + slot_of_label: HashMap, + /// Ordered mesh identifiers (index = the stable `mesh_index` used by slots and the bbox buffer). + mesh_order: Vec, } impl Scene { @@ -88,6 +131,9 @@ impl Scene { lights: Lights::new(), ambient: [1.0, 1.0, 1.0], shadow_caster: None, + entity_slots: Vec::new(), + slot_of_label: HashMap::new(), + mesh_order: Vec::new(), } } @@ -227,6 +273,7 @@ impl Scene { mesh.set_material(mat); } self.meshes.insert(id.to_string(), Arc::new(mesh)); + self.mesh_order.push(id.to_string()); Ok(id.to_string()) } @@ -399,6 +446,7 @@ impl Scene { return Err(format!("Mesh ID '{}' already exists.", id)); } self.meshes.insert(id.to_string(), mesh); + self.mesh_order.push(id.to_string()); Ok(id.to_string()) } @@ -440,6 +488,40 @@ impl Scene { } self.entities .insert(label.to_string(), Entity::new(mesh_id, transform)); + // Keep the stable slot in sync (Phase 3): a re-added label reuses its slot (stable index); + // a new label appends a slot. The mesh index + draw metadata are read from the mesh. + let mesh_index = self + .mesh_order + .iter() + .position(|id| id == mesh_id) + .expect("mesh validated above") as u32; + let mesh = &self.meshes[mesh_id]; + let has_index = mesh.index_buffer.is_some(); + let draw_count = if has_index { + mesh.num_indices + } else { + mesh.num_vertices + }; + let slot_index = match self.slot_of_label.get(label) { + Some(&i) => i, + None => { + let i = self.entity_slots.len(); + self.slot_of_label.insert(label.to_string(), i); + self.entity_slots.push(EntitySlot { + label: label.to_string(), + mesh_index: 0, + draw_count: 0, + has_index: false, + }); + i + } + }; + self.entity_slots[slot_index] = EntitySlot { + label: label.to_string(), + mesh_index, + draw_count, + has_index, + }; Ok(label.to_string()) } @@ -497,6 +579,91 @@ impl Scene { pub fn entity_count(&self) -> usize { self.entities.len() } + + // ======================================================================== + // Phase 3 — GPU-driven entity slot accessors (Step 15) + // ======================================================================== + // These expose the stable slot system to the Renderer: the packed transform slots (uploaded + // to the GPU each frame), the per-slot draw descriptors (for the indirect render loop), the + // per-mesh bounding boxes (uploaded once), and the slot/mesh counts. The world matrices are + // derived on the GPU; these methods only feed the CPU→GPU inputs and the draw-loop metadata. + + /// Packs the stable entity slots into GPU [`TransformSlot`]s (one per slot; tombstones → + /// inactive). The renderer uploads this to the transform buffer each frame (Phase 3, Step 15). + pub fn packed_transform_slots(&self) -> Vec { + self.entity_slots + .iter() + .map(|slot| match self.entities.get(&slot.label) { + Some(entity) => TransformSlot::from_transform( + entity.transform(), + slot.mesh_index, + slot.draw_count, + slot.has_index, + ), + None => TransformSlot::inactive(), + }) + .collect() + } + + /// Iterates the stable entity slots as per-slot draw descriptors (Phase 3). Each item carries + /// the slot index (→ indirect-args/matrix buffer offset), whether the slot is active (tombstones + /// are skipped on the CPU), the mesh, and whether it is indexed. Used by the indirect render loop. + pub fn iter_slot_draws(&self) -> impl Iterator + '_ { + self.entity_slots + .iter() + .enumerate() + .map(|(i, slot)| SlotDraw { + slot_index: i, + active: self.entities.contains_key(&slot.label), + mesh: self.meshes[&self.mesh_order[slot.mesh_index as usize]].clone(), + has_index: slot.has_index, + }) + } + + /// The local-space bounding boxes for all registered meshes, in `mesh_index` order (one per + /// mesh). Uploaded once to the GPU bbox buffer (Phase 3). Meshes without a bounding box get a + /// degenerate (all-zero) box, which the cull pass treats as a zero-radius sphere. + pub fn mesh_bboxes(&self) -> Vec { + self.mesh_order + .iter() + .map(|id| { + self.meshes[id] + .geometry() + .bbox() + .map(|b| BBoxSlot::from_bbox(&b)) + .unwrap_or_else(BBoxSlot::empty) + }) + .collect() + } + + /// Stable index of `mesh_id` in the ordered mesh list (its slot in the GPU bbox buffer), if + /// the mesh is registered. + pub fn mesh_index_of(&self, mesh_id: &str) -> Option { + self.mesh_order + .iter() + .position(|id| id == mesh_id) + .map(|p| p as u32) + } + + /// The registered mesh at a stable `mesh_index` (panics if the index is out of range; in + /// practice it is always valid, being derived from `mesh_order` positions). + pub fn mesh_by_index(&self, index: u32) -> &Arc { + &self.meshes[&self.mesh_order[index as usize]] + } + + /// Number of entity slots (tombstones included) — the `num_slots` written to the cull uniforms + /// (slots at/beyond this are no-ops on the GPU). + pub fn num_slots(&self) -> usize { + self.entity_slots.len() + } + + /// Number of *active* entity slots (tombstones excluded) — the live entity count. + pub fn num_active_slots(&self) -> usize { + self.entity_slots + .iter() + .filter(|s| self.entities.contains_key(&s.label)) + .count() + } } #[cfg(test)] @@ -586,4 +753,44 @@ mod tests { assert!(!scene.set_entity_transform("missing", Transform::identity())); assert!(!scene.remove_entity("missing")); } + + #[test] + fn gpu_driven_slot_bookkeeping_is_empty_when_no_entities() { + // The GPU-driven slot system starts empty; the full slot/mesh interplay requires a + // wgpu device (to build meshes) and is validated by the examples. + let scene = Scene::new(); + assert_eq!(scene.num_slots(), 0); + assert_eq!(scene.num_active_slots(), 0); + assert!(scene.packed_transform_slots().is_empty()); + assert!(scene.mesh_bboxes().is_empty()); + assert!(scene.iter_slot_draws().next().is_none()); + assert!(scene.mesh_index_of("nope").is_none()); + } + + #[test] + fn packed_slot_roundtrips_transform_and_is_inactive_when_tombstoned() { + // Exercises the CPU→GPU packing (GPU-independent): an active slot packs the transform + + // mesh index + draw count + index flag; a tombstoned slot packs to `inactive()`. + let t = Transform { + translation: glam::Vec3::new(1.0, 2.0, 3.0), + ..Transform::identity() + }; + let slot = TransformSlot::from_transform(&t, 7, 36, true); + assert_eq!(slot.mesh_index(), 7, "mesh index packed in flags.x"); + assert!(slot.is_active(), "active = 1"); + assert_eq!(slot.draw_count(), 36, "draw count packed in flags.z"); + assert!(slot.has_index(), "indexed flag packed in flags.w"); + assert!( + slot.translation + .iter() + .zip([1.0, 2.0, 3.0].iter()) + .all(|(a, b)| (a - b).abs() < 1e-5) + ); + + let inactive = TransformSlot::inactive(); + assert!(!inactive.is_active(), "inactive active-flag = 0"); + assert_eq!(inactive.mesh_index(), 0); + assert_eq!(inactive.draw_count(), 0); + assert!(!inactive.has_index()); + } } diff --git a/lib/src/shaders/gpu_driven.wgsl b/lib/src/shaders/gpu_driven.wgsl new file mode 100644 index 0000000..ed6cd9b --- /dev/null +++ b/lib/src/shaders/gpu_driven.wgsl @@ -0,0 +1,188 @@ +// # GPU-driven rendering compute shader (Phase 3, Step 15) +// +// Two compute entry points run sequentially in a single command encoder, before the render passes: +// 1. `compute_matrices` derives each entity's world matrix on the GPU from its transform slot. +// 2. `cull` decides per-entity visibility (bounding sphere vs frustum) and fills the indirect +// draw arguments (the vertex/index count, zeroed when the entity is culled or inactive). +// +// The main and shadow render passes are then 100% indirect: they read the draw slots (zero count +// = no-op) instead of a CPU-side per-entity loop. +// +// All buffers are fixed-capacity (MAX_ENTITIES = 256, see DRAFT D12) and allocated once. Per frame the CPU +// rewrites only the transform slots and the cull uniforms; everything else is GPU-driven. +// +// GPU buffer layouts mirror the bytemuck structs in `resources::uniform` (byte-for-byte): +// TransformSlot (64B), MatSlot (256B), BBoxSlot (32B), DrawSlot (80B), CullUniforms (112B). +// +// GOTCHA — WGSL `select` argument order: `select(reject, accept, cond)` returns the SECOND +// argument when `cond` is true and the FIRST when false (the reverse of HLSL's +// `select(trueVal, falseVal, cond)`). The original cull pass wrote +// `select(u32(t.flags.z), 0u, visible)`, which zeroed the count of every VISIBLE entity (and +// would have drawn the culled ones) — the source of the black-window bug. + +// 64 bytes: packed TRS, the single source of truth for world matrices. +struct TransformSlot { + translation : vec3f, + flags : vec4f, // x = mesh index, y = active, z = draw count, w = has_index + rotation : vec4f, + scale : vec3f, +}; + +// 256 bytes: a 64-byte world matrix followed by 192 bytes of padding. The padding is REQUIRED — +// the render pipelines read this slot through the `uniform` object group with a per-slot dynamic +// offset, and WebGPU demands that offset be a multiple of `min_uniform_buffer_offset_alignment` +// (256 bytes). A bare 64-byte matrix can never be individually addressable via a uniform offset, +// so each slot is padded to a 256-byte boundary (capacity is capped at 256 = 64 KB / 256 B). +struct MatSlot { + m : mat4x4f, + pad : array, // 192 bytes, alignment only — never read +}; + +// 32 bytes: local-space axis-aligned bounding box (uploaded once per mesh). +struct BBoxSlot { + min : vec3f, + max : vec3f, +}; + +// 80 bytes: indirect draw arguments for one entity (a 4-u32 non-indexed / 5-u32 indexed block). +// Only `.a` is written by the shader; the rest stays zero (instance count is the constant 1 in `.a.y`). +struct DrawSlot { + a : vec4u, + b : vec4u, + c : vec4u, + d : vec4u, + e : vec4u, +}; + +// 112 bytes: frustum planes + control flags, rewritten by the CPU each frame. +struct CullUniforms { + planes : array, // unit (normal, d); inside iff dot(p, normal) + d >= 0 + num_slots : u32, // number of live entity slots + culling : u32, // 0 = culling disabled, 1 = enabled + _pad : vec2u, +}; + +// ---- Bind groups (Step 15.5) ---- +// Group 0 (transforms) is shared by both entry points; group 1 (matrices) by `compute_matrices`; +// group 2 (cull uniforms + bboxes + draw args) by `cull`. Each pipeline infers the subset it uses. +@group(0) @binding(0) var transforms : array; +@group(1) @binding(0) var matrices : array; +@group(2) @binding(0) var cull_u : CullUniforms; +@group(2) @binding(1) var bboxes : array; +@group(2) @binding(2) var draw_args : array; + +// ---- Shared helpers ---- + +// Builds a rotation mat4x4f from a quaternion (x, y, z, w) in column-major form. +fn quat_to_mat4(q : vec4f) -> mat4x4f { + let x = q.x; + let y = q.y; + let z = q.z; + let w = q.w; + return mat4x4f( + vec4f(1.0 - 2.0 * (y * y + z * z), 2.0 * (x * y + w * z), 2.0 * (x * z - w * y), 0.0), + vec4f(2.0 * (x * y - w * z), 1.0 - 2.0 * (x * x + z * z), 2.0 * (y * z + w * x), 0.0), + vec4f(2.0 * (x * z + w * y), 2.0 * (y * z - w * x), 1.0 - 2.0 * (x * x + y * y), 0.0), + vec4f(0.0, 0.0, 0.0, 1.0) + ); +} + +// Rotates a local-space vector by a quaternion (via the rotation matrix). +fn rotate_by_quat(v : vec3f, q : vec4f) -> vec3f { + let m = quat_to_mat4(q); + return (m * vec4f(v, 0.0)).xyz; +} + +// World matrix = T * R * S, column-major (matches the CPU `Transform::to_matrix`, D13). +// Returns just the 4x4 matrix; the caller stores it in `matrices[i].m` (the slot's 192-byte pad +// is left at its zero-initialised value). +fn world_matrix(t : TransformSlot) -> mat4x4f { + let r = quat_to_mat4(t.rotation); + return mat4x4f( + r[0] * t.scale.x, + r[1] * t.scale.y, + r[2] * t.scale.z, + vec4f(t.translation.x, t.translation.y, t.translation.z, 1.0) + ); +} + +// The identity world matrix (used for inactive slots so any stale read is harmless). +fn identity_mat() -> mat4x4f { + return mat4x4f( + vec4f(1.0, 0.0, 0.0, 0.0), + vec4f(0.0, 1.0, 0.0, 0.0), + vec4f(0.0, 0.0, 1.0, 0.0), + vec4f(0.0, 0.0, 0.0, 1.0) + ); +} + +// Fills a draw slot with a non-zero (visible) count of `count`, or zero (culled / inactive). +// The count lands in `.a.x`; `.a.y` (instance count) is the constant 1; the rest stays zero. +fn set_draw_count(i : u32, count : u32) { + draw_args[i].a = vec4u(count, 1u, 0u, 0u); +} + +// ---- Pass 1: derive world matrices (Step 15.5) ---- +// Dispatched for MAX_ENTITIES; inactive slots get the identity matrix (a harmless stale read). +@compute +@workgroup_size(64) +fn compute_matrices(@builtin(global_invocation_id) gid : vec3u) { + let i = gid.x; + let t = transforms[i]; + if (t.flags.y < 0.5) { + matrices[i].m = identity_mat(); + } else { + matrices[i].m = world_matrix(t); + } +} + +// ---- Pass 2: cull + fill indirect draw args (Step 15.6) ---- +// Dispatched for MAX_ENTITIES. Slots at or beyond `num_slots` (and inactive slots) are zeroed so +// the indirect render passes skip them; visible slots keep their packed count. +@compute +@workgroup_size(64) +fn cull(@builtin(global_invocation_id) gid : vec3u) { + let i = gid.x; + let t = transforms[i]; + + // Beyond the live slots: zero the count so the indirect draw is a no-op. + if (i >= cull_u.num_slots) { + set_draw_count(i, 0u); + return; + } + + // Inactive slot (tombstone): no draw. + if (t.flags.y < 0.5) { + set_draw_count(i, 0u); + return; + } + + // Culling disabled: every active entity is visible, with its packed count. + if (cull_u.culling == 0u) { + set_draw_count(i, u32(t.flags.z)); + return; + } + + // Culling enabled: test the entity's world bounding sphere against the frustum planes. + let b = bboxes[u32(t.flags.x)]; + let center_local = (b.min + b.max) * 0.5; + // World center = translation + rotation * local center (no scale; the radius carries the scale). + let center_world = t.translation + rotate_by_quat(center_local, t.rotation); + let half_extents = (b.max - b.min) * 0.5; + let radius = length(half_extents) * max(t.scale.x, max(t.scale.y, t.scale.z)); + + var visible = true; + for (var p = 0u; p < 6u; p = p + 1u) { + let plane = cull_u.planes[p]; + let dist = dot(plane.xyz, center_world) + plane.w; + if (dist < -radius) { + visible = false; + break; + } + } + + // NOTE: WGSL `select(reject, accept, cond)` — the accept value is the SECOND argument. + // visible => full count; culled => 0. + let count = select(0u, u32(t.flags.z), visible); + set_draw_count(i, count); +} diff --git a/lib/src/utils/conf.rs b/lib/src/utils/conf.rs index aa9231f..33c20eb 100644 --- a/lib/src/utils/conf.rs +++ b/lib/src/utils/conf.rs @@ -35,6 +35,32 @@ pub const SHADOW_SHADER_PATH: &str = "assets/shaders/shadow_shader.wgsl"; /// (Step 14, D4). Serves as the fallback when `SHADOW_SHADER_PATH` cannot be read. pub const SHADOW_SHADER: &str = include_str!("../shaders/shadow_shader.wgsl"); +/// The GPU-driven rendering compute shader source (Phase 3, Step 15), embedded at compile time. +/// It carries two compute entry points — `compute_matrices` (derive world matrices) and `cull` +/// (per-entity visibility + indirect draw args) — compiled directly by the renderer (internal to +/// the library; no external file is read). +pub const GPU_DRIVEN_SHADER: &str = include_str!("../shaders/gpu_driven.wgsl"); + +/// Fixed capacity of the GPU-driven entity slot buffers (Phase 3). The transform, matrix, bbox and +/// indirect-draw-args buffers are all sized to this capacity and allocated once; per frame the CPU +/// rewrites only the transform slots and the cull uniforms. +/// +/// **Why 256 (and why the matrix slot is padded to 256 B):** the matrix buffer is bound to the +/// render pipeline's `uniform` object slot (group 1), which imposes TWO limits: +/// (1) a single uniform binding is capped at `max_uniform_buffer_binding_size` (64 KB on most +/// backends), and (2) a uniform offset must be a multiple of `min_uniform_buffer_offset_alignment` +/// (256 B). A 64-byte matrix can therefore never be individually addressable via a per-slot +/// dynamic offset, so each matrix slot is padded to 256 B (see `MatSlot`); with 256-byte slots, +/// 256 slots × 256 B is exactly 64 KB — the largest the single-buffer / dynamic-offset design +/// can address. (The transform / bbox / draw-args buffers are `storage` bindings with a 128 MB +/// limit and no 256-B offset rule, so they keep their natural 64 / 32 / 80 B slot sizes.) +/// 256 is a multiple of the 64-wide workgroup size, giving a whole number of workgroups. +pub const MAX_ENTITIES: u32 = 256; + +/// Workgroup size of the GPU-driven compute shaders (matches the `@workgroup_size` in +/// `gpu_driven.wgsl`). The compute dispatch is `MAX_ENTITIES / WORKGROUP_SIZE` workgroups. +pub const GPU_WORKGROUP_SIZE: u32 = 64; + /// Default shadow-map resolution in pixels per side (square, D2). A 1024² depth map is a good /// quality/cost trade-off for the dedicated `shadow_test` example and most simple scenes. pub const SHADOW_MAP_SIZE: u32 = 1024; diff --git a/lib/tests/wgsl_validate.rs b/lib/tests/wgsl_validate.rs index 2e16254..7e80864 100644 --- a/lib/tests/wgsl_validate.rs +++ b/lib/tests/wgsl_validate.rs @@ -54,3 +54,34 @@ fn shadow_shader_is_valid_wgsl() { .collect(); assert_eq!(entry_names, vec!["vs_main"], "only vs_main expected"); } + +/// Parses and fully validates the embedded `gpu_driven.wgsl` compute shader (Phase 3, Step 15) +/// via naga. The renderer compiles it directly into two `ComputePipeline`s (one per entry point), +/// so this offline validation is the guarantee of its validity. The contract expects exactly the +/// two compute entry points: `compute_matrices` and `cull`. +#[test] +fn gpu_driven_shader_is_valid_wgsl() { + let src = include_str!("../src/shaders/gpu_driven.wgsl"); + let module = naga::front::wgsl::parse_str(src) + .unwrap_or_else(|e| panic!("gpu_driven.wgsl: parsing error: {e:?}")); + + let mut validator = naga::valid::Validator::new( + naga::valid::ValidationFlags::all(), + naga::valid::Capabilities::all(), + ); + validator + .validate(&module) + .unwrap_or_else(|e| panic!("gpu_driven.wgsl: validation failed: {e:?}")); + + let mut entry_names: Vec<&str> = module + .entry_points + .iter() + .map(|ep| ep.name.as_str()) + .collect(); + entry_names.sort(); + assert_eq!( + entry_names, + vec!["compute_matrices", "cull"], + "the two compute entry points are expected" + ); +}