GPU culling

docs/DRAFT.md

@@ -1,24 +1,404 @@
-# Prochaine étape
+# DRAFT — Étape 17 : Rendu GPU-driven (ROADMAP 3.1 / 3.2 / 3.3)
 
-> Étape 16 (Phase 5 — Documentation & Polish) **terminée** le 2026-07-19.
+> **Statut** : **implémenté et validé** (voir la section « Acceptation » plus bas).
+> Conventions : cette étape couvre les 3 items de la Phase 3 de `docs/ROADMAP.md`.
+> ROADMAP / README / docs user sont mis à jour ; ce DRAFT est conservé comme référence de
+> conception (les décisions D1–D14 y sont documentées).
 >
-> **Tuning caméra 2026-07-19** — sur rétroaction utilisateur (caméra trop sensible, orbit permanent) :
-> (1) `CameraController` a gagné deux champs configurables `orbit_sensitivity`/`zoom_factor` (defaults : 0.01→0.005
-> rad/px, 0.9/tic) ; (2) `InputState` normalise `PixelDelta`/32 en unités « notch » (le Wayland renvoyait ~100 px/tic,
-> donc `0.9^100` → zoom au clamping en un geste) ; (3) le `demo` orbite désormais **sur clic gauche enfoncé** (arc-rotate
-> classique). Docs `camera-input.md`/`examples.md` mises à jour. 51 tests OK.
+> **Note de validation runtime (correction D12)** : la capacité est passée de **4096 à 256**
+> entités, et le slot de matrice mondes est **padded de 64 à 256 o**. Le buffer de matrices est
+> lié au slot `uniform` « object » du pipeline de rendu, et WebGPU impose **deux** contraintes :
+> (1) une binding `uniform` unique est plafonnée à `max_uniform_buffer_binding_size` (64 ko), et
+> (2) un offset de buffer `uniform` doit être un multiple de `min_uniform_buffer_offset_alignment`
+> (256 o). Une matrice de 64 o ne peut donc jamais être adressée individuellement par un offset
+> dynamique `uniform` : chaque slot de matrice est **padded à 256 o** (`MatSlot { m, pad }`), et
+> 256 slots × 256 o = 64 ko est le maximum adressable. Par ailleurs, le bind group « object » lie
+> une **slice de 64 o** (une matrice) plutôt que le buffer entier — une binding sur tout le buffer
+> plafonnerait l'offset dynamique à 0. Les mentions « 4096 » / « 1024 » ci-dessous renvoient au
+> plan initial ; la valeur réelle (et les tailles de buffer dérivées) est 256 (cf. D2 / D4 / D12).
 >
-> **Bug fix 2026-07-19** — `InputState::begin_frame()` remettait à zéro `mouse_delta`/`scroll` AVANT que
-> `update()` ne les lise, alors que les événements winit (CursorMoved/MouseWheel) s'accumulent ENTRE deux
-> frames. Résultat : la caméra orbitale du `demo` ne bougeait jamais (souris/molette toujours (0,0) dans
-> `update`). Corrigé par rotation des accumulateurs (`frame_mouse_delta`/`frame_scroll` → queryables à
-> `begin_frame`), aligné sur le pattern clavier/boutons. Tests mis à jour (ordre réel winit) + 45/50 tests OK.
->
-> Traduction anglaise de toute la documentation (hors `docs/tech/`, DRAFT/PLAN/ROADMAP) **terminée** le 2026-07-19 :
-> `docs/user/*`, `README.md`, READMEs de modules, doc/rustdoc de tous les `.rs` (src + examples + tests),
-> `Étape`→`Step` global. Vérifications : 50 tests OK, `cargo fmt` clean, aucun lien cassé, 0 accent restant hors zone franche.
+> **Note du 2026-09-22 (D14, correction post-implémentation)** : le `demo` (culling activé) affichait
+> une **fenêtre noire** — readback GPU : tous les comptes de draw args étaient à 0 alors que les
+> transforms, bboxes et plans de frustum vus par le GPU étaient corrects. Cause racine : l'ordre des
+> arguments de `select` en WGSL (`select(reject, accept, cond)` renvoie le **second** argument quand
+> `cond` est vrai — l'inverse de la convention HLSL). Corrigé et vérifié par readback (cf. D14).
 
-## Prochaines options
-- **Phase 3 — GPU-driven rendering** (ROADMAP 3.1/3.2/3.3) : indirect draw, buffers de paramètres GPU, culling GPU. C'est le gros morceau performance qui reste.
-- **Phase 4.4 — Performance** : LOD, instancing/multi-instancing, occlusion culling, batching par matériau (4.3).
-- Multi-caméras (`scene.set_active_camera`) — restant de la Phase 2.1.
+## Objectifs
+
+1. **3.1 — Matrices mondes sur le GPU** : un compute shader dérive la matrice monde de
+   chaque entité à partir de ses données de transform (T·R·S), au lieu d'un `to_matrix()`
+   CPU par entité chaque frame.
+2. **3.2 — Culling GPU** : un compute shader teste la visibilité de chaque entité
+   (approximation sphère vs les 6 plans du frustum) et écrit un slot d'arguments de draw.
+3. **3.3 — Draw indirect** : le rendu de la scène passe en `draw_indexed_indirect` /
+   `draw_indirect` par slot d'entité, piloté par les arguments produits par le culling.
+
+## Contraintes & principes
+
+- **API publique stable** : les exemples existants (cube, simple, shadow_test, spot_test,
+  manual) restent fonctionnels **sans modification de leurs appels**. En pratique aucun n'est
+  touché (cf. D11) ; seul `demo.rs` change (activation du culling, cf. D8).
+- **Aucune régression visuelle** : le culling est **désactivé par défaut** (D8). Le chemin
+  GPU-driven (matrices + draw indirect) est actif pour `render_scene` mais produit un rendu
+  **identique** au chemin CPU actuel.
+- **`standard_shader.wgsl` non modifié** (D2) : le binding group 1 reste
+  `var<uniform> object: ObjectUniform`. Seuls les commentaires en français sont traduits en
+  anglais (convention : toute la doc est en anglais).
+
+## Décisions (à valider)
+
+### D1 — Un draw indirect par entité (pas de multi-instancing par groupe de mesh)
+
+Chaque entité possède son **slot GPU** (index stable `0..N`). Un draw indirect par entité lit
+les arguments de son slot.
+
+- ✅ Couvre les 3 items de la roadmap : matrices mondes GPU (3.1), culling GPU (3.2),
+  draw indirect (3.3).
+- ✅ API publique inchangée ; les entités hétérogènes (meshs différents) restent trivialement
+  supportées.
+- ❌ Coût CPU : 1 `draw_indirect` par entité (vs 1 draw multi-instanced par groupe de mesh).
+  Pour la scale (milliers d'instances d'un même mesh), le multi-instancing par groupe est plus
+  efficace → **reporté en Phase 4.4** (instancing / multi-instancing).
+- **Capacité fixe : 4096 entités** (`MAX_GPU_ENTITIES`). `Scene::add_entity` renvoie `Err`
+  au-delà. Les slots sont **append-only avec tombstones** : la suppression ne décale pas les
+  indices (stabilité GPU) ; l'entité est marquée inactive (flag dans le slot de transform) et
+  le culling met ses arguments à 0.
+
+### D2 — Le bind group « object » par entité devient une slice du buffer GPU-computé
+
+- Le binding group layout existant (group 1, `var<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).
+
+### D3 — Layout des slots d'indirect draw : 80 o (pdc(16,20))
+
+- `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)
+
+```
+[CPU] write_buffer TransformBuffer  (64 o × num_slots, depuis Scene)
+[CPU] write_buffer CullUniforms     (6 plans + num_slots + culling)
+[CPU] (si gpu_generation changée) write_buffer BBoxBuffer
+[Compute 1] compute_matrices : m[i] = T·R·S          (i in 0..num_slots)
+[Compute 2] cull           : draws[i] = visible ? args : 0   (sphère vs 6 plans)
+[Render pass ombre]  (si ombres activées) : draw_indexed_indirect par caster (slice mat + args slot)
+[Render pass main]                      : draw_indexed_indirect / draw_indirect par slot vivant
+[Queue] submit(encoder)
+```
+
+Même encoder → exécution séquentielle garantie (write → compute → render).
+
+## Layouts GPU (WGSL)
+
+```wgsl
+struct TransformSlot { translation: vec3f, flags: vec4f, rotation: vec4f, scale: vec3f } // 64 o
+  // flags.x = mesh_index (index dans BBoxBuffer), flags.y = active (1/0),
+  // flags.z = draw count (compte de sommets/indices du mesh), flags.w = has_index (1/0)
+struct MatSlot  { m: mat4x4<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`.
+
+## Fichiers touchés
+
+| Fichier | Action |
+|---------|--------|
+| `lib/assets/shaders/gpu_driven.wgsl` | **Nouveau** — compute `compute_matrices` + `cull` |
+| `lib/src/math/frustum.rs` | **Nouveau** — extraction des 6 plans (Gribb-Hartmann [0,1]) + tests |
+| `lib/src/math/geometry.rs` | + struct `BBox` + `Geometry::bbox()` + tests |
+| `lib/src/math/mod.rs` | + `pub mod frustum;` |
+| `lib/src/resources/uniform.rs` | + `TransformSlot`, `BBoxSlot`, `CullUniforms` (Pod) + constantes `MAX_GPU_ENTITIES`, `MAX_MESH_BBOXES` + tailles de slots |
+| `lib/src/resources/mesh.rs` | + champ `bbox: BBox` (calculé dans `from_geometry`) + accès `mesh.bbox()` |
+| `lib/src/scene/scene.rs` | `entities` → slots append-only + index + génération ; `create_mesh`/`add_mesh` capturent le bbox ; + `iter_entity_slots()`, `gpu_generation()`, `gpu_bbox_slots()` |
+| `lib/src/core/renderer.rs` | + 4 buffers GPU + buffer cull-uniforms + 2 pipelines compute + layout compute ; `render_scene` → `&mut self`, passe indirect ; + `set_culling`/`culling_enabled` ; shadow pass indirect |
+| `lib/src/pipeline/pipeline_cache.rs` | + `create_compute_pipelines()` (module gpu_driven, 2 entry points) + `create_compute_bind_group_layout()` |
+| `lib/src/app.rs` | + `AppBuilder::with_culling(bool)` + champ `culling` ; `resumed()` applique `renderer.set_culling` |
+| `lib/src/utils/conf.rs` | + fallback embarqué `include_str!` de `gpu_driven.wgsl` (comme standard/shadow) |
+| `lib/examples/demo.rs` | + `AppBuilder::with_culling(true)` (démo du culling) |
+| `lib/tests/wgsl_validate.rs` | + test `gpu_driven_shader_is_valid_wgsl` |
+| `lib/src/assets/shaders/standard_shader.wgsl`, `shadow_shader.wgsl` | Commentaires FR → EN (convention doc anglaise) |
+| `docs/user/gpu-driven.md` | **Nouveau** — doc user (activation du culling, limites, comportement) |
+| `docs/user/README.md` | + lien vers `gpu-driven.md` |
+| `README.md` | + ligne « rendu GPU-driven (indirect draw, culling GPU) » dans les features |
+| `docs/ROADMAP.md` | Coche 3.1, 3.2, 3.3 |
+| `docs/PLAN.md` | + ligne Étape 17 |
+
+## Périmètre exclus (reportés)
+
+- **Multi-instancing par groupe de mesh** (Phase 4.4) — 1 draw par groupe de mesh identique.
+- **AABB exact transformé** (8 sommets) — v1 utilise une sphère conservative.
+- **Ring-buffering** des buffers GPU (D4) — v1 fait un `write_buffer` par frame.
+- **Occlusion culling**, **LOD**, **batching par matériau** (Phase 4.3 / 4.4).
+- **Multi-caméras** (2.1 restant).
+
+## Risques & mitigations
+
+| Risque | Mitigation |
+|--------|-----------|
+| **Alignement indirect draw** (offset multiple de 16/20/8/4 o) | D3 : slot 80 o = pdc(16,20), multiple de 8 et 4 → valide pour les 2 variants de draw. |
+| **Buffer undefined à la création** | `IndirectArgsBuffer` zéro-initialisé (`create_buffer_init` zéros) → `c..e` restent 0. |
+| **Ordre compute → render** | Même `CommandEncoder` → séquentiel. `write_buffer` + compute + render dans le même encoder. |
+| **Bind groups par slot : mémoire** | 4096 bind groups ≈ trivial (état driver partagé via le layout). |
+| **Bug de culling (objet qui disparaît)** | D8 : OFF par défaut ; le `demo` l'active → bug visible immédiatement. Sphère conservative (D5). **Réalisé en 2026-09-22 (D14)** : le bug s'est produit (fenêtre noire) et a été traçable grâce à cette mitigation + readback. |
+| **Builtins WGSL aux conventions d'arguments non intuitives** (ex. `select`, l'inverse de HLSL) | D14 : piège documenté en tête du shader + dans `AGENTS.md` ; diagnostic par **readback des slots** (`debug_dump`) plutôt que par l'image seule. |
+| **Ordre non-déterministe des entités** | D9 : `Vec` append-only (stable) remplace le `HashMap` (ordre aléatoire). |
+| **`render_scene` → `&mut self`** | Aucun exemple n'appelle `Renderer::render_scene` directement (cf. D11) → non-régression. |
+| **WGSL compute non validé** | Test `wgsl_validate` sur le nouveau shader + `cargo build` (wgpu compile à l'exécution). |
+| **Shadow pass indirect** | Les casters utilisent les mêmes slots d'args → cohérent ; testé dans le `demo` (ombres + culling actifs). |
+| **Alignement WGSL ≠ Rust** | `#[repr(align(16))]` + static asserts de taille sur chaque struct Pod (cf. section « Layouts GPU »). |
+
+## Plan de vérification
+
+1. `cargo build --workspace` — OK.
+2. `cargo test --workspace` — 45 tests unitaires + **nouveaux tests frustum/geometry** +
+   2→3 tests WGSL + 3 doctests, tous verts.
+3. `cargo fmt --all -- --check` — clean.
+4. **`demo` en headless** : `WGPU_BACKEND=vulkan timeout 10 cargo run -p examples --example
+   demo` → exit 0, **culling actif** (`with_culling(true)`), pas d'objet qui disparaît.
+5. **Tous les autres exemples** (cube, simple, shadow_test, spot_test) en headless → exit 0,
+   **rendu identique** (culling OFF par défaut, chemin indirect actif).
+6. **`manual.rs`** en headless → exit 0 (chemin `render()` bas niveau inchangé, UBO CPU).
+7. **Check liens** : 0 lien cassé (nouveaux `docs/user/gpu-driven.md` + liens README).
+8. **Aucun caractère accentué** dans les fichiers modifiés (sauf `docs/tech/`, DRAFT, PLAN,
+   ROADMAP, DOCUMENTATION) — les commentaires WGSL traduits.
+9. **WGSL valide** : le nouveau `gpu_driven.wgsl` compile (test `wgsl_validate`).
+
+## Critères d'acceptation (definition of done)
+
+> **Implémenté le 2026-07-20.** Tous les critères sont remplis (avec les écarts de nommage
+> notés ci-dessous). `cargo build`/`cargo test --workspace` : 56 lib + 3 WGSL + 3 doctests verts,
+> clippy sans avertissement dans `renderer.rs`.
+
+- [x] `compute_matrices` + `cull` dans `gpu_driven.wgsl` (2 entry points, 1 module).
+      *Écart : layout compute explicite à 3 groupes partagé par les 2 pipelines (évite les gaps de
+      groupes du layout inféré) — cf. D11.*
+- [x] 4 buffers GPU (transform, matrices, bboxes, draw-args) + 1 buffer cull-uniforms dans le `Renderer`.
+- [x] `render_scene` 100 % indirect (main + shadow).
+      *Écart : `render_scene` reste `&self` (buffers persistants créés dans `new()` + `Cell<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.

docs/PLAN.md

@@ -63,7 +63,7 @@ Une fois la plomberie encapsulée, nous devons rendre l'assemblage des objets co
 ### Gestion des Matériaux et Shaders
 
 - [X] S'assurer que chaque Mesh possède une référence vers un Material (à l'heure actuelle le lien est porté par l'entité `(mesh_id, material_id)` de la Scene, pas par le Mesh lui-même). *(fait — 2026-09-17, DRAFT Étape 7 : `Mesh.material: Option<Arc<Material>>` ; `Entity { mesh_id, transform }`, plus de `material_id`)*
-- [X] Implémenter le comportement par défaut : si aucun matériau n'est assigné, le moteur injecte automatiquement le `standard_shader` (variante unlit) (non implémenté). *(fait — 2026-09-17, DRAFT Étape 7.3.5 : `Scene::default_material()` injecte `standard` ; le flat reste piloté par `Renderer::set_unlit`)*
+- [X] Implémenter le comportement par défaut : si aucun matériau n'est assigné, le moteur injecte automatiquement le `standard_shader` (variante unlit). *(fait — 2026-09-17, DRAFT Étape 7.3.5 : `Scene::default_material()` injecte `standard` ; le flat reste piloté par `Renderer::set_unlit`)*
 
 ## Phase 3 : Documentation et Interface (API "User-Friendly")
 

docs/ROADMAP.md

@@ -117,24 +117,25 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z }
 
 ---
 
-## Phase 3️⃣ — GPU-Driven Rendering
+## Phase 3️⃣ — GPU-Driven Rendering ✅ (2026-07-20)
 
 **Objectif** : Déléguer les calculs de transformation et culling au GPU (suivre ARCHI_CPU_GPU.md).
+**Statut** : implémenté (DRAFT Étape 17, décisions D1–D14). Culling **désactivé par défaut** (non-régression), opt-in `AppBuilder::with_culling(true)`. **Correction 2026-09-22** : bug « fenêtre noire » avec culling ON (arguments de `select` WGSL écrits à la convention HLSL — toutes les entités visibles étaient remises à 0) ; corrigé et vérifié par readback GPU (DRAFT D14).
 
 ### 3.1 Compute Shader
-- [ ] Buffer `TransformBuffer` (CPU → GPU) : positions/rotations/échelles brutes
-- [ ] Buffer `MatrixBuffer` (GPU calculé) : World Matrices finales
-- [ ] Compute shader : calcul des World Matrices pour tous les meshes
+- [x] Buffer `TransformBuffer` (CPU → GPU) : positions/rotations/échelles brutes (`TransformSlot`, 64 B)
+- [x] Buffer `MatrixBuffer` (GPU calculé) : World Matrices finales (`MatSlot`, `STORAGE|UNIFORM`)
+- [x] Compute shader : calcul des World Matrices pour tous les meshes (`compute_matrices`)
 
 ### 3.2 Frustum Culling GPU
-- [ ] Ajouter `BBox` dans `Geometry` (center + extents)
-- [ ] Buffer `BoundingBoxBuffer` (CPU → GPU, statique)
-- [ ] Compute shader : culling basé sur la frustum de caméra
-- [ ] Buffer `IndirectDrawBuffer` rempli par le GPU
+- [x] Ajouter `BBox` dans `Geometry` (coins min/max locaux) + `math::Frustum` (Gribb–Hartmann `[0,1]`)
+- [x] Buffer `BoundingBoxBuffer` (CPU → GPU, ré-upload quand l'ensemble des meshes change, peu coûteux)
+- [x] Compute shader : culling sphère vs frustum (`cull`), **désactivé par défaut** *(bug « fenêtre noire » corrigé le 2026-09-22 — ordre des arguments de `select` WGSL inversé ; cf. DRAFT D14)*
+- [x] Buffer `IndirectDrawBuffer` rempli par le GPU (`DrawSlot`, 80 B, zéro = no-op)
 
 ### 3.3 Rendu Indirect
-- [ ] `draw_indexed_indirect()` au lieu de draw calls individuels
-- [ ] Un seul command draw pour tous les objets visibles
+- [x] `draw_indexed_indirect()`/`draw_indirect()` au lieu de draw calls individuels
+- [x] Un draw indirect **par slot actif** (décision D1 — pas un draw fusionné unique) ; le shadow pass est aussi indirect
 
 ---
 
@@ -198,3 +199,4 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z }
 | **String IDs pour le MVP, slotmap reporté** | Le code et le README utilisent des String IDs (simples, sûrs, figés avant la boucle de rendu) ; `ARCHI_ARENES.md` reste la cible "handles typés" pour plus tard. La dépendance `slotmap` a été retirée tant qu'elle est inutilisée |
 | **Present mode FIFO figé pour l'instant** | Le swapchain utilise `PresentMode::Fifo` avec `desired_maximum_frame_latency: 2` (double buffering vsync) — défaut sûr : pas de tearing, énergie minimale, zéro artefact. On **gèle ce choix** ; `Mailbox` (triple buffering) pourra être exposé en option et `Immediate` restera réservé à l'offscreen, **on s'occupera du present mode le moment venu** (quand le pipeline GPU-driven arrivera, Phase 3) — ce n'est pas bloquant pour les étapes 1-2 |
 | **Resize géré (avec recréation de la depth texture), acté en D3 (2026-09-18), réalisé en Étape 11 (2026-09-18)** | L'app reconfigure désormais la surface et recrée la depth texture **en même temps** à chaque `Resized` (`App::resize` → `Context::configure` + `Renderer::resize_depth`), via le helper `create_depth_texture` isolé. Vérifié au runtime (exemple `cube`) : pas de crash, pas d'artefact, aspect correct |
+| **WGSL `select(reject, accept, cond)`** | L'ordre des arguments est l'inverse de la convention HLSL : le **second** argument est retenu quand la condition est vraie. L'avoir écrit à la convention HLSL a produit le bug « fenêtre noire » du culling (comptes remis à 0 pour les entités visibles), corrigé le 2026-09-22 (DRAFT D14). Piège documenté en tête de `gpu_driven.wgsl` + `AGENTS.md` |

docs/tech/ARCHI_APP.md

@@ -15,17 +15,20 @@ stale_after: 2027-01-31
 
 wsg_lib est un moteur de rendu modulaire basé sur wgpu. Il adopte une architecture à deux niveaux : une façade de haut niveau pour la productivité et un accès bas niveau pour un contrôle total.
 
-> **État du document : ACTUEL pour la façade (`App`/`AppHandler`, §3, §4A) ; CIBLE pour la partie
-> GPU-driven (§1, §4B, §5, §6).** La façade `AppBuilder`/`App`/`AppHandler` est livrée et est le
-> **workflow recommandé** : `setup` (déclaration de la scène) → par frame `update` (mutation) →
+> **État du document : ACTUEL** — façade (`App`/`AppHandler`, §3, §4A) et pipeline GPU-driven
+> (§1, §4B, §5, §6) **implémenté en Phase 3** du ROADMAP (2026-07-20, DRAFT Étape 17, décisions
+> D1–D14). La façade `AppBuilder`/`App`/`AppHandler` est livrée et est le **workflow recommandé** :
+> `setup` (déclaration de la scène) → par frame `update` (mutation) →
 > `render` (défaut : `App::render_scene` = itération des entités + **rendu groupé en une passe**,
 > un `CommandEncoder`/soumission par frame ; passe d'ombre en tête si un caster est actif).
 > Exemples : `simple` (2D unlit), `cube` (3D éclairé), `demo` (vitrine : primitives, lumières,
-> ombres, caméra orbitale). Le workflow **manuel** (exemple `manual`) coexiste pour le contrôle fin.
-> Les sections §1, §4B, §5 et §6 décrivent la **cible** : pipeline GPU-driven à deux passes
-> (Compute Pass → `draw_indexed_indirect`), buffers persistants en VRAM (Transform/Matrix/BBox/Indirect)
-> et synchronisation single/double buffer. **Rien de tout cela n'existe encore dans le code** — c'est
-> la trajectoire ROADMAP Phase 3.
+> ombres, caméra orbitale, culling GPU activé). Le workflow **manuel** (exemple `manual`) coexiste
+> pour le contrôle fin.
+> Les sections §1, §4B, §5 et §6 décrivent le pipeline GPU-driven **tel qu'implémenté**, avec les
+> écarts documentés (DRAFT Étape 17) : un draw indirect par slot (D1), table fixe de 256 slots
+> (D12), culling par sphère conservative (D5), single buffer (D4), et le piège de l'ordre des
+> arguments de `select` en WGSL (D14, bug « fenêtre noire » corrigé le 2026-09-22). La section
+> « Notes pour l'implémentation future » (double buffering) reste **CIBLE**.
 
 ## 1. Philosophie et Principes
 
@@ -107,8 +110,8 @@ Le moteur gère la renderloop interne via un pipeline à **deux passes séquenti
 
 1. **Update** (`AppHandler::update`) — L'utilisateur modifie la scène (transformations, entités). Ces changements sont synchronisés vers le GPU via un **single buffer** Transform avant la passe de calcul.
    > La synchronisation est assurée par le pipeline wgpu : `queue.submit()` après le compute pass garantit que les données Transform sont valides avant le render pass suivant. Aucun double buffering n'est nécessaire tant que la latence maximale de la surface (via `desired_maximum_frame_latency`) est ≥ 3.
-2. **Compute Pass** — Un compute shader lit les Transform bruts, calcule les World Matrices finales, effectue le Frustum Culling par AABB, et remplit l'Indirect Draw Buffer avec les identifiants des objets visibles.
-3. **Render Pass** — Le CPU émet une unique commande `draw_indexed_indirect`. Le GPU pioche dans l'Indirect Draw Buffer et dessine uniquement les objets visibles, sans intervention du CPU.
+2. **Compute Pass** — Deux entry points compute séquentiels (`compute_matrices` puis `cull`, un seul module WGSL) lisent les Transform bruts, calculent les World Matrices finales, effectuent le Frustum Culling par **sphère conservative** (D5), et remplissent l'Indirect Draw Buffer avec les **comptes** de draw des objets visibles (0 si cullé/inactif).
+3. **Render Pass** — Le CPU émet **un draw indirect par slot** (écart D1 — la cible initiale prévoyait une commande unique fusionnée). Le GPU pioche les comptes dans l'Indirect Draw Buffer et dessine uniquement les objets non cullés et actifs, sans intervention du CPU.
 4. **Présentation** — La surface est présentée à l'écran.
 
 L'ordre d'appel des méthodes sur le `CommandEncoder` (`begin_compute_pass` puis `begin_render_pass`) garantit l'exécution séquentielle. Les barrières de mémoire entre passes sont insérées automatiquement par le pilote.
@@ -117,10 +120,11 @@ L'ordre d'appel des méthodes sur le `CommandEncoder` (`begin_compute_pass` puis
 
 | Buffer | Rôle | Type wGPU | Direction du flux |
 |--------|------|-----------|-------------------|
-| Transform Buffer | Positions/rotations/échelles brutes | Storage Buffer | CPU → GPU |
-| Matrix Buffer | World Matrices finales calculées | Storage Buffer | GPU (Calculé) → GPU (Lu par Render) |
-| Bounding Box Buffer | AABB de chaque mesh pour culling | Storage Buffer | CPU → GPU (Statique) |
-| Indirect Draw Buffer | Liste dynamique des objets à dessiner | Indirect + Storage | GPU (Rempli par Compute) → GPU (Lu par Render) |
+| Transform Buffer | Positions/rotations/échelles brutes + flags par entité (64 o/slot) | Storage Buffer | CPU → GPU (chaque frame, `write_buffer`) |
+| Matrix Buffer | World Matrices finales calculées (256 o/slot, padded — D12) | Storage + Uniform | GPU (Calculé) → GPU (Lu par Render) |
+| Bounding Box Buffer | Coins min/max de l'AABB de chaque mesh (32 o/slot) | Storage Buffer | CPU → GPU (quand l'ensemble des meshes change) |
+| Indirect Draw Buffer | Comptes de draw par slot (80 o/slot, zéro = no-op) | Indirect + Storage | GPU (Rempli par Compute) → GPU (Lu par Render) |
+| CullUniforms | 6 plans du frustum + `num_slots` + `culling` (112 o) | Uniform Buffer | CPU → GPU (chaque frame) — réservé au compute (group 2) |
 
 > **Synchronisation single buffer** : Les buffers Transform et Matrix utilisent un **single buffer** en phase initiale. Le CPU écrit dans le buffer pendant `update()`, puis le compute shader lit les données au frame suivant via `queue.submit()` qui garantit la séquence d'exécution. Cette approche fonctionne correctement tant que la surface a une latence maximale ≥ 2 frames (configuré via `desired_maximum_frame_latency`). Le double buffering sera ajouté uniquement si des artefacts visuels apparaissent à haute fréquence (typiquement > 90 fps sur machines rapides).
 

docs/tech/ARCHI_CPU_GPU.md

@@ -7,7 +7,7 @@ actor: person/jerome
 sources: []
 generated: { by: human:jerome, at: 2026-07-31T00:00:00Z }
 verified: true
-status: target
+status: current
 stale_after: 2027-01-31
 ---
 
@@ -16,13 +16,19 @@ Bonnes Pratiques & Guide d'Implémentation
 
 Ce document sert de spécification technique et de trame d'implémentation pour l'architecture de rendu 3D pilotée par le GPU (GPU-Driven Rendering) utilisant wgpu. L'objectif est de déléguer un maximum de charges de calcul au GPU pour soulager le CPU et maximiser les performances de parallélisme.
 
-> **État du document : CIBLE (spécification du pipeline GPU-driven, non implémenté).**
+> **État du document : ACTUEL (implémenté — Phase 3 du ROADMAP, 2026-07-20, DRAFT Étape 17).**
 > La répartition CPU/GPU, le compute pass (World Matrices + Frustum Culling), l'Indirect Draw Buffer
-> et les buffers persistants en VRAM décrits ici correspondent à la **Phase 3 du ROADMAP** et aux
-> README étapes 2-3. **Aucun de ces mécanismes n'existe encore dans le code.** Aujourd'hui le rendu est
-> piloté par le CPU, **objet par objet** (une soumission par mesh, voir README.md et l'exemple `manual`).
-> Considérez ce document comme la spécification de référence pour l'implémentation future du pipeline
-> GPU-driven, pas comme une description de l'état actuel.
+> et les buffers persistants en VRAM décrits ici sont en place : `shaders/gpu_driven.wgsl`
+> (deux entry points `compute_matrices` + `cull`, un module, layout explicite à 3 groupes) et les
+> buffers de slots du `Renderer` (`TransformSlot`/`MatSlot`/`BBoxSlot`/`DrawSlot`/`CullUniforms`,
+> capacité fixe de 256 slots).
+> **Écarts documentés** (cf. DRAFT Étape 17) : (D1) un draw indirect **par slot** plutôt qu'une
+> commande unique fusionnée ; (D12) 256 slots, slot matrice padded à 256 o (plafond `uniform` WebGPU) ;
+> (D5) culling par **sphère** conservative dérivée de l'AABB locale du mesh, pas par l'AABB transformée
+> exacte ; (D4) single buffer, pas de double-buffering.
+> **Piège connu (2026-09-22, D14)** : l'ordre des arguments de `select` en WGSL est l'inverse de la
+> convention HLSL — l'avoir inversé a produit un bug « fenêtre noire » (entités visibles remises à 0),
+> corrigé et vérifié par readback GPU. Documenté en tête de `gpu_driven.wgsl` et dans `AGENTS.md`.
 
 1. Répartition des Rôles : CPU vs GPU (La Source de Vérité)
 
@@ -53,10 +59,10 @@ L'exécution des tâches s'appuie sur une structure séquentielle stricte au sei
 - Mise à jour CPU (Minimaliste) : Le CPU écrit les transformations brutes (Transform) modifiées dans un buffer GPU mappé (single buffer en phase initiale — la synchronisation est assurée par `queue.submit()` qui garantit la séquence d'exécution). Double buffering sera ajouté uniquement si des artefacts apparaissent à haute fréquence (> 90 fps).
 - Pass de Calcul (Compute Pass) :
   - Calcul des World Matrices : Un compute shader lit les transformations brutes et génère la matrice 4x4 finale pour chaque mesh.
-  - Frustum Culling GPU : Le même compute shader (ou un compute pass dédié) compare la Bounding Box (AABB) de chaque objet avec les plans de la caméra (matrice de projection/vue).
-  - Remplissage du Buffer Indirect : Si l'objet est visible, son identifiant est injecté dans un buffer de commandes de dessin indirect (Indirect Draw Buffer).
+  - Frustum Culling GPU : Un compute pass dédié (`cull`) compare la **sphère bounding** de chaque objet (D5 — conservative, dérivée de l'AABB locale du mesh et de l'échelle de l'entité) avec les 6 plans du frustum de la caméra.
+  - Remplissage du Buffer Indirect : le pass `cull` écrit le **compte de sommets/indices** de chaque objet dans son `DrawSlot` (80 o) — mis à 0 si l'objet est cullé ou inactif (no-op).
 - Pass de Rendu (Render Pass) :
-  - Le CPU émet une unique commande globale : draw_indexed_indirect.
+  - Le CPU émet **un draw indirect par slot** (écart D1 — la spécification initiale prévoyait une commande unique fusionnée) ; les slots à compte 0 (cullés/inactifs/vides) sont des no-ops.
   - Le GPU pioche directement dans le buffer préparé par le compute pass et dessine uniquement les objets visibles, sans intervention du CPU.
   
 3. Stratégie de Synchronisation
@@ -66,12 +72,15 @@ L'exécution des tâches s'appuie sur une structure séquentielle stricte au sei
 
 4. Synthèse des Structures de Données en VRAM
 
-Pour implémenter cette architecture, prévoyez l'utilisation des buffers wGPU suivants :
-Nom du Buffer,Rôle,Type wGPU,Direction du flux
-Transform Buffer,Stocke les positions/rotations/échelles brutes.,Storage Buffer,CPU → GPU
-Matrix Buffer,Stocke les World Matrices finales calculées.,Storage Buffer,GPU (Calculé) → GPU (Lu par le Render)
-Bounding Box Buffer,Stocke les AABB de chaque mesh pour le culling.,Storage Buffer,CPU → GPU (Statique)
-Indirect Draw Buffer,Contient la liste dynamique des objets à dessiner.,Indirect Buffer + Storage,GPU (Rempli par Compute) → GPU (Lu par Render)
+L'implémentation utilise les buffers wGPU suivants (tous créés par le `Renderer` à l'initialisation, capacité fixe de 256 slots) :
+
+| Buffer | Rôle | Type wGPU | Direction du flux |
+|--------|------|-----------|-------------------|
+| Transform Buffer | Positions/rotations/échelles brutes + flags par entité (64 o/slot) | Storage Buffer | CPU → GPU (chaque frame, `write_buffer`) |
+| Matrix Buffer | World Matrices finales calculées (256 o/slot, padded — D12) | Storage + Uniform Buffer | GPU (Calculé) → GPU (Lu par le Render) |
+| Bounding Box Buffer | Coins min/max de l'AABB de chaque mesh (32 o/slot) | Storage Buffer | CPU → GPU (quand l'ensemble des meshes change) |
+| Indirect Draw Buffer | Comptes de draw par slot (80 o/slot, zéro = no-op) | Indirect + Storage Buffer | GPU (Rempli par Compute) → GPU (Lu par le Render) |
+| CullUniforms | 6 plans du frustum + `num_slots` + `culling` (112 o) | Uniform Buffer | CPU → GPU (chaque frame) — réservé au compute (group 2) |
 
 ## Liens
 

docs/tech/FRAME_LOOP.md

@@ -68,10 +68,11 @@ Avec notre nouvelle architecture "Atelier", la distinction est devenue encore pl
 | CommandEncoder | Par-Frame | Ton "carnet de notes" temporaire pour les ordres du GPU. |
 | TextureView | Par-Frame | Fenêtre temporaire sur la texture active du swapchain. |
 
-> **Ressources GPU persistantes (single buffer) — CIBLE, non implémenté** : À l'état **visé**, les
-> buffers Transform et Matrix vivent en VRAM avec un single buffer en phase initiale (le CPU écrit
-> pendant `update()`, le compute shader lit au frame suivant, séquencé par `queue.submit()`), puis un
-> double buffering si des artefacts apparaissent à haute fréquence. **Aucune de ces ressources n'existe
-> encore dans le code** — c'est la cible GPU-driven (ROADMAP Phase 3 / ARCHI_CPU_GPU).
+> **Ressources GPU persistantes (single buffer) — implémenté (Phase 3, 2026-07-20)** : les buffers
+> Transform, Matrix, BBox et Indirect Draw vivent en VRAM (créés à l'initialisation du `Renderer`,
+> capacité fixe de 256 slots). Le CPU écrit les transforms chaque frame par `queue.write_buffer`
+> **dans le même `CommandEncoder`** que les compute passes, qui les lisent **dans la même frame**
+> (l'ordre est garanti par l'encoder, pas par `queue.submit()` inter-frames). Le double buffering
+> reste la **cible** si des artefacts apparaissent à haute fréquence (voir ARCHI_CPU_GPU / ARCHI_APP).
 
 ---

docs/user/README.md

@@ -20,6 +20,7 @@ GPU graphics background is required.
 | [Materials & textures](materials.md) | Appearance: the `standard` shader, unlit mode, diffuse textures |
 | [Lights](lights.md) | Directional, point, spot, ambient, `MAX_LIGHTS` |
 | [Shadows](shadows.md) | Shadow mapping: picking the casting light, the packed-index pitfall |
+| [GPU-driven rendering](gpu-driven.md) | GPU world matrices + indirect draws, opt-in frustum culling |
 | [Camera & input](camera-input.md) | Active camera, orbital controller, unified keyboard/mouse state |
 | [Examples](examples.md) | The 7 repo examples, the advanced `manual` workflow, adding your own example |
 

docs/user/gpu-driven.md

@@ -0,0 +1,132 @@
+# GPU-driven rendering
+
+WSG's scene rendering is **GPU-driven**: the per-entity world matrices and the indirect draw
+arguments are computed on the GPU each frame, so the CPU no longer loops over entities to issue
+draw calls. This page explains what that means for you and how to opt into **frustum culling**.
+
+## What runs on the GPU
+
+Each frame, before the render passes, two compute passes run over a fixed-capacity slot table
+(256 entities, allocated once):
+
+1. **`compute_matrices`** derives each entity's world matrix from its transform
+   (translation / rotation / scale). The result feeds the render pipelines as the per-entity
+   model matrix.
+2. **`cull`** decides per-entity visibility and fills the **indirect draw arguments** (the
+   vertex/index count, zeroed when the entity is culled or inactive).
+
+The main and shadow render passes are then **100 % indirect**: each active slot issues one
+indirect draw that reads its own count and world matrix. A culled or inactive slot has a zero
+count, so its draw is a no-op. The CPU only rewrites the transform slots and the cull uniforms
+each frame — it never iterates the entities to issue draws.
+
+You do not need to do anything special to get this: `render_scene` is GPU-driven by default.
+
+## Frustum culling (opt-in)
+
+Culling is **off by default**. The culling pass still runs, but with culling disabled it marks
+every active entity visible — so the rendered image is **identical** to a CPU-culled scene.
+This protects you from a culling bug (an object that should be visible vanishing) becoming a
+silent correctness issue.
+
+To enable culling, build your `App` with `.with_culling(true)`:
+
+```rust
+let app = AppBuilder::new()
+    .title("My app")
+    .with_culling(true)   // skip entities whose bounding sphere leaves the frustum
+    .build()
+    .await?;
+```
+
+Or toggle it at runtime on the renderer:
+
+```rust
+app.renderer().set_culling(true);   // enable
+app.renderer().set_culling(false);  // disable again
+```
+
+## How culling works
+
+When culling is on, each entity's **local-axis-aligned bounding box** (computed once from its
+geometry, `Geometry::bbox()`) is treated as a **bounding sphere**:
+
+- **center** = the box center, transformed by the entity's world transform (rotation +
+  translation; scale is folded into the radius),
+- **radius** = the box's circumradius scaled by the entity's largest scale component.
+
+The sphere is tested against the six camera frustum planes. If it is **fully outside** (beyond
+a plane by more than its radius), the entity is culled; otherwise it is drawn.
+
+The sphere is a **conservative** approximation of the box: it can draw an object that is partly
+out of view (false negative), but it will **never cull an object that is actually visible**
+(false positive). For tight culling you would need per-mesh sphere fitting or per-face tests,
+which are out of scope for v1.
+
+## Debugging the GPU path
+
+If something looks wrong — a missing object, a black window — the GPU-side slot tables can be
+read back and printed. The `Renderer` ships a debug helper (intentionally **not** part of the
+documented API):
+
+```rust
+app.renderer().debug_dump(8);   // prints the first 8 GPU slots to stderr
+```
+
+It dumps exactly what the GPU sees: the transform slots, the derived world matrices, the
+indirect draw arguments, the mesh bounding boxes and the cull uniforms. A slot whose vertex
+count reads `0` was zeroed by the cull pass (culled, inactive, or beyond `num_slots`); a full
+count means the entity is drawn. In the `demo` example the dump is opt-in via an environment
+variable, so the showcase stays silent by default:
+
+```sh
+WSG_DEBUG_DUMP=120 cargo run -p wsg-lib --example demo
+```
+
+`WSG_DEBUG_DUMP=N` dumps for the first *N* frames. The demo stays **silent** when the variable is
+unset; a set-but-non-numeric value (e.g. `WSG_DEBUG_DUMP=on`) gives 3 frames.
+
+**How to verify culling is actually working** (a correct culling pass is invisible — culled
+objects were off-screen anyway — so the proof is in the counts, not the image):
+
+1. Launch the demo with `WSG_DEBUG_DUMP=120` (the demo has culling **on** and an orbiting
+   camera — drag the mouse to orbit).
+2. Note first that orbiting/zooming this camera **cannot cull the entity ring**: the camera
+   always looks at the origin, so each entity's angular offset from the view axis is bounded
+   by `atan(ring radius / camera distance)` = `atan(1.7/6.1)` ≈ 15.5°, under the ~22° vertical
+   half-FOV. The seven demo entities therefore keep their **full** counts (cube `36`, sphere
+   `3840`, …) in every orientation — that is the expected and correct behaviour (verified
+   2026-09-22: 600-frame camera sweep, GPU cull verdicts matched an independent CPU sphere
+   test on all 6000 entity frames, zero flips on the ring).
+3. To see the counts actually flip to **`0`**, you need an entity well **off the target axis**
+   — e.g. one placed far away so it ends up behind the near plane. Its count then toggles
+   `0` ↔ full as the camera orbits, while the on-axis entities stay full. (This off-axis test
+   is the one that verified the cull path end-to-end, positive and negative.)
+4. Optional A/B: temporarily build with `.with_culling(false)` and repeat — with culling off,
+   every entity keeps its full count in **every** orientation (the off-axis one included).
+
+This readback is the reference truth when a shader bug is suspected: it shows both the computed
+counts and the raw inputs of the cull pass, independently of what ends up on screen. (It is how
+the 2026-09-22 « black window » bug — an inverted WGSL `select` argument order — was diagnosed
+and verified fixed, see the DRAFT D14 in `docs/DRAFT.md`.)
+
+## Limitations
+
+- **Culling is all-or-nothing per entity.** There is no partial (per-triangle) culling.
+- **The sphere is a coarse bound** for elongated meshes (a long thin box gets a large sphere).
+  If your scene is dominated by such shapes, culling may bring little gain.
+- **Capacity is 256 entities per render pass.** Beyond that, extra entities are not drawn.
+  This is the largest a single-buffer design can address under WebGPU's two `uniform` rules: a
+  single `uniform` binding is capped at 64 KB, *and* a `uniform` offset must be a multiple of 256 B.
+  A 64-byte matrix can never be individually addressable by a `uniform` offset, so each matrix
+  slot is padded to 256 B — and 256 slots × 256 B = 64 KB is the maximum. It is amply generous
+  for a simple scene (the demo has 7).
+- **Mesh bounding boxes are recomputed when meshes are added**; a scene whose mesh set changes
+  at runtime simply re-uploads the small bbox table (a few bytes per mesh).
+
+Culling is a **performance** feature, not a visual one: with it off you get the same image with
+the indirect-draw machinery still active.
+
+---
+
+Next: [Examples](examples.md) · Back to [User documentation index](README.md)