material batching
This commit is contained in:
Jérôme Bousquié
+149
-363
@@ -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<uniform> 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<Material>` 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<clé, index>` + `Vec<groupe>` :
|
||||
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<uniform> 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<K: Eq + Hash>(keys: &[K]) -> Vec<Vec<usize>>`
|
||||
(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<u32>`. 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<f32>` (normale + d), normalisé.
|
||||
- Uploadé dans `CullUniformsBuffer` (group **2** du compute, binding 0, `var<uniform>`) 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<Option<Entity>>` append-only + index label→slot
|
||||
|
||||
- `Scene::entities` passe de `HashMap<String, Entity>` à :
|
||||
- `entity_slots: Vec<Option<Entity>>` (append-only, tombstones)
|
||||
- `entity_labels: Vec<Option<String>>` (parallèle)
|
||||
- `entity_index: HashMap<String, usize>` (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<Mesh>, &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<vec4f,12> }` = 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<f32>, pad: array<vec4<f32>, 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<u32>, b: vec4<u32>, c: vec4<u32>, d: vec4<u32>, e: vec4<u32> } // 80 o
|
||||
struct CullUniforms {
|
||||
planes: array<vec4<f32>, 6>, num_slots: u32, culling: u32, _pad: vec2<u32> // 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<bool>`
|
||||
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.
|
||||
- [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).
|
||||
|
||||
+5
-5
@@ -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<Material>` 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` |
|
||||
|
||||
@@ -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**.
|
||||
|
||||
@@ -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é)
|
||||
|
||||
|
||||
@@ -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**
|
||||
|
||||
+15
-1
@@ -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
|
||||
|
||||
|
||||
Reference in New Issue
Block a user