From 531c43a4573efa090003a16e91a53124777f2bec Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?J=C3=A9r=C3=B4me=20Bousqui=C3=A9?= Date: Tue, 22 Sep 2026 17:07:52 +0200 Subject: [PATCH] material batching --- README.md | 10 +- docs/DRAFT.md | 512 ++++++++++---------------------- docs/ROADMAP.md | 10 +- docs/tech/ARCHI_APP.md | 4 +- docs/tech/ARCHI_CPU_GPU.md | 12 +- docs/tech/FRAME_LOOP.md | 2 +- docs/user/gpu-driven.md | 16 +- lib/src/core/renderer.rs | 149 ++++++++-- lib/src/shaders/gpu_driven.wgsl | 5 +- 9 files changed, 313 insertions(+), 407 deletions(-) diff --git a/README.md b/README.md index df25c73..96ad0f2 100644 --- a/README.md +++ b/README.md @@ -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) | ✅ 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) | +| 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, 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). @@ -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`) 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_APP](docs/tech/ARCHI_APP.md) — engine architecture. ✅ **Current** — facade (`App`/`AppHandler`) and GPU-driven two-pass pipeline (implemented in Phase 3, 2026-09-22, 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-09-22, É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, 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.)* +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-09-22 — see item 17. Deviation from the original spec: one indirect draw **per slot** rather than a single fused draw, D1 — see ARCHI_CPU_GPU.)* 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,4 +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.) +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-09-22; WGSL + frustum + scene-slot tests, 57 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 D14 in ARCHI_CPU_GPU.md.) diff --git a/docs/DRAFT.md b/docs/DRAFT.md index db13a0c..a03f685 100644 --- a/docs/DRAFT.md +++ b/docs/DRAFT.md @@ -1,404 +1,190 @@ -# DRAFT — Étape 17 : Rendu GPU-driven (ROADMAP 3.1 / 3.2 / 3.3) +# DRAFT — Étape 18 : Batching par Material (ROADMAP 4.3) -> **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). +> **Statut** : brouillon de conception (à valider avant implémentation). +> Couvre le 1er item de la Phase 4.3 de `docs/ROADMAP.md` : +> « Batching par Material (réduction des state changes GPU) ». > -> **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). -> -> **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). +> **Archive** : le draft Étape 17 (rendu GPU-driven) a été vidé après validation (2026-09-22). +> Référence durable : `docs/tech/ARCHI_CPU_GPU.md` (écarts D1/D4/D5/D12 + piège D14) ; +> texte intégral : git `3a424af` (`git show 3a424af:docs/DRAFT.md`). -## Objectifs +## Contexte — où on en est -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. +Le rendu est 100 % indirect (Étape 17) : un draw indirect **par slot**, dans l'ordre d'insertion +des entités. À chaque draw, la passe principale met à jour (`Renderer::render_scene`, +`renderer.rs` §7) : + +| Appel | Coût | +|---|---| +| `set_pipeline(material.pipeline)` | **changement d'état** (swap de pipeline côté driver) | +| `set_bind_group(0, frame)` | constant dans la passe (re-set inutile mais pas cher) | +| `set_bind_group(1, object, [offset dynamique])` | paramètre de draw — **pas** un changement d'état | +| `set_bind_group(2, material.texture_bind_group)` | **changement d'état** (un bind group par `Material`) | +| `set_bind_group(3, shadow)` | constant | +| `set_vertex_buffer` / `set_index_buffer` | par mesh, pas cher | +| `draw_*_indirect` | le draw lui-même | + +→ Le nombre de changements d'état (pipeline + bind group @2) est proportionnel au **nombre +d'entités**, même quand des dizaines d'entités partagent le même `Material`. La doc du module +promet déjà « Entity sorting ... minimizes pipeline switches (batching by material) » — c'est ce +que fait cette étape. + +Le pass d'ombre n'a qu'**un** pipeline et l'appelle **une seule fois** avant la boucle +(`renderer.rs` §5) → déjà « batché » sur l'état : rien à y faire. + +## Objectif + +Réduire les changements d'état de la passe principale de **O(entités)** à **O(matériaux +distincts)**, sans changer le rendu, sans changer l'API publique, et sans toucher au chemin +GPU-driven (compute, culling, buffers, indirect) ni au pass d'ombre. ## 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). +- **Aucune rupture d'API** : optimisation interne du `Renderer` ; aucun type/méthode publique + nouveau ; les exemples ne sont pas modifiés. +- **Aucune régression visuelle** : tous les pipelines de l'engine sont **opaques** + (`BlendState::REPLACE`, `pipeline_cache.rs`) → le depth buffer résout l'ordre → réordonner + les draws est visuellement neutre (D4). +- **Culling/indirect intacts** : les verdicts GPU (args = 0 → no-op) et les buffers ne changent + pas ; seul l'**ordre d'émission** des draws change. +- **Déterminisme** : l'ordre des draws doit rester reproductible frame après frame (slots + append-only stables, Étape 17 D9). +- Capacité ≤ 256 slots → le groupage (O(N), `HashMap` ≤ 256 entrées) est du bruit ; on le + refait **chaque frame** (toujours correct, aucun cache à invalider). -## Décisions (à valider) +## Décisions -### D1 — Un draw indirect par entité (pas de multi-instancing par groupe de mesh) +### D1 — Clé de groupage : l'identité du `Material` (pointeur `Arc`), pas le shader_id +- Le vrai « état » qui change entre deux draws partageant un pipeline est le **bind group @2** + (texture/sampler) : chaque `Material` en possède un propre (`Material::build`). Deux + materials qui partagent le même `shader_id` partagent le pipeline (Arc, `PipelineCache`) + mais **pas** le bind group @2 → grouper au niveau pipeline ferait alterner le @2. +- Clé retenue : `Arc::as_ptr(&material)` — même pointeur ⟺ même objet `Material` ⟺ même + pipeline **et** même bind group @2 → les deux changements d'état sont figés dans le groupe. +- Le fallback `scene.default_material()` est un `Arc` unique en cache (`RefCell`) → toutes les + entités sans material forment un groupe. +- **Subtilité de durée de vie** : le groupage matérialise d'abord les `Arc` dans un + `Vec` (un par slot actif) ; les pointeurs-clés ne sont dérivés qu'ensuite. Les `Arc` restent + donc vivants pendant toute la passe → aucun pointeur ne pend (le `Arc` retourné par + `default_material()` est un clone : sans le `Vec`, il serait libéré en fin de closure). -Chaque entité possède son **slot GPU** (index stable `0..N`). Un draw indirect par entité lit -les arguments de son slot. +### D2 — Ordre des groupes : première apparition en ordre de slots ; intra-groupe : ordre de slots +- On parcourt les slots dans leur ordre stable (insertion, Étape 17 D9) ; le groupe d'un slot + est créé à la **première** apparition de sa clé. `HashMap` + `Vec` : + O(N), sans tri, **déterministe** et stable frame à frame (scène inchangée → ordre inchangé). +- En pratique, le premier draw de chaque groupe est le même que l'ancien premier draw du slot + → l'impression visuelle est préservée ; seuls les draws de matériaux *différents* + s'intercalent moins. -- ✅ 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. +### D3 — Les slots cullés (no-op) restent émis dans leur groupe +- Le CPU ne connaît pas le verdict GPU du culling (un readback par frame stallerait la boucle + — cf. `debug_dump`) : le draw d'un slot cullé est un indirect **zéro count ≈ gratuit** ; on + l'émet quand même, dans son groupe. +- Le nombre de **draw calls** est donc inchangé (1 par slot actif) ; seul le nombre de + **changements d'état** baisse. Réduire aussi les draw calls = multi-instancing (exclu — + Périmètre). -### D2 — Le bind group « object » par entité devient une slice du buffer GPU-computé +### D4 — Réordonnancement sûr : pipelines 100 % opaques +- `pipeline_cache.rs` crée toutes les pipelines avec `blend: Some(BlendState::REPLACE)` + (aucun alpha blending dans l'engine) et le depth write est actif partout (Étape 9) → + l'ordre de rasterisation n'a pas d'impact visuel. +- **Contrainte à documenter** (docs user + rustdoc) : si du blending transparent est ajouté un + jour, il faudra isoler les matériaux transparents (trier back-to-front en fin de passe) — + signalé ici comme prérequis d'un futur `Material.blend`. Le groupage par Material reste + correct en l'état ; seule l'ordre inter-groupes devra évoluer. -- 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). +### D5 — Groupage en fonction pure, testable sans GPU +- Le groupage est factorisé en fonction libre : + `fn batch_slots(keys: &[K]) -> Vec>` + (groupes dans l'ordre de première apparition de la clé ; indices dans l'ordre d'entrée ; + chaque indice apparaît exactement une fois). +- Le `Renderer` l'appelle avec `keys = [*const Material]` (D1) ; les **tests unitaires** + l'appellent avec des clés `u32` → testable sans instance/device wgpu (un `Material` exige + un pipeline compilé = device ; les tests de la crate restent headless/CI-safe). +- Les tombstones (`active == false`) sont filtrés **avant** l'appel (comme aujourd'hui) : + `batch_slots` ne voit que les slots actifs. -### D3 — Layout des slots d'indirect draw : 80 o (pdc(16,20)) +### D6 — Pass d'ombre : inchangé (déjà batché) +- Un seul `shadow_pipeline`, `set_pipeline` une seule fois avant la boucle ; l'état résiduel + par slot (offset dynamique @1 + vertex/index buffers) n'est pas un changement d'état + driver → `render_shadow_map` n'est pas modifié par cette étape. -- `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) +## Mécanisme par frame (seul le point 7 change) ``` -[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) +[CPU] write_buffer TransformBuffer / CullUniforms / (BBoxBuffer si génération changée) +[Compute 1] compute_matrices [Compute 2] cull (inchangés) +[Ombre] (si caster) draw indirect par slot, 1 pipeline (inchangé — D6) +[Main] slots groupés par Material (D1/D2) : + groups = batch_slots(keys) + pour chaque groupe G : + set_pipeline(G) + set_bind_group(0) + set_bind_group(2, G) + set_bind_group(3) + pour chaque slot s de G : + set_bind_group(1, [offset(s)]) + set_vertex/set_index + draw_indirect(s) +[Queue] submit ``` -Même encoder → exécution séquentielle garantie (write → compute → render). +## Gain attendu (changements d'état / frame, passe principale) -## Layouts GPU (WGSL) +| Scène | Avant (par entité) | Après (par matériau distinct) | +|---|---|---| +| `demo` (7 entités, ~5 materials distincts) | 7 × (pipeline + @2) | 5 × (pipeline + @2) | +| 200 cubes / 1 material | 200 × (pipeline + @2) | **1** × (pipeline + @2) | -```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`. +Le `set_bind_group(1, offset dynamique)` et les vertex/index buffers restent par draw +(paramètres de draw, pas d'état driver). Le gain est maximal quand peu de matériaux distincts +pour beaucoup d'entités — le cas « instancé » en attendant le multi-instancing. ## 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 | +| `lib/src/core/renderer.rs` | `render_scene` : boucle plate → `batch_slots` + boucle par groupe ; fonction libre `batch_slots` + tests unitaires ; rustdoc du module alignée | +| `docs/user/gpu-driven.md` | + note : draws groupés par matériau (interne, sans effet API) ; contrainte blending (D4) | +| `docs/tech/ARCHI_CPU_GPU.md` | + note : la passe main émet les draws groupés par Material (Étape 18) | +| `docs/ROADMAP.md` | Coche « Batching par Material » (4.3) + date | ## 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). +- **Multi-instancing** (1 draw par groupe mesh+material, matrice par instance) : demande un + changement de shader (matrice instanciée) + draw instancié par groupe — étape distincte, + plus lourde (déjà reportée depuis l'Étape 17). +- **LOD, HDR + tone mapping** (items 2-3 de la Phase 4.3) : étapes distinctes. +- **Sauter les no-ops par readback** (réduire les draw calls, pas seulement les états) : un + readback synchro stalle la boucle de rendu → rejeté en v1. +- **Blending/transparent** : hors engine actuel (D4). ## 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 »). | +| **Réordonnancement → rendu différent** | D4 : pipelines opaques (`REPLACE`) + depth write → le depth buffer résout l'ordre ; vérifié headless + visuellement (plan 4-5), draw args `debug_dump` inchangés (plan 6). | +| **Pointeur-clé `Arc::as_ptr` en pend** | D1 : les `Arc` sont matérialisés dans un `Vec` vivant pendant la passe ; les groupes reconstruits chaque frame → aucune hypothèse de stabilité entre frames. | +| **Groupage O(N) par frame** | N ≤ 256, `HashMap` ≤ 256 entrées → bruit ; pas de cache (D5 : toujours correct). | +| **Ordre des groupes non déterministe** | D2 : première apparition sur des slots stables (Étape 17 D9) → déterministe ; verrouillé par test (D5). | +| **Matériau transparent futur** | D4 documenté comme contrainte ; le groupage par Material reste correct, seule l'ordre inter-groupes devra évoluer. | ## 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. +1. `cargo build --workspace` — OK, sans avertissement. +2. `cargo test --workspace` — vert, y compris les nouveaux tests `batch_slots` (première + apparition, ordre intra-groupe, chaque indice une fois, entrées vides, un seul groupe, + tout distinct). 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`). +4. **`demo` headless** (`WGPU_BACKEND=vulkan timeout 10 ... --example demo`) → exit 0. +5. **A/B changements d'état** : compteur temporaire de `set_pipeline` par frame (sous + `WSG_DEBUG_DUMP`) — avant : 7 dans le `demo` ; après : nombre de materials distincts. + (Compteur retiré après mesure, ou conservé sous `WSG_DEBUG_DUMP` au choix.) +6. **Rendu identique** : `demo` avant/après → même image (opaque, D4) ; `debug_dump` : + draw args inchangés (le culling n'est pas touché). +7. Check liens doc — 0 lien cassé. ## 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 +- [x] `render_scene` émet les draws **groupés par Material** (D1/D2) ; pass d'ombre inchangé (D6). +- [x] `batch_slots` fonction pure testable (D5) + tests unitaires verts (6 tests, CI-safe). +- [x] Aucune rupture d'API publique ; exemples non modifiés. +- [x] Rendu identique (D4) — vérifié headless ; draw args `debug_dump` inchangés (6/36/3840/960/384/192/2304). +- [x] Changements d'état réduits : compteur `set_pipeline` du `demo` = 3 = nb de materials distincts (avant : 7). +- [x] Docs : `gpu-driven.md` § « Batching by material » + `ARCHI_CPU_GPU.md` + ROADMAP 4.3 coché. +- [ ] DRAFT.md vidé après validation utilisateur (convention de la maison). diff --git a/docs/ROADMAP.md b/docs/ROADMAP.md index 4321d5c..ff83094 100644 --- a/docs/ROADMAP.md +++ b/docs/ROADMAP.md @@ -117,10 +117,10 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z } --- -## Phase 3️⃣ — GPU-Driven Rendering ✅ (2026-07-20) +## Phase 3️⃣ — GPU-Driven Rendering ✅ (2026-09-22) **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). +**Statut** : implémenté (Étape 17, validé 2026-09-22, décisions D1–D14 — référence durable : `docs/tech/ARCHI_CPU_GPU.md`, texte intégral du draft : git `3a424af`). 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 (D14). ### 3.1 Compute Shader - [x] Buffer `TransformBuffer` (CPU → GPU) : positions/rotations/échelles brutes (`TransformSlot`, 64 B) @@ -130,7 +130,7 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z } ### 3.2 Frustum Culling 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] 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. D14 dans `docs/tech/ARCHI_CPU_GPU.md`)* - [x] Buffer `IndirectDrawBuffer` rempli par le GPU (`DrawSlot`, 80 B, zéro = no-op) ### 3.3 Rendu Indirect @@ -161,7 +161,7 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z } par défaut**. Exemple `shadow_test` : cube projetant une ombre douce sur un sol.)* ### 4.3 Optimisations -- [ ] Batching par Material (réduction des state changes GPU) +- [x] Batching par Material (réduction des state changes GPU) — 2026-09-22 (Étape 18 : draws groupés par `Arc` dans la passe principale, 1 `set_pipeline` par matériau distinct — le démo passe de 7 à 3 ; pass d'ombre inchangé) - [ ] Level of Detail (LOD) - [ ] HDR + Tone Mapping (optionnel) @@ -199,4 +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` | +| **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 (D14). Piège documenté en tête de `gpu_driven.wgsl`, dans `AGENTS.md` et `docs/tech/ARCHI_CPU_GPU.md` | diff --git a/docs/tech/ARCHI_APP.md b/docs/tech/ARCHI_APP.md index 39c99ca..6c17283 100644 --- a/docs/tech/ARCHI_APP.md +++ b/docs/tech/ARCHI_APP.md @@ -16,7 +16,7 @@ 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** — 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 +> (§1, §4B, §5, §6) **implémenté en Phase 3** du ROADMAP (Étape 17, 2026-09-22, 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**, @@ -25,7 +25,7 @@ wsg_lib est un moteur de rendu modulaire basé sur wgpu. Il adopte une architect > 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 +> écarts documentés (cf. `ARCHI_CPU_GPU.md`) : 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**. diff --git a/docs/tech/ARCHI_CPU_GPU.md b/docs/tech/ARCHI_CPU_GPU.md index 6e0d5d4..3b40fc2 100644 --- a/docs/tech/ARCHI_CPU_GPU.md +++ b/docs/tech/ARCHI_CPU_GPU.md @@ -16,19 +16,27 @@ 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 : ACTUEL (implémenté — Phase 3 du ROADMAP, 2026-07-20, DRAFT Étape 17).** +> **État du document : ACTUEL (implémenté — Phase 3 du ROADMAP, Étape 17, validé 2026-09-22).** > 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 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 +> Ce document est la **référence durable** de la conception : le draft d'origine de l'Étape 17 +> (décisions D1–D14, layouts, plan de validation) a été vidé de `docs/DRAFT.md` après validation +> et vit dans le git (`git show 3a424af:docs/DRAFT.md`) ; l'essentiel en est repris ci-dessous. +> **Écarts documentés** (numérotation du draft d'origine) : (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`. +> **Batching par material (Étape 18, 2026-09-22)** : la passe principale émet désormais les draws +> groupés par `Material` (1 `set_pipeline` + 1 bind group @2 par matériau distinct, pas par entité ; +> le pass d'ombre — un seul pipeline — est inchangé). Réordonnancement sûr car tous les pipelines +> sont opaques (`BlendState::REPLACE`) ; les no-ops cullés restent émis dans leur groupe. +> Détail : `docs/user/gpu-driven.md` § « Batching by material ». 1. Répartition des Rôles : CPU vs GPU (La Source de Vérité) diff --git a/docs/tech/FRAME_LOOP.md b/docs/tech/FRAME_LOOP.md index 9818797..b66a99e 100644 --- a/docs/tech/FRAME_LOOP.md +++ b/docs/tech/FRAME_LOOP.md @@ -68,7 +68,7 @@ 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) — implémenté (Phase 3, 2026-07-20)** : les buffers +> **Ressources GPU persistantes (single buffer) — implémenté (Phase 3, 2026-09-22)** : 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** diff --git a/docs/user/gpu-driven.md b/docs/user/gpu-driven.md index bc5a223..adbaf39 100644 --- a/docs/user/gpu-driven.md +++ b/docs/user/gpu-driven.md @@ -22,6 +22,20 @@ 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. +## Batching by material + +The main render pass batches the draws by material: all entities sharing the same material are +drawn back to back, so the GPU pipeline and the material's texture bind group are switched **once +per distinct material**, not once per entity (the per-draw work — matrix offset, vertex/index +buffers, the indirect draw itself — is unchanged). The grouping is internal: it does not change +the rendered image and there is nothing to configure. + +> **Constraint:** the batching reorders the draws, which is safe here because every pipeline in +> the engine is **opaque** (`BlendState::REPLACE`, no alpha blending) — the depth buffer resolves +> the draw order. If transparent materials are ever added, the transparent draws must be isolated +> (sorted back-to-front at the end of the pass) and must not interleave with the grouped opaque +> draws. + ## Frustum culling (opt-in) Culling is **off by default**. The culling pass still runs, but with culling disabled it marks @@ -108,7 +122,7 @@ objects were off-screen anyway — so the proof is in the counts, not the image) 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`.) +and verified fixed, see the D14 note in `docs/tech/ARCHI_CPU_GPU.md`.) ## Limitations diff --git a/lib/src/core/renderer.rs b/lib/src/core/renderer.rs index 25d3a63..20881ed 100644 --- a/lib/src/core/renderer.rs +++ b/lib/src/core/renderer.rs @@ -15,7 +15,8 @@ //! ## Architecture Notes (per ARCHI_APP.md) //! - **Execution Phase**: Renderer executes per-frame render loops. During this phase it iterates Scene entities //! and draws each one by binding the appropriate Material+Mesh pair. -//! - **Performance**: Entity sorting within the render loop minimizes pipeline switches (batching by material). +//! - **Performance**: the main pass batches draws by material (Étape 18), so the pipeline + +//! texture state changes happen once per distinct material, not once per entity. //! - **Low-Level Access**: Advanced users can bypass Scene and call Renderer directly for custom rendering paths. use crate::core::Context; @@ -40,6 +41,9 @@ use crate::utils::conf::{ }; use glam::{Mat4, Vec3, Vec4}; use std::cell::Cell; +use std::collections::HashMap; +use std::hash::Hash; +use std::sync::Arc; /// 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 @@ -118,6 +122,11 @@ pub struct Renderer { /// 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, + /// Number of `set_pipeline` calls in the LAST `render_scene` main pass. Since Étape 18 the + /// main pass batches by material, so this equals the number of DISTINCT materials drawn that + /// frame. Interior-mutable (all-`&self` API); exposed through `debug_dump` for the + /// state-change A/B measurement (Étape 18 verification). + debug_pipeline_switches: Cell, } impl Renderer { @@ -460,6 +469,7 @@ impl Renderer { cull_bundle_bg, matrix_object_bg, cull_enabled: Cell::new(false), + debug_pipeline_switches: Cell::new(0), }; // 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. @@ -755,8 +765,10 @@ impl Renderer { // 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. + // 7. Main render pass: one indirect draw per active slot, batched by material (Étape 18). + // 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. State changes (pipeline + texture bind group @2) + // are hoisted out of the slot loop: one per DISTINCT material, not one per entity. { let mut render_pass = encoder.begin_render_pass(&wgpu::RenderPassDescriptor { label: Some("scene render pass"), @@ -782,34 +794,54 @@ impl Renderer { ..Default::default() }); - 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_offset = (slot.slot_index as u64 * MAT_SLOT_SIZE) as u32; - let indirect_offset = slot.slot_index as u64 * DRAW_SLOT_SIZE; + // Batching by material (Étape 18): the material's Arc pointer is the group key — the + // same Material shares one pipeline AND one texture bind group (@2), so both state + // changes are frozen within a group. Groups appear in order of first appearance in + // stable slot order (D2), so the draw order stays deterministic frame to frame. All + // pipelines are opaque (BlendState::REPLACE), so reordering draws is visually neutral + // (D4 — if transparent blending is ever added, see the constraint in the user docs). + let slots: Vec<_> = scene.iter_slot_draws().filter(|s| s.active).collect(); + // Materialize the Material Arcs first: the pointer keys below must stay valid for the + // whole pass, and `default_material()` returns a clone that would otherwise be dropped + // at the end of the closure (D1). + let materials: Vec> = slots + .iter() + .map(|s| { + s.mesh + .material() + .cloned() + .unwrap_or_else(|| scene.default_material()) + }) + .collect(); + let keys: Vec<*const Material> = materials.iter().map(Arc::as_ptr).collect(); + let groups = batch_slots(&keys); + self.debug_pipeline_switches.set(groups.len() as u32); + for group in &groups { + // The first slot of a group carries the group's material (all slots in the group + // share the same Arc pointer). + let material = &materials[group[0]]; 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); + for &pos in group { + let slot = &slots[pos]; + let object_offset = (slot.slot_index as u64 * MAT_SLOT_SIZE) as u32; + let indirect_offset = slot.slot_index as u64 * DRAW_SLOT_SIZE; + // 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_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); } - render_pass.draw_indexed_indirect(&self.draw_args_buffer, indirect_offset); - } else { - render_pass.draw_indirect(&self.draw_args_buffer, indirect_offset); } } } @@ -958,6 +990,10 @@ impl Renderer { for (i, p) in c.planes.iter().enumerate() { eprintln!("[dbg] plane[{i}] = {p:?}"); } + eprintln!( + "[dbg] pipeline switches (this frame's main pass) = {}", + self.debug_pipeline_switches.get() + ); eprintln!("[dbg] done"); } } @@ -1157,3 +1193,64 @@ fn draw_entity( pass.draw(0..mesh.num_vertices, 0..1); } } + +/// Groups the positions of a key slice for material batching (Étape 18, D2/D5). Groups appear in +/// order of first key occurrence; indices within a group keep input order; every input index +/// appears exactly once. Pure and GPU-free, so it is unit-testable with integer keys. The `Clone` +/// bound only serves to keep an owned copy of each group's key (the real keys are `*const T`, +/// i.e. `Copy`). +fn batch_slots(keys: &[K]) -> Vec> { + let mut groups: Vec<(K, Vec)> = Vec::new(); + let mut index: HashMap<&K, usize> = HashMap::new(); + for (i, k) in keys.iter().enumerate() { + let g = *index.entry(k).or_insert_with(|| { + groups.push((k.clone(), Vec::new())); + groups.len() - 1 + }); + groups[g].1.push(i); + } + groups.into_iter().map(|(_, idxs)| idxs).collect() +} + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn batch_slots_groups_by_first_occurrence() { + assert_eq!( + batch_slots(&[1u32, 2, 1, 3, 2]), + vec![vec![0, 2], vec![1, 4], vec![3]] + ); + } + + #[test] + fn batch_slots_single_group() { + assert_eq!(batch_slots(&[7u32, 7, 7]), vec![vec![0, 1, 2]]); + } + + #[test] + fn batch_slots_all_distinct() { + assert_eq!(batch_slots(&[1u32, 2, 3]), vec![vec![0], vec![1], vec![2]]); + } + + #[test] + fn batch_slots_empty() { + assert_eq!(batch_slots::(&[]), Vec::>::new()); + } + + #[test] + fn batch_slots_each_index_exactly_once() { + let keys: Vec = (0..50).map(|i| i % 4).collect(); + let groups = batch_slots(&keys); + let mut all: Vec = groups.iter().flatten().copied().collect(); + all.sort(); + assert_eq!(all, (0..50).collect::>()); + } + + #[test] + fn batch_slots_deterministic_repeated() { + let keys: Vec = vec![2, 0, 1, 2, 0, 1, 3]; + assert_eq!(batch_slots(&keys), batch_slots(&keys)); + } +} diff --git a/lib/src/shaders/gpu_driven.wgsl b/lib/src/shaders/gpu_driven.wgsl index ed6cd9b..a3d07d5 100644 --- a/lib/src/shaders/gpu_driven.wgsl +++ b/lib/src/shaders/gpu_driven.wgsl @@ -8,8 +8,9 @@ // 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. +// All buffers are fixed-capacity (MAX_ENTITIES = 256, see ARCHI_CPU_GPU.md D12) and are +// allocated once. Each frame the CPU rewrites the transform slots and 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).