LOD GPU
This commit is contained in:
Jérôme Bousquié
+288
-158
@@ -1,190 +1,320 @@
|
||||
# DRAFT — Étape 18 : Batching par Material (ROADMAP 4.3)
|
||||
# Étape 19 — LOD (Level of Detail) par entité
|
||||
|
||||
> **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) ».
|
||||
>
|
||||
> **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`).
|
||||
## Contexte
|
||||
|
||||
## Contexte — où on en est
|
||||
Le culling (Étape 17) supprime les objets **hors écran**. Le LOD supprime le travail **invisible**
|
||||
sur les objets qui sont sur l'écran : un maillage de 3840 indices qui n'occupe que 20×20 pixels
|
||||
à l'écran gaspille du GPU — presque tous ses triangles ne produisent aucun pixel.
|
||||
|
||||
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) :
|
||||
Principe : chaque mesh peut porter **plusieurs niveaux de géométrie** (L0 = détaillée, L1, L2 =
|
||||
de plus en plus simplifiée, silhouettes proches). Chaque frame, on mesure la **taille perçue**
|
||||
de chaque entité (rayon de sa bounding sphere projeté en pixels) et on choisit le niveau le plus
|
||||
grossier **suffisant**. La transition est protégée par une **hystérésis** (bande morte) pour
|
||||
éviter le scintillement d'un objet oscillant autour d'un seuil.
|
||||
|
||||
| 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 |
|
||||
ROADMAP : item 4.3 « Level of Detail (LOD) ». Le DRAFT de l'Étape 18 (batching par material) est
|
||||
remplacé par ce draft — il est validé et son contenu est dans `ARCHI_CPU_GPU.md` + git history.
|
||||
|
||||
→ 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.
|
||||
## La contrainte d'architecture qui tout détermine
|
||||
|
||||
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
|
||||
|
||||
- **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).
|
||||
`set_vertex_buffer` est un **état de passe posé par le CPU** — le GPU ne peut pas choisir entre
|
||||
deux buffers par draw indirect. Donc **plusieurs buffers par niveau est exclu** : tous les
|
||||
niveaux d'un mesh vivent dans **un seul buffer de vertices** (et un seul buffer d'indices), et
|
||||
le GPU sélectionne le niveau en écrivant l'**offset** (`baseVertex` / `firstVertex`) et le
|
||||
`count` dans les indirect draw args — des champs que le draw indirect lit déjà. C'est ce qui
|
||||
rend le LOD compatible avec l'architecture GPU-driven existante sans aucun nouveau mécanisme de
|
||||
rendu : le pass `cull` écrit déjà les draw args, il écrit juste les bonnes valeurs.
|
||||
|
||||
## Décisions
|
||||
|
||||
### 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).
|
||||
### D1 — Le CPU décide du niveau, le GPU mappe niveau → draw args
|
||||
|
||||
### 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.
|
||||
Le CPU calcule chaque frame le niveau par entité (rayon projeté + hystérésis) et l'uploade dans
|
||||
un petit buffer par slot ; le pass `cull` existe déjà et lit ce niveau pour choisir la ligne du
|
||||
tableau par mesh (count + offset).
|
||||
|
||||
### 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).
|
||||
Raisonnement :
|
||||
- **Testabilité** : la décision est une fonction pure Rust `lod_level(radius_px, last, max,
|
||||
thresholds)` → testable sans GPU, comme `batch_slots` (maison).
|
||||
- Le CPU **réécrit déjà** les transform slots chaque frame — un niveau de plus dans la même
|
||||
passe coûte ~1 flop multipli par entité.
|
||||
- Aucune extension de `CullUniforms` (pas de view-projection à ajouter) — le CPU a déjà
|
||||
caméra + viewport.
|
||||
- L'alternative (décision GPU : viewProj + hauteur dans les cull uniforms, état hystérésis GPU)
|
||||
est documentée comme **extension future** — elle n'apporte qu'un gain marginal (1 KB de
|
||||
upload/frame) au prix d'un test par readback.
|
||||
|
||||
### 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.
|
||||
### D2 — Un seul buffer packé par mesh, **niveau 0 à l'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.
|
||||
Les vertices de tous les niveaux sont concaténés dans le `vertex_buffer` du mesh (idem indices).
|
||||
Niveau 0 en tête : la voie basse `draw_entity` (`draw(0..num_vertices)`, sans LOD) et tout draw
|
||||
direct existant restent **inchangés** — ils lisent le début du buffer, qui est encore L0.
|
||||
|
||||
### 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.
|
||||
### D3 — Indices rebasés par niveau
|
||||
|
||||
## Mécanisme par frame (seul le point 7 change)
|
||||
Chaque niveau est une géométrie indépendante (indices 0-based sur ses propres vertices). Le
|
||||
packer décale les indices du niveau k de `vertex_offset(k)` ; le draw indirect du niveau k pose
|
||||
`baseVertex = vertex_offset(k)` → l'index buffer concaténé fonctionne tel quel. (Un niveau reste
|
||||
< 65 536 vertices, sinon erreur à l'ajout — `u16`.)
|
||||
|
||||
### D4 — Hystérésis **asymétrique**
|
||||
|
||||
Seuils descendants `t[0] > t[1] > …` (pixels) : `t[k]` = rayon **au-dessus** duquel le niveau
|
||||
k+1 est exigé (équiv. : le niveau k+1 est suffisant tant que `r ≤ t[k]`). Le niveau cible sans
|
||||
hystérésis est le **plus grossier** dont la borne est encore respectée (L0 n'a pas de borne).
|
||||
- Passer à un niveau **plus grossier** : seulement si `r ≤ borne × 0.8` (bande morte 20 %).
|
||||
- Passer à un niveau **plus fin** : immédiatement dès que `r` franchit la borne.
|
||||
Si un mesh a plus de niveaux que de seuils, les niveaux excédentaires partagent la dernière
|
||||
borne (clamp) — avec `[48, 12]` seuls les 3 premiers niveaux sont distincts.
|
||||
Le pop « perte de détail » (le plus visible) est retardé ; le pop « retour au détail » est
|
||||
immédiat. C'est la pratique standard des moteurs — et c'est testable en série de rayons
|
||||
oscillants autour d'un seuil.
|
||||
|
||||
### D5 — Deux nouveaux petits buffers, layout des slots existants intact
|
||||
|
||||
Les 4 composants de `TransformSlot.flags` sont tous occupés (x=mesh, y=active, z=count,
|
||||
w=has_index) — étendre le slot à 68 B ferait ripple dans tout le contrat documenté. À la place :
|
||||
- `lod_levels` : `array<u32>` par **slot** (256 × 4 B = 1 KB), re-uploadé chaque frame (comme
|
||||
les transforms). Nouveau `@group(2) @binding(3)` du pass `cull`.
|
||||
- `lod_tables` : par **mesh** (80 B = count + pad + 4 lignes de 16 B), re-uploadé chaque frame
|
||||
(mêmes motifs que le bbox buffer : petit, et le mapping mesh-index → tableau reste correct si
|
||||
des meshes/niveaux sont ajoutés à chaud).
|
||||
|
||||
`flags.z` (count plein) reste dans le slot (métadonne ; le GPU ne l'utilise plus pour la
|
||||
branche visible, qui passe par le tableau — voir D6).
|
||||
|
||||
### D6 — Les deux branches visibles écrivent depuis le tableau LOD
|
||||
|
||||
Dans `cull`, la branche « culling désactivé » et la branche « visible » remplacent
|
||||
`set_draw_count(i, flags.z)` par un helper `write_level_args(i, t)` : lire `lod_levels[i]`
|
||||
(clampé au count du mesh), indexer la ligne du mesh, écrire
|
||||
`.a = (count_ligne, 1, 0, baseVertex_ligne)` (indexé : `baseVertex = .a.w` ; non indexé :
|
||||
`firstVertex = .a.z`). Les branches zéro (slot ≥ num_slots, inactive) sont inchangées.
|
||||
Conséquence : un mesh **sans** LOD (1 niveau) a un tableau d'une seule ligne = count + offset 0
|
||||
→ comportement **bit-identique** à aujourd'hui.
|
||||
|
||||
### D7 — Le pass d'ombre partage les draw args → il dessine le niveau sélectionné
|
||||
|
||||
Le pass d'ombre lit le même buffer de draw args : les ombres utilisent **automatiquement** le
|
||||
niveau LOD (ombres moins chères, cohérent). Accepté pour v1 ; documenté. (Si un jour on veut
|
||||
des ombres au niveau fin, ce serait un deuxième buffer de draw args — hors périmètre.)
|
||||
|
||||
### D8 — Niveaux ≤ 4, seuils `[48, 12]` px en `conf.rs` (v1)
|
||||
|
||||
`MAX_LOD_LEVELS = 4`. Un mesh avec 1 seul niveau = LOD éteint pour lui (aucun changement de
|
||||
comportement). Les seuils par défaut sont des constantes (pas encore configurables par scene —
|
||||
extension triviale plus tard). Rayon projeté = le **même** rayon sphere que le culling
|
||||
(circonradius × max scale, centre transformé) → pas de nouvelle donnée géométrique.
|
||||
|
||||
### D9 — Interrupteur global LOD au niveau du Renderer
|
||||
|
||||
`Renderer::set_lod_enabled(bool)` (défaut `true`) — coupure générale, orthogonale au mode
|
||||
par mesh. Quand il est éteint, le calcul par entité est **court-circuité** (niveaux tous à
|
||||
0, aucune projection calculée) : tout se dessine au niveau 0, le GPU n'est pas touché
|
||||
(il lit simplement le niveau 0). Appelable à l'instanciation (juste après `Renderer::new`)
|
||||
ou **à chaud** (toggle de debug au runtime). Aucun changement de signature existante.
|
||||
|
||||
### D10 — Niveaux **générés automatiquement** pour les géométries utilisateur
|
||||
|
||||
L'utilisateur déclare sa géométrie comme aujourd'hui (`positions`/`indices`/`normals`/`uvs`,
|
||||
y compris générée procéduralement par lui) et la bibliothèque **calcule les niveaux LOD sous le
|
||||
capot**. Mécanisme : décimation gloutonne du maillage indexé —
|
||||
|
||||
```
|
||||
[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
|
||||
Geometry::decimated(&self, target_triangles: u32) -> Geometry (pure, sans GPU)
|
||||
```
|
||||
- Triangles (non dégénérés) triés par aire croissante ; on retire le plus petit jusqu'à la
|
||||
cible — **sans carte de topologie** (« Règle A ») : tout sous-ensemble de faces d'un mesh valide
|
||||
est un mesh valide (pas de trou à boucher, pas de flip). La cible est donc toujours atteignable
|
||||
(clamp `[1, T]`).
|
||||
- Sortie : re-indexation (weld par position exacte), **normales lisses recalculées** sur les faces
|
||||
survivantes (normales de faces accumulées par coin soudé, puis normalisées), UVs/couleurs =
|
||||
valeur du premier vertex de chaque groupe soudé (documenté : évite le bleeding entre seams UV ;
|
||||
le LOD sacrifie la précision UV au profit de la silhouette).
|
||||
- Non indexé à l'entrée → weld par position préalable (la sortie est **toujours indexée**).
|
||||
Triangles dégénérés jetés.
|
||||
- **Déterministe** (tri stable, tie-break par index d'origine) → counts reproductibles en tests.
|
||||
- Fallback : la cible est clampée à `[1, T]` — un mesh n'est jamais décimé sous un triangle, et
|
||||
le count **réel** va dans le tableau LOD (`validate()` OK, jamais de géométrie corrompue).
|
||||
|
||||
**L0 reste la géométrie exacte de l'utilisateur, byte-pour-byte** — seule la voie basse et les
|
||||
niveaux générés sont concernés par la simplification. Coût : setup uniquement (quelques ms pour
|
||||
des milliers de triangles), jamais par frame. Ratios par niveau : 0.5^k (50 %, 25 %, 12.5 %).
|
||||
|
||||
`add_mesh_lod(id, level, geometry)` (niveaux explicites) est la voie **générale** : elle
|
||||
**remplace** un niveau existant ou **ajoute** le niveau libre suivant (≤ 4 ; même jeu
|
||||
**d'attributs** et même **indexation** que L0 — validé). Applicable à **tout** mesh, y compris
|
||||
Auto — ex. : remplacer un niveau décimé d'une primitive de la maison par une régénération à
|
||||
la tessellation divisée par 2 (plus propre que la décimation pour les meshes réguliers).
|
||||
`create_mesh_with_lod` est l'ergonomie qui génère les niveaux 1..n à la création ; les deux
|
||||
voies sont combinables.
|
||||
|
||||
### D11 — Mode LOD par mesh (`Off` / `Auto` / `Manual`) — contrat pour les meshes updatables
|
||||
|
||||
Le mesh porte un mode interne : `Off` (1 niveau = comportement d'aujourd'hui), `Auto` (niveaux
|
||||
décimés, D10), `Manual` (niveaux fournis, `add_mesh_lod`).
|
||||
|
||||
Question posée : **un mesh updatable** (positions/normals réécrits par frame) — faut-il pouvoir
|
||||
désactiver le LOD au niveau mesh ? **Oui, et automatiquement** : les niveaux générés sont
|
||||
**dérivés de l'ancien L0** ; si la géométrie change, ils représentent une forme obsolète. Contrat
|
||||
réservé (pas d'API d'update dans v1, mais le design le prévoit) : toute future
|
||||
`mesh.update_geometry(…)` **réinitialise le mode à `Off`** (niveaux jetés, buffer packé reconstruit
|
||||
avec le seul nouveau L0, tableau à 1 ligne) — l'utilisateur n'a rien à désactiver : un mesh
|
||||
mis à jour par frame est simplement un mesh mono-niveau. Documenter cette règle quand l'API
|
||||
d'update existera.
|
||||
|
||||
## Flux de données (changement en surgraisse)
|
||||
|
||||
```
|
||||
CPU (chaque frame) GPU
|
||||
───────────────────── ───
|
||||
Transforms (slots) ──upload──► transforms
|
||||
CullUniforms (planes,…) ──upload──► cull_u
|
||||
BBoxes (par mesh) ──upload──► bboxes
|
||||
** LodLevels (par slot) ──upload──► lod_levels (binding 3)
|
||||
** LodTables (par mesh) ──upload──► lod_tables (binding 4)
|
||||
│
|
||||
compute_matrices (inchangé) │
|
||||
cull : zéro si culled/inactive│
|
||||
sinon write_level_args └──► draw_args
|
||||
│
|
||||
main pass (Étape 18 : groupé par material, set_vertex_buffer
|
||||
du buffer packé par slot, draw indirect lit count+offset)
|
||||
```
|
||||
|
||||
## Gain attendu (changements d'état / frame, passe principale)
|
||||
## Niveaux dans le démo
|
||||
|
||||
| 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) |
|
||||
Le démo exerce l'**API utilisateur demandée** : géométrie L0 déclarée, niveaux générés sous le
|
||||
capot (D10). Les primitives de la maison sont des « géométries utilisateur » pour le LOD :
|
||||
|
||||
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.
|
||||
| Mesh | L0 (déclaré) | L1, L2 (auto : décimation 50 % / 25 %) |
|
||||
|-----------|--------------|-----------------------------------------|
|
||||
| sphere | 32×20 | générés par `decimated` |
|
||||
| cylinder | 32 | générés |
|
||||
| cone | 32 | générés |
|
||||
| torus | 24×16 | générés |
|
||||
| icosphere | subdiv 2 | générés (triangles quasi uniformes → décimation propre) |
|
||||
| cube,plane| — | `create_mesh` simple, 1 niveau |
|
||||
|
||||
## Fichiers touchés
|
||||
Setup du démo : `create_mesh_with_lod(id, geo, mat, 3)` pour les 5 meshes tessellés, `create_mesh`
|
||||
pour cube/plane. (Les niveaux « re-tessellés » du draft initial servent aux tests unitaires de
|
||||
l'issue de secours `add_mesh_lod`, pas au démo.)
|
||||
|
||||
| Fichier | Action |
|
||||
|---------|--------|
|
||||
| `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 |
|
||||
## API publique (ajouts — aucune rupture)
|
||||
|
||||
## Périmètre exclus (reportés)
|
||||
```rust
|
||||
// Scene
|
||||
/// Crée un mesh dont les `levels` niveaux (2..=4) sont générés automatiquement
|
||||
/// par décimation (D10) de la géométrie fournie. L0 = la géométrie, exacte.
|
||||
pub fn create_mesh_with_lod(
|
||||
&mut self, id: &str, geometry: Geometry, material: Option<&str>, levels: u8,
|
||||
) -> Result<String, String>;
|
||||
|
||||
- **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).
|
||||
/// Remplace (ou ajoute, si `level` est le niveau libre suivant) un niveau **explicite** du mesh
|
||||
/// `id` (≤ 4 niveaux ; même jeu d'attributs et même indexation que L0 ; pack total < 65 536
|
||||
/// vertices). Reconstruit les buffers packés (D2). Combinable avec `create_mesh_with_lod`.
|
||||
pub fn add_mesh_lod(&mut self, id: &str, level: u8, geometry: Geometry) -> Result<(), String>;
|
||||
|
||||
## Risques & mitigations
|
||||
// Renderer
|
||||
/// Interrupteur global LOD (D9). Défaut `true`. `false` → tout se dessine au niveau 0,
|
||||
/// calcul par entité court-circuité. Appelable à l'instanciation ou à chaud.
|
||||
pub fn set_lod_enabled(&self, enabled: bool);
|
||||
```
|
||||
|
||||
| Risque | Mitigation |
|
||||
|--------|-----------|
|
||||
| **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. |
|
||||
`create_mesh` (sans LOD) reste la voie par défaut : un mesh simple se comporte exactement comme
|
||||
aujourd'hui. Le LOD s'active **explicitement** (`create_mesh_with_lod`) ou par niveaux explicites
|
||||
(`add_mesh_lod`) — pas de changement de comportement pour un utilisateur qui ne demande rien.
|
||||
|
||||
## Plan de vérification
|
||||
## Plan de code
|
||||
|
||||
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` 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é.
|
||||
| Fichier | Changement |
|
||||
|---|---|
|
||||
| `resources/uniform.rs` | `LodTable` (80 B : `count` + 4 × `LodRow{base, count}` 16 B) + `LOD_TABLE_SIZE` ; export. |
|
||||
| `shaders/gpu_driven.wgsl` | structs `LodTable`/`LodRow` ; bindings 3+4 group(2) ; helper `write_level_args` ; `cull` : les 2 branches visibles l'appellent ; commentaire header mis à jour. |
|
||||
| `math/geometry.rs` | `Geometry::decimated(target) -> Geometry` (D10, pure) + tests ; `Geometry::generate_lod_levels(levels)` (ratios 0.5^k). |
|
||||
| `resources/mesh.rs` | champ `lod_geometries: Vec<Arc<Geometry>>` + `lod_mode` (D11) ; **packer** : concatène `to_vertices()` des niveaux → `vertex_buffer` packé ; indices rebasés concaténés → `index_buffer` packé ; `num_*` = niveau 0 ; `LodTable::from_geometry` par mesh. |
|
||||
| `scene/scene.rs` | `create_mesh_with_lod` + `add_mesh_lod` (validation + rebuild packé via le contexte GPU de la scene) ; `packed_lod_levels()` (par frame, par slot — décision D1) ; `mesh_lod_tables()` (par mesh, comme `mesh_bboxes`). |
|
||||
| `core/renderer.rs` | buffer `lod_levels_buffer` (1 KB) + `lod_table_buffer` ; layout group(2) +2 bindings ; upload par frame ; **D1** : `lod_levels: Vec<u32>` état par slot + fn pures `lod_level`/`projected_radius_px` (view + proj[1][1] + hauteur viewport, déjà disponibles) ; **D9** : champ `lod_enabled` + `set_lod_enabled` (court-circuit du calcul) ; `debug_dump` : readback du buffer de niveaux. |
|
||||
| `utils/conf.rs` | `MAX_LOD_LEVELS = 4` ; `LOD_THRESHOLDS: [f32; 2] = [48.0, 12.0]`. |
|
||||
| `examples/demo.rs` | `create_mesh_with_lod(id, geo, mat, 3)` pour les 5 meshes tessellés (D10) ; cube/plane en `create_mesh`. |
|
||||
|
||||
## Critères d'acceptation (definition of done)
|
||||
### Fonction pure (testable, module `math` ou `resources::lod`)
|
||||
|
||||
- [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é.
|
||||
```rust
|
||||
/// Niveau cible : le plus grossier dont la borne est respectée, avec hystérésis asymétrique (D4).
|
||||
/// `thresholds[k]` = rayon px au-dessus duquel le niveau k+1 est exigé (descendants) ;
|
||||
/// niveaux au-delà du nb de seuils : borne clamped (dernier seuil).
|
||||
/// `last` = niveau de la frame précédente. Renvoie un niveau ≤ `max_level`.
|
||||
pub fn lod_level(radius_px: f32, last: u32, max_level: u32, thresholds: &[f32]) -> u32;
|
||||
|
||||
/// Rayon projeté en pixels de la bounding sphere (même sphere que le culling, D8).
|
||||
/// `depth <= eps` (objet dans la caméra) → f32::INFINITY (niveau 0).
|
||||
pub fn projected_radius_px(center_world: Vec3, radius: f32, view: Mat4, proj: Mat4, height_px: f32) -> f32;
|
||||
|
||||
/// Décimation gloutonne (D10, Règle A) : renvoie un `Geometry` indexé avec `target_triangles`
|
||||
/// triangles (clamp `[1, T]`) — les plus petits partent en premier (faible impact visuel),
|
||||
/// normales lisses recalculées. Déterministe.
|
||||
pub fn decimated(&self, target_triangles: u32) -> Geometry; // sur Geometry
|
||||
```
|
||||
|
||||
## Vérification
|
||||
|
||||
1. **Unit tests `lod_level`** (sans GPU) : seuils simples (r > t0 → 0, entre → 1, < t1 → 2) ;
|
||||
**anti-scintillement** : série de rayons oscillant ±10 % autour d'un seuil → le niveau reste
|
||||
stable (montée immédiate, descente retardée par le facteur 0.8) ; clamp `max_level` ;
|
||||
`INFINITY → 0`.
|
||||
2. **Unit tests `projected_radius_px`** : objet à l'origine, caméra sur Z → valeur analytique ;
|
||||
invariance d'échelle (objets 10× plus gros 10× plus loin → même rayon px).
|
||||
3. **Unit tests du packer** : offsets/counts par niveau sur un cas 2 niveaux ; niveau 0 à
|
||||
offset 0 ; indices rebasés.
|
||||
4. **Unit tests `decimated`** : grille 2×2 (4 triangles de tailles différentes) → cible 2 :
|
||||
count exact, positions ⊆ entrée, les plus petits retirés en premier ; entrée non indexée →
|
||||
sortie indexée weldée (coins dupliqués fondus, UV = premier rencontré) ; triangle dégénéré
|
||||
en entrée → jeté ; **déterminisme** (deux appels → sortie identique) ; icosahèdre 20 → 10 :
|
||||
count exact + bbox dans l'originale ; icosphere subdiv 1 → 50 % : count exact ; clamp : cible
|
||||
0 ou > T → clone ; `validate()` OK.
|
||||
5. **GPU A/B** : entités temporaires à z = +6 / +30 / +60 (démo) → readback `debug_dump` :
|
||||
niveaux 0 / 1 / 2 selon la distance et **draw args = count de la ligne du tableau
|
||||
CPU-uploadé** (self-cohérent, pas de count codé en dur — les niveaux auto ont des counts
|
||||
déterministes mais calculés) ; entité culled (hors frustum) → 0 quel que soit le niveau.
|
||||
6. **D9** : `set_lod_enabled(false)` au démarrage → draw args **bit-identiques** à l'exécutable
|
||||
pré-LOD (tous niveaux 0) ; toggle à chaud → bascule visible au readback.
|
||||
7. **D11** (contrat) : `add_mesh_lod` s'applique à tout mesh (y compris Auto — le niveau
|
||||
explicite remplace le niveau décimé au même index) ; le mode `Off` (1 niveau) = comportement
|
||||
d'aujourd'hui ; un mesh updatable futur repasse en `Off` (D11).
|
||||
8. **Régression** : 69 tests verts ; démo silencieux par défaut ; batching (Étape 18) intact
|
||||
(compteur de switches inchangé) ; mesh à 1 niveau → draw args **bit-identiques** à avant
|
||||
l'étape (D6).
|
||||
9. **Visuel** : zoom arrière sur le démo → les primitives basses du ring passent au niveau
|
||||
simplifié sans pop visible (silhouettes décimées proches + hystérésis D4).
|
||||
|
||||
## Critères d'acceptation
|
||||
|
||||
- [ ] `lod_level` + `projected_radius_px` pures, unit-testées (dont la série oscillante).
|
||||
- [ ] `decimated` pure, déterministe, unit-testée (counts, retire-les-plus-petits, weld + UV
|
||||
premier rencontré, clamp `[1, T]`, dégénérés, `validate()` OK).
|
||||
- [ ] Packer multi-niveaux : un buffer packé, niveau 0 à offset 0, indices rebasés.
|
||||
- [ ] `create_mesh_with_lod` (niveaux auto) + `add_mesh_lod` (niveaux explicites, ≤ 4,
|
||||
validation attributs/indexation/65536, combinables) ; aucune rupture d'API ;
|
||||
`create_mesh` seul = comportement d'aujourd'hui.
|
||||
- [ ] `Renderer::set_lod_enabled` (D9) : `false` → bit-identique à pré-LOD ; toggle à chaud.
|
||||
- [ ] Pass `cull` : les 2 branches visibles écrivent depuis le tableau LOD ; mesh 1 niveau
|
||||
bit-identique au comportement actuel.
|
||||
- [ ] A/B GPU : niveaux 0/1/2 choisis selon la distance, counts correspondants, culling intact.
|
||||
- [ ] 69+ tests verts, `cargo fmt` clean, démo silencieux par défaut.
|
||||
- [ ] Docs : `gpu-driven.md` (section LOD + contrainte buffers) ; `ARCHI_CPU_GPU.md` ; ROADMAP
|
||||
4.3 coché.
|
||||
- [ ] DRAFT.md vidé après validation utilisateur (convention de la maison).
|
||||
|
||||
## Hors périmètre (suivant)
|
||||
|
||||
- **API d'update de géométrie** (mesh updatable par frame) — D11 réserve le contrat
|
||||
(update ⇒ LOD réinitialisé à `Off`) mais l'API d'update elle-même est un autre sujet.
|
||||
- Décision du niveau **côté GPU** (viewProj dans CullUniforms + état GPU) — option D1 future.
|
||||
- Seuils **configurables par scene/mesh** (v1 : constantes `conf.rs`).
|
||||
- Réglage fin du **qualité/cout** de la décimation (poids par arête, quadric edge collapse)
|
||||
— v1 : aire + flip, suffisant pour les silhouettes.
|
||||
- Geometric morphing / transitions douces entre niveaux (les pops restent discrets avec les
|
||||
primitives procédurales ; le morphing est un sujet à part entière).
|
||||
- LOD de **shader** (simplification du lighting par distance) — autre item (HDR/tone mapping).
|
||||
- Ombres au niveau fin (D7 accepte le niveau LOD dans les ombres).
|
||||
|
||||
+1
-1
@@ -162,7 +162,7 @@ generated: { by: human:jerome, at: 2026-07-31T00:00:00Z }
|
||||
|
||||
### 4.3 Optimisations
|
||||
- [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)
|
||||
- [x] Level of Detail (LOD) — 2026-09-23 (Étape 19 : ≤ 4 niveaux/mesh — L0 exacte, L1–L3 par décimation gloutonne au setup (`Geometry::decimated`/`generate_lod_levels`, Règle A : retrait des plus petits triangles, weld par position exacte, rebase u16), packés dans les buffers vertex/index du mesh (offsets en unités d'élément, plafond 65 535 sommets) ; décision par frame **côté CPU** (sphère bounding projetée en pixels + hystérésis asymétrique ×0.8 — `math/lod.rs` pure, unit-testée), exécution **côté GPU** (le pass `cull` mappe niveau → ligne de la table LOD → args indirects) ; **activé par défaut**, `set_lod_enabled(false)` → rendu bit-à-bit identique au pré-LOD. Vérifié par readback GPU : zoom 4,6× → tous les meshes multi-niveaux passent au niveau 1 avec exactement leurs lignes L1 (ex. sphère 3840 → 1824 indices), stable frame à frame)
|
||||
- [ ] HDR + Tone Mapping (optionnel)
|
||||
|
||||
### 4.4 Gestion du Resize (cycle de vie Surface + Depth)
|
||||
|
||||
@@ -16,7 +16,7 @@ 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, Étape 17, validé 2026-09-22).**
|
||||
> **État du document : ACTUEL (implémenté — Phase 3 du ROADMAP, Étapes 17–19, 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
|
||||
@@ -37,6 +37,16 @@ Ce document sert de spécification technique et de trame d'implémentation pour
|
||||
> 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 ».
|
||||
> **LOD (Étape 19, 2026-09-23)** : le pass `cull` remplit désormais les arguments indirects à partir
|
||||
> du **niveau de détail** du slot, et non d'un seul jeu de comptes. Le choix du niveau est fait **côté CPU**
|
||||
> (rayon de la sphère bounding projeté en pixels + hystérésis asymétrique — `math/lod.rs`, pur et unit-testé) ;
|
||||
> le GPU n'effectue que le mappage niveau → ligne de la table LOD du mesh. Les niveaux d'un mesh sont
|
||||
> générés par décimation gloutonne au setup (`Geometry::decimated` : suppression des plus petits triangles,
|
||||
> soudure par position exacte, rebase u16) et **empilés dans les buffers vertex/index du mesh** (offsets en
|
||||
> unités d'élément, pas d'octet — c'est ce qu'exigent les arguments `drawIndirect*` de WebGPU ; plafond u16 :
|
||||
> 65 535 sommets/mesh, 4 niveaux max). LOD activé par défaut ; `set_lod_enabled(false)` restaure un rendu
|
||||
> bit-à-bit identique au pré-LOD (niveau 0 partout = comptes complets). Détail : `docs/user/gpu-driven.md`
|
||||
> § « Level of Detail ».
|
||||
|
||||
1. Répartition des Rôles : CPU vs GPU (La Source de Vérité)
|
||||
|
||||
@@ -45,7 +55,8 @@ Pour éviter les goulets d'étranglement dus aux allers-retours sur le bus PCIe,
|
||||
Côté CPU (Source de Vérité)
|
||||
- Ce qu'il conserve : Les données logiques et les transformations brutes des objets (ex: Vec<Transform> contenant la position, la rotation, et l'échelle).
|
||||
- Ce qu'il fait : Il gère la logique de jeu, l'IA, le réseau et les interactions globales.
|
||||
- Ce qu'il ne fait plus : Il ne calcule plus les matrices de transformation mondiales (World Matrices) en masse, et ne fait plus de tests de visibilité unitaires.
|
||||
- Ce qu'il ne fait plus : Il ne calcule plus les matrices de transformation mondiales (World Matrices) en masse, et ne fait plus de tests de visibilité unitaires (le culling frustum reste 100 % GPU).
|
||||
- Ce qu'il fait en plus (LOD, Étape 19) : le **choix du niveau de détail** par entité — O(N) projections de sphères en pixels + hystérésis, coût négligeable. C'est l'unique décision de visibilité/détail conservée côté CPU : elle dépend de la taille écran (un choix artistique), pas de la géométrie, et l'hystérésis a besoin de l'état de la frame précédente.
|
||||
|
||||
Côté GPU (Exécutant Autonome)
|
||||
- Ce qu'il calcule : Les World Matrices, le Frustum Culling, et la génération des listes de dessin indirectes.
|
||||
@@ -64,11 +75,11 @@ L'exécution des tâches s'appuie sur une structure séquentielle stricte au sei
|
||||
```
|
||||
|
||||
Étape par étape :
|
||||
- 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).
|
||||
- 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), **et le niveau LOD de chaque slot** (1 u32/slot, Étape 19). 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 : 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).
|
||||
- Remplissage du Buffer Indirect : le pass `cull` lit le **niveau LOD** du slot, en choisit la ligne dans la table LOD du mesh (`LodTable` : 4 lignes d'offsets/comptes en unités d'élément) et écrit les arguments dans le `DrawSlot` (80 o) — mis à 0 si l'objet est cullé ou inactif (no-op). Le niveau 0 porte les comptes du mesh complet, donc LOD désactivé ≡ pré-LOD bit-à-bit.
|
||||
- Pass de Rendu (Render Pass) :
|
||||
- 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.
|
||||
@@ -89,6 +100,15 @@ L'implémentation utilise les buffers wGPU suivants (tous créés 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) |
|
||||
| Lod Levels | Niveau LOD par slot choisi par le CPU (4 o/slot) | Storage Buffer (read-only) | CPU → GPU (chaque frame) |
|
||||
| Lod Tables | Table par mesh : `count` + 4 lignes de 16 o (offset/compte d'éléments) (80 o/mesh) | Storage Buffer (read-only) | CPU → GPU (quand l'ensemble des meshes change) |
|
||||
|
||||
**Buffers de géométrie LOD (Étape 19)** : les niveaux d'un mesh sont **empilés** — un seul buffer vertex et un
|
||||
seul buffer index par mesh, contenant les niveaux concaténés (L0, L1, …). Les lignes de la table LOD portent
|
||||
les offsets en **unités d'élément** (premier vertex / premier index), car les arguments `drawIndirect*` de
|
||||
WebGPU s'expriment en éléments, et le buffer est lié en entier à l'offset 0. Conséquences : un mesh LOD ne
|
||||
peux dépasser 65 535 sommets au total (indices u16) et 4 niveaux (`MAX_LOD_LEVELS`) ; le mélange indexé/non-indexé
|
||||
dans un même mesh est supporté (la commande de draw par slot suit le niveau choisi par le CPU).
|
||||
|
||||
## Liens
|
||||
|
||||
|
||||
+87
-5
@@ -2,7 +2,8 @@
|
||||
|
||||
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**.
|
||||
draw calls. This page explains what that means for you, how to opt into **frustum culling**, and how the
|
||||
**Level of Detail (LOD)** system works.
|
||||
|
||||
## What runs on the GPU
|
||||
|
||||
@@ -13,7 +14,9 @@ Each frame, before the render passes, two compute passes run over a fixed-capaci
|
||||
(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).
|
||||
vertex/index count, zeroed when the entity is culled or inactive). With LOD on (the default), the
|
||||
count it writes comes from the mesh's **LOD table** at the level the CPU chose for the slot this
|
||||
frame — see [Level of Detail](#level-of-detail-lod).
|
||||
|
||||
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
|
||||
@@ -77,6 +80,77 @@ out of view (false negative), but it will **never cull an object that is actuall
|
||||
(false positive). For tight culling you would need per-mesh sphere fitting or per-face tests,
|
||||
which are out of scope for v1.
|
||||
|
||||
## Level of Detail (LOD)
|
||||
|
||||
LOD is **on by default**: distant entities automatically draw a coarser version of their mesh, so
|
||||
the GPU stops spending fillrate and vertex work on detail the eye cannot see. It is a quality
|
||||
feature with a performance payoff — unlike culling, it is safe to leave on because the
|
||||
worst case (a level chosen too fine) is exactly what you would have drawn anyway.
|
||||
|
||||
### How it works
|
||||
|
||||
LOD is a **CPU-decided, GPU-executed** split (the one deliberate per-entity decision kept on the
|
||||
CPU):
|
||||
|
||||
1. **Setup (once per mesh).** Each mesh can carry up to 4 levels. Levels 1..3 are generated
|
||||
automatically from level 0 by greedy decimation (`Geometry::generate_lod_levels`): the smallest
|
||||
triangles are removed first (no edge map, no crease handling — a subset of the faces of a valid
|
||||
mesh is valid), duplicate corners are welded, and the levels are **packed into the mesh's single
|
||||
vertex/index buffers** (see the constraint below). Level 0 is always your exact geometry.
|
||||
2. **Per frame (CPU).** For each entity, the bounding sphere used by culling is projected to screen
|
||||
pixels (its *perceived size*); that radius picks a level with **asymmetric hysteresis** — going
|
||||
finer is immediate, going coarser only below 80 % of the bound (a 20 % dead band) — which is what
|
||||
prevents flicker when an entity hovers around a threshold. Default thresholds: 48 px and 12 px
|
||||
(bigger than 48 px → full detail; smaller than 12 px → coarsest).
|
||||
3. **Per frame (GPU).** The `cull` pass reads the slot's level, looks up the matching row of the
|
||||
mesh's LOD table (element-unit offsets + counts), and writes the indirect draw arguments from it.
|
||||
|
||||
### Using it
|
||||
|
||||
```rust
|
||||
// One level (the default): create_mesh is unchanged.
|
||||
let id = scene.create_mesh("hero", &geometry, &material)?;
|
||||
|
||||
// Auto-generated levels 1..3 (decimated at half, quarter, eighth the triangle count).
|
||||
let id = scene.create_mesh_with_lod("hero", &geometry, &material, 4)?;
|
||||
|
||||
// Or supply your own levels (same attributes, same indexed-ness as level 0).
|
||||
scene.add_mesh_lod("hero", 1, &my_coarse_geometry)?;
|
||||
```
|
||||
|
||||
Toggle at runtime (off = every slot forced to level 0 = byte-identical rendering to the pre-LOD
|
||||
engine — the level-0 rows carry the full-mesh counts, so nothing else changes):
|
||||
|
||||
```rust
|
||||
app.renderer().set_lod_enabled(false);
|
||||
```
|
||||
|
||||
### Constraint: packed LOD buffers
|
||||
|
||||
A level is **not a separate buffer**: the mesh's levels are concatenated into its one vertex buffer
|
||||
and one index buffer, and the per-mesh LOD table stores each level's offsets/counts. Two
|
||||
consequences:
|
||||
|
||||
- **u16 indices** → the *sum* of all levels must stay under 65 535 vertices (the scene rejects a
|
||||
level set that would not fit, with a clear error);
|
||||
- **at most 4 levels** per mesh (`MAX_LOD_LEVELS`, also the size of the GPU table row).
|
||||
|
||||
Indexed-ness: levels supplied through `add_mesh_lod` must match level 0's indexed-ness (validated).
|
||||
Auto-generated levels from a **non-indexed** level 0 are indexed anyway (decimation rebuilds with
|
||||
indices), and the packed buffer supports that mix — the per-slot draw command follows the level the
|
||||
CPU chose (the shadow pass always uses the level-0 command, so casters stay at full detail).
|
||||
|
||||
### How to verify LOD with the debug dump
|
||||
|
||||
The debug dump (below) prints, per frame: the per-slot **levels** and each mesh's **LOD table**
|
||||
(rows = `vertex_offset / vertex_count / index_offset / index_count`, element units). The clean test
|
||||
is to **zoom the camera out**: the entities' perceived size drops below the thresholds, the levels
|
||||
step up (0 → 1 → 2), and the indirect argument counts shrink to the corresponding rows — e.g. the
|
||||
demo's 3 840-index sphere drops to 1 824, then 912 — while the levels stay **stable frame to frame**
|
||||
(hysteresis holding). Verified 2026-09-23: at the demo's default distance every entity sits at
|
||||
level 0 with full counts; zoomed to 4.6×, all multi-level meshes select level 1 with exactly their
|
||||
L1 rows, stable across frames.
|
||||
|
||||
## Debugging the GPU path
|
||||
|
||||
If something looks wrong — a missing object, a black window — the GPU-side slot tables can be
|
||||
@@ -88,9 +162,10 @@ 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
|
||||
indirect draw arguments, the mesh bounding boxes, the cull uniforms, the per-slot **LOD levels**
|
||||
and the per-mesh **LOD tables**. 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 — and with LOD
|
||||
on, the *row* the count comes from tells you the selected level (see above). In the `demo` example the dump is opt-in via an environment
|
||||
variable, so the showcase stays silent by default:
|
||||
|
||||
```sh
|
||||
@@ -137,6 +212,13 @@ and verified fixed, see the D14 note in `docs/tech/ARCHI_CPU_GPU.md`.)
|
||||
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).
|
||||
- **LOD levels are packed into the mesh's own buffers**: u16 indices cap the *total* across all
|
||||
levels at 65 535 vertices, and there are at most 4 levels. The decimation (smallest-triangle
|
||||
removal + welding) is a setup-time cost only (a few ms for thousands of triangles); the
|
||||
per-frame cost is one sphere projection per entity on the CPU.
|
||||
- **LOD detail loss is visible by design** — the hysteresis dead band makes the pop rare and
|
||||
one-directional (immediate when gaining detail, delayed when losing it), but a coarse level is
|
||||
coarser. `set_lod_enabled(false)` is the escape hatch.
|
||||
|
||||
Culling is a **performance** feature, not a visual one: with it off you get the same image with
|
||||
the indirect-draw machinery still active.
|
||||
|
||||
+25
-6
@@ -11,7 +11,11 @@
|
||||
//! * `R` resets the view, keys `1`/`2`/`3` jump to front / side / top presets,
|
||||
//! * a **directional** light (the shadow caster) + a **point** light + a **spot** light,
|
||||
//! so the shadow of the cube and the colored light halos are all visible,
|
||||
//! * the primitives slowly rotate in `update`, so depth, lighting and shadows read clearly.
|
||||
//! * the primitives slowly rotate in `update`, so depth, lighting and shadows read clearly,
|
||||
//! * **LOD** (Step 19): the rounded primitives are created with three levels each
|
||||
//! (`create_mesh_with_lod`, auto-decimated by halving targets); the CPU picks each entity's
|
||||
//! level from its projected screen size (with hysteresis) — zoom in/out with the wheel and
|
||||
//! the sphere/cylinder/cone/torus visibly lose detail as they shrink on screen.
|
||||
//!
|
||||
//! Doc (this header) follows the English convention used for examples; internal comments stay
|
||||
//! concise and French where helpful. Run with:
|
||||
@@ -118,23 +122,38 @@ impl AppHandler for Demo {
|
||||
app.scene.add_entity("ground", "ground_mesh").unwrap();
|
||||
|
||||
// 4. One mesh per primitive, each assigned to a textured (or stripe) material.
|
||||
// The cube + ground stay single-level (tiny meshes — LOD would buy nothing); the
|
||||
// rounded primitives get three LOD levels each (Step 19): level 0 is the full mesh,
|
||||
// levels 1.. are auto-generated by greedy decimation at halving targets (D10), all
|
||||
// packed into the mesh's single vertex/index buffers (D7). Zooming with the wheel
|
||||
// switches levels on the fly (asymmetric hysteresis, D4).
|
||||
app.scene
|
||||
.create_mesh("cube_mesh", cube(0.8), Some("solid_mat"))
|
||||
.unwrap();
|
||||
app.scene
|
||||
.create_mesh("sphere_mesh", uv_sphere(0.55, 32, 20), Some("stripes_mat"))
|
||||
.create_mesh_with_lod(
|
||||
"sphere_mesh",
|
||||
uv_sphere(0.55, 32, 20),
|
||||
Some("stripes_mat"),
|
||||
3,
|
||||
)
|
||||
.unwrap();
|
||||
app.scene
|
||||
.create_mesh("ico_mesh", icosphere(0.5, 2), Some("solid_mat"))
|
||||
.create_mesh_with_lod("ico_mesh", icosphere(0.5, 2), Some("solid_mat"), 3)
|
||||
.unwrap();
|
||||
app.scene
|
||||
.create_mesh("cyl_mesh", cylinder(0.4, 0.9, 32), Some("stripes_mat"))
|
||||
.create_mesh_with_lod("cyl_mesh", cylinder(0.4, 0.9, 32), Some("stripes_mat"), 3)
|
||||
.unwrap();
|
||||
app.scene
|
||||
.create_mesh("cone_mesh", cone(0.45, 0.9, 32), Some("solid_mat"))
|
||||
.create_mesh_with_lod("cone_mesh", cone(0.45, 0.9, 32), Some("solid_mat"), 3)
|
||||
.unwrap();
|
||||
app.scene
|
||||
.create_mesh("torus_mesh", torus(0.42, 0.16, 24, 16), Some("solid_mat"))
|
||||
.create_mesh_with_lod(
|
||||
"torus_mesh",
|
||||
torus(0.42, 0.16, 24, 16),
|
||||
Some("solid_mat"),
|
||||
3,
|
||||
)
|
||||
.unwrap();
|
||||
|
||||
place("cube_e", "cube_mesh", app, 0);
|
||||
|
||||
+2
-2
@@ -7,11 +7,11 @@ This is the source tree for `wsg-lib`, a Rust library wrapping [wgpu](https://gi
|
||||
| Module | Responsibility |
|
||||
|--------|---------------|
|
||||
| **core** | Manager (Context) + Executor (Renderer) layers — GPU lifecycle and draw call orchestration (incl. the GPU-driven compute passes + opt-in frustum culling); also `InputState` (unified keyboard/mouse input, Step 15.B) |
|
||||
| **resources** | Data types: Vertex (CPU-side), Mesh (GPU geometry + bounding box), Material (appearance descriptor), Texture, Lights, Camera + CameraController, and the uniform slot types (`TransformSlot`/`MatSlot`/`BBoxSlot`/`DrawSlot`/`CullUniforms`) |
|
||||
| **resources** | Data types: Vertex (CPU-side), Mesh (GPU geometry + bounding box, multi-level LOD via packed vertex/index buffers), Material (appearance descriptor), Texture, Lights, Camera + CameraController, and the uniform slot types (`TransformSlot`/`MatSlot`/`BBoxSlot`/`DrawSlot`/`CullUniforms`/`LodRow`/`LodTable`) |
|
||||
| **pipeline** | PipelineCache — WGSL shader loading and RenderPipeline compilation cache |
|
||||
| **shaders** | Embedded WGSL sources (`standard`, `shadow`, `gpu_driven`) loaded via the `include_str!` fallback in `utils::conf` |
|
||||
| **scene** | Scene — resource depot and slot-based entity graph for declarative rendering setup (Step 17) |
|
||||
| **math** | Transform, Geometry (per-attribute mesh data + AABB), `Frustum` (Gribb–Hartmann, WebGPU `[0,1]` z) and `primitives` (procedural mesh generators) |
|
||||
| **math** | Transform, Geometry (per-attribute mesh data + AABB, greedy decimation `decimated`/`generate_lod_levels`), `Frustum` (Gribb–Hartmann, WebGPU `[0,1]` z), `lod` (per-frame level selection: `projected_radius_px` + `lod_level` with asymmetric hysteresis) and `primitives` (procedural mesh generators) |
|
||||
| **utils** | Configuration constants and WsgError type |
|
||||
| **app** | App facade — high-level application orchestration with window lifecycle, event loop, and render automation |
|
||||
| **handler** | AppHandler trait — user-defined game logic interface injected into the render loop |
|
||||
|
||||
+230
-13
@@ -22,25 +22,26 @@
|
||||
use crate::core::Context;
|
||||
use crate::core::Frame;
|
||||
use crate::math::Frustum;
|
||||
use crate::math::lod::{lod_level, projected_radius_px};
|
||||
use crate::pipeline::{
|
||||
DEPTH_FORMAT, build_shadow_pipeline, create_shadow_map_bind_group_layout,
|
||||
create_shadow_uniform_layout, create_uniform_bind_group_layouts,
|
||||
};
|
||||
use crate::resources::uniform::{
|
||||
BBOX_SLOT_SIZE, BBoxSlot, CULL_UNIFORMS_SIZE, DRAW_SLOT_SIZE, DrawSlot, FRAME_UNIFORMS_SIZE,
|
||||
MAT_SLOT_SIZE, MAX_LIGHTS, MatSlot, OBJECT_UNIFORM_SIZE, SHADOW_UNIFORM_SIZE,
|
||||
TRANSFORM_SLOT_SIZE, TransformSlot,
|
||||
LOD_TABLE_SIZE, LodTable, MAT_SLOT_SIZE, MAX_LIGHTS, MatSlot, OBJECT_UNIFORM_SIZE,
|
||||
SHADOW_UNIFORM_SIZE, TRANSFORM_SLOT_SIZE, TransformSlot,
|
||||
};
|
||||
use crate::resources::{
|
||||
Camera, CullUniforms, FrameUniforms, Lights, Material, Mesh, ObjectUniform, ShadowUniform,
|
||||
};
|
||||
use crate::scene::Scene;
|
||||
use crate::utils::conf::{
|
||||
GPU_DRIVEN_SHADER, GPU_WORKGROUP_SIZE, MAX_ENTITIES, SHADOW_DEPTH_BIAS, SHADOW_MAP_SIZE,
|
||||
SHADOW_SCENE_CENTER, SHADOW_SCENE_RADIUS,
|
||||
GPU_DRIVEN_SHADER, GPU_WORKGROUP_SIZE, LOD_THRESHOLDS, MAX_ENTITIES, MAX_LOD_LEVELS,
|
||||
SHADOW_DEPTH_BIAS, SHADOW_MAP_SIZE, SHADOW_SCENE_CENTER, SHADOW_SCENE_RADIUS,
|
||||
};
|
||||
use glam::{Mat4, Vec3, Vec4};
|
||||
use std::cell::Cell;
|
||||
use glam::{Mat4, Quat, Vec3, Vec4};
|
||||
use std::cell::{Cell, RefCell};
|
||||
use std::collections::HashMap;
|
||||
use std::hash::Hash;
|
||||
use std::sync::Arc;
|
||||
@@ -108,6 +109,12 @@ pub struct Renderer {
|
||||
draw_args_buffer: wgpu::Buffer,
|
||||
/// GPU cull uniforms (`UNIFORM | COPY_DST`): frustum planes + control flags; rewritten each frame.
|
||||
cull_uniform_buffer: wgpu::Buffer,
|
||||
/// GPU per-slot LOD levels (`STORAGE | COPY_DST`): one u32 per entity slot, the CPU's per-frame
|
||||
/// level decision (Step 19, D8); read by `cull` (group 2, binding 3).
|
||||
lod_levels_buffer: wgpu::Buffer,
|
||||
/// GPU per-mesh LOD tables (`STORAGE | COPY_DST`): one 80-byte [`LodTable`] per mesh in
|
||||
/// `mesh_order` order; read by `cull` (group 2, binding 4) to map a level to its draw args.
|
||||
lod_tables_buffer: wgpu::Buffer,
|
||||
/// `compute_matrices`/`cull` group 0 (transform buffer, storage read) — shared by both compute passes.
|
||||
transform_bg: wgpu::BindGroup,
|
||||
/// `compute_matrices` group 1 (matrix buffer, storage read_write).
|
||||
@@ -127,6 +134,17 @@ pub struct Renderer {
|
||||
/// frame. Interior-mutable (all-`&self` API); exposed through `debug_dump` for the
|
||||
/// state-change A/B measurement (Étape 18 verification).
|
||||
debug_pipeline_switches: Cell<u32>,
|
||||
/// Whether LOD is enabled (Step 19, D8). When `false` the CPU writes level 0 for every slot
|
||||
/// each frame, and the GPU indirect args are byte-identical to the pre-LOD behavior
|
||||
/// (level-0 rows carry the full-mesh counts). Interior-mutable (all-`&self` API).
|
||||
lod_enabled: Cell<bool>,
|
||||
/// Per-slot level of the PREVIOUS frame — the hysteresis state of [`lod_level`] (Step 19, D4):
|
||||
/// going coarser requires a 20 % dead band measured against this value. Interior-mutable
|
||||
/// (all-`&self` API); resized when the slot count grows (entity append).
|
||||
last_lod_levels: RefCell<Vec<u32>>,
|
||||
/// Viewport height in pixels (Step 19, D9): the unit of the LOD projected-size test. Set from
|
||||
/// the initial surface size in `new` and refreshed by `resize_depth` on window resize.
|
||||
viewport_height: u32,
|
||||
}
|
||||
|
||||
impl Renderer {
|
||||
@@ -310,6 +328,27 @@ impl Renderer {
|
||||
},
|
||||
count: None,
|
||||
},
|
||||
// Step 19 (LOD): per-slot levels (storage read) + per-mesh tables (storage read).
|
||||
wgpu::BindGroupLayoutEntry {
|
||||
binding: 3,
|
||||
visibility: wgpu::ShaderStages::COMPUTE,
|
||||
ty: wgpu::BindingType::Buffer {
|
||||
ty: wgpu::BufferBindingType::Storage { read_only: true },
|
||||
has_dynamic_offset: false,
|
||||
min_binding_size: None,
|
||||
},
|
||||
count: None,
|
||||
},
|
||||
wgpu::BindGroupLayoutEntry {
|
||||
binding: 4,
|
||||
visibility: wgpu::ShaderStages::COMPUTE,
|
||||
ty: wgpu::BindingType::Buffer {
|
||||
ty: wgpu::BufferBindingType::Storage { read_only: true },
|
||||
has_dynamic_offset: false,
|
||||
min_binding_size: None,
|
||||
},
|
||||
count: None,
|
||||
},
|
||||
],
|
||||
});
|
||||
let gpu_pipeline_layout = device.create_pipeline_layout(&wgpu::PipelineLayoutDescriptor {
|
||||
@@ -386,6 +425,25 @@ impl Renderer {
|
||||
| wgpu::BufferUsages::COPY_SRC,
|
||||
mapped_at_creation: false,
|
||||
});
|
||||
// Step 19 (LOD): per-slot levels (one u32 per entity slot, CPU-written each frame) and
|
||||
// per-mesh tables (one 80-byte LodTable per mesh, mesh_order order). `COPY_SRC` lets
|
||||
// `debug_dump` read them back. The cull pass maps each slot's level to its draw args.
|
||||
let lod_levels_buffer = device.create_buffer(&wgpu::BufferDescriptor {
|
||||
label: Some("GPU LOD levels"),
|
||||
size: MAX_ENTITIES as u64 * 4,
|
||||
usage: wgpu::BufferUsages::STORAGE
|
||||
| wgpu::BufferUsages::COPY_DST
|
||||
| wgpu::BufferUsages::COPY_SRC,
|
||||
mapped_at_creation: false,
|
||||
});
|
||||
let lod_tables_buffer = device.create_buffer(&wgpu::BufferDescriptor {
|
||||
label: Some("GPU LOD tables"),
|
||||
size: MAX_ENTITIES as u64 * LOD_TABLE_SIZE,
|
||||
usage: wgpu::BufferUsages::STORAGE
|
||||
| wgpu::BufferUsages::COPY_DST
|
||||
| wgpu::BufferUsages::COPY_SRC,
|
||||
mapped_at_creation: false,
|
||||
});
|
||||
// Bind groups against the explicit layouts. `transform_bg` is shared by both compute passes
|
||||
// (group 0); `matrices_bg` by `compute_matrices` (group 1); `cull_bundle_bg` by `cull`
|
||||
// (group 2). `matrix_object_bg` uses the render pipeline's dynamic object layout (group 1) and
|
||||
@@ -423,6 +481,15 @@ impl Renderer {
|
||||
binding: 2,
|
||||
resource: draw_args_buffer.as_entire_binding(),
|
||||
},
|
||||
// Step 19 (LOD): the level + table buffers (whole-buffer bindings).
|
||||
wgpu::BindGroupEntry {
|
||||
binding: 3,
|
||||
resource: lod_levels_buffer.as_entire_binding(),
|
||||
},
|
||||
wgpu::BindGroupEntry {
|
||||
binding: 4,
|
||||
resource: lod_tables_buffer.as_entire_binding(),
|
||||
},
|
||||
],
|
||||
});
|
||||
let matrix_object_bg = device.create_bind_group(&wgpu::BindGroupDescriptor {
|
||||
@@ -464,12 +531,17 @@ impl Renderer {
|
||||
bbox_buffer,
|
||||
draw_args_buffer,
|
||||
cull_uniform_buffer,
|
||||
lod_levels_buffer,
|
||||
lod_tables_buffer,
|
||||
transform_bg,
|
||||
matrices_bg,
|
||||
cull_bundle_bg,
|
||||
matrix_object_bg,
|
||||
cull_enabled: Cell::new(false),
|
||||
debug_pipeline_switches: Cell::new(0),
|
||||
lod_enabled: Cell::new(true),
|
||||
last_lod_levels: RefCell::new(Vec::new()),
|
||||
viewport_height: height,
|
||||
};
|
||||
// Seed the shared frame buffer with an identity camera + current unlit flag so the low-level
|
||||
// `render` path (which has no window/camera) sees coherent values before `render_scene` runs.
|
||||
@@ -511,6 +583,8 @@ impl Renderer {
|
||||
let (depth_texture, depth_view) = create_depth_texture(&self.device, width, height);
|
||||
self._depth_texture = depth_texture;
|
||||
self.depth_view = depth_view;
|
||||
// Step 19 (D9): refresh the viewport height — the unit of the LOD projected-size test.
|
||||
self.viewport_height = height;
|
||||
}
|
||||
|
||||
/// Updates the stored surface texture format after a surface reconfigure (ROADMAP Phase 4.4).
|
||||
@@ -704,16 +778,40 @@ impl Renderer {
|
||||
bytemuck::cast_slice(&transform_slots),
|
||||
);
|
||||
|
||||
// 2b. LOD (Step 19, D8): the CPU picks each slot's detail level from the projected
|
||||
// bounding-sphere radius (asymmetric hysteresis in `math::lod`), then uploads the
|
||||
// per-slot levels + per-mesh tables the `cull` pass maps level → indirect args.
|
||||
// LOD disabled ⇒ every level is 0, and level-0 rows carry the full-mesh counts, so
|
||||
// the GPU args are byte-identical to the pre-LOD behavior (D8 compatibility).
|
||||
let camera = scene.camera();
|
||||
let cam_view = camera.view_matrix();
|
||||
let cam_proj = camera.projection_matrix(aspect);
|
||||
let bboxes = scene.mesh_bboxes();
|
||||
let lod_levels: Vec<u32> = if self.lod_enabled.get() {
|
||||
self.compute_lod_levels(scene, &cam_view, &cam_proj, &transform_slots, &bboxes)
|
||||
} else {
|
||||
vec![0u32; transform_slots.len()]
|
||||
};
|
||||
self.queue.write_buffer(
|
||||
&self.lod_levels_buffer,
|
||||
0,
|
||||
bytemuck::cast_slice(&lod_levels),
|
||||
);
|
||||
let lod_tables = scene.mesh_lod_tables();
|
||||
self.queue.write_buffer(
|
||||
&self.lod_tables_buffer,
|
||||
0,
|
||||
bytemuck::cast_slice(&lod_tables),
|
||||
);
|
||||
|
||||
// 3. Upload the local-space bounding boxes (small; the mesh set is static in practice, but
|
||||
// re-uploading each frame keeps the mesh-index → bbox mapping correct if meshes are added).
|
||||
let bboxes = scene.mesh_bboxes();
|
||||
self.queue
|
||||
.write_buffer(&self.bbox_buffer, 0, bytemuck::cast_slice(&bboxes));
|
||||
|
||||
// 4. Compute the view frustum from the camera's view-projection and write the cull uniforms
|
||||
// (six unit planes + the num_slots / culling control flags).
|
||||
let camera = scene.camera();
|
||||
let view_proj = camera.projection_matrix(aspect) * camera.view_matrix();
|
||||
let view_proj = cam_proj * cam_view;
|
||||
let frustum = Frustum::from_view_proj(&view_proj);
|
||||
let cull_uniforms =
|
||||
CullUniforms::from_frustum(&frustum, scene.num_slots() as u32, self.cull_enabled.get());
|
||||
@@ -831,7 +929,16 @@ impl Renderer {
|
||||
// Group 1 (dynamic): the 64-byte matrix slice for this slot.
|
||||
render_pass.set_bind_group(1, &self.matrix_object_bg, &[object_offset]);
|
||||
render_pass.set_vertex_buffer(0, slot.mesh.vertex_buffer.slice(..));
|
||||
if slot.has_index {
|
||||
// Step 19: the draw command follows the CHOSEN level's indexedness, not L0's —
|
||||
// an Auto-mode mesh may mix indexed levels (e.g. L0 indexed, L1+ non-indexed).
|
||||
// Level 0 (LOD off, or a single-level mesh) reproduces the pre-LOD command.
|
||||
let row = slot
|
||||
.mesh
|
||||
.lod_rows()
|
||||
.get(lod_levels[slot.slot_index] as usize)
|
||||
.copied()
|
||||
.unwrap_or_default();
|
||||
if row.index_count > 0 {
|
||||
if let Some(index_buffer) = &slot.mesh.index_buffer {
|
||||
render_pass.set_index_buffer(
|
||||
index_buffer.slice(..),
|
||||
@@ -879,6 +986,8 @@ impl Renderer {
|
||||
let cull_buf = self.cull_uniform_buffer.clone();
|
||||
let transform_buf = self.transform_buffer.clone();
|
||||
let bbox_buf = self.bbox_buffer.clone();
|
||||
let lod_levels_buf = self.lod_levels_buffer.clone();
|
||||
let lod_tables_buf = self.lod_tables_buffer.clone();
|
||||
let n = n.min(MAX_ENTITIES as u32).max(1);
|
||||
{
|
||||
// Creates a MAP_READ staging buffer and records a copy of `size` bytes from `src`
|
||||
@@ -900,12 +1009,14 @@ impl Renderer {
|
||||
read
|
||||
}
|
||||
|
||||
let specs: [(&wgpu::Buffer, u64); 5] = [
|
||||
let specs: [(&wgpu::Buffer, u64); 7] = [
|
||||
(&matrix_buf, n as u64 * MAT_SLOT_SIZE),
|
||||
(&draw_args_buf, n as u64 * DRAW_SLOT_SIZE),
|
||||
(&transform_buf, n as u64 * TRANSFORM_SLOT_SIZE),
|
||||
(&bbox_buf, n as u64 * BBOX_SLOT_SIZE),
|
||||
(&cull_buf, CULL_UNIFORMS_SIZE),
|
||||
(&lod_levels_buf, n as u64 * 4),
|
||||
(&lod_tables_buf, n as u64 * LOD_TABLE_SIZE),
|
||||
];
|
||||
let (tx, rx) = std::sync::mpsc::channel::<()>();
|
||||
let mut reads = Vec::with_capacity(specs.len());
|
||||
@@ -953,8 +1064,17 @@ impl Renderer {
|
||||
for b in &reads {
|
||||
b.unmap();
|
||||
}
|
||||
let (mat_data, args_data, tr_data, bb_data, cull_data) =
|
||||
(&data[0], &data[1], &data[2], &data[3], &data[4]);
|
||||
let (
|
||||
mat_data,
|
||||
args_data,
|
||||
tr_data,
|
||||
bb_data,
|
||||
cull_data,
|
||||
lod_levels_data,
|
||||
lod_tables_data,
|
||||
) = (
|
||||
&data[0], &data[1], &data[2], &data[3], &data[4], &data[5], &data[6],
|
||||
);
|
||||
for i in 0..n {
|
||||
let off = (i as u64 * TRANSFORM_SLOT_SIZE) as usize;
|
||||
let t: TransformSlot =
|
||||
@@ -990,6 +1110,28 @@ impl Renderer {
|
||||
for (i, p) in c.planes.iter().enumerate() {
|
||||
eprintln!("[dbg] plane[{i}] = {p:?}");
|
||||
}
|
||||
// Step 19 (LOD): the CPU-decided per-slot levels and the per-mesh tables the GPU
|
||||
// maps level → indirect args from.
|
||||
let levels: Vec<u32> = (0..n)
|
||||
.map(|i| {
|
||||
bytemuck::pod_read_unaligned(
|
||||
&lod_levels_data[i as usize * 4..i as usize * 4 + 4],
|
||||
)
|
||||
})
|
||||
.collect();
|
||||
eprintln!(
|
||||
"[dbg] lod: enabled={} viewport_height={} levels={:?}",
|
||||
self.lod_enabled.get(),
|
||||
self.viewport_height,
|
||||
levels
|
||||
);
|
||||
for i in 0..n {
|
||||
let off = (i as u64 * LOD_TABLE_SIZE) as usize;
|
||||
let t: LodTable = bytemuck::pod_read_unaligned(
|
||||
&lod_tables_data[off..off + LOD_TABLE_SIZE as usize],
|
||||
);
|
||||
eprintln!("[dbg] lod_table[{}] count={} rows={:?}", i, t.count, t.rows);
|
||||
}
|
||||
eprintln!(
|
||||
"[dbg] pipeline switches (this frame's main pass) = {}",
|
||||
self.debug_pipeline_switches.get()
|
||||
@@ -1090,6 +1232,81 @@ impl Renderer {
|
||||
pub fn set_culling(&self, enabled: bool) {
|
||||
self.cull_enabled.set(enabled);
|
||||
}
|
||||
|
||||
/// Enables or disables LOD (Step 19, D8). When disabled the CPU writes level 0 for every slot
|
||||
/// each frame; level-0 rows carry the full-mesh draw counts, so the GPU indirect args are
|
||||
/// byte-identical to the pre-LOD behavior (the scene renders exactly as before). When enabled
|
||||
/// the CPU picks each slot's level from the projected bounding-sphere radius (with the
|
||||
/// asymmetric hysteresis of `math::lod::lod_level`) and the `cull` pass maps level → args.
|
||||
/// Enabled by default. Inputs: enabled (true = LOD on, false = always level 0).
|
||||
pub fn set_lod_enabled(&self, enabled: bool) {
|
||||
self.lod_enabled.set(enabled);
|
||||
}
|
||||
|
||||
/// Computes the per-slot LOD levels for this frame (Step 19, D8): for each ACTIVE slot, the
|
||||
/// entity's bounding sphere — the **same sphere** the GPU frustum culling uses (D8: bbox
|
||||
/// center + max half-extent × max scale component, rotated by the entity's quaternion) — is
|
||||
/// projected to screen pixels ([`projected_radius_px`]), and [`lod_level`] turns that radius
|
||||
/// into a level with asymmetric hysteresis (D4: the `last` level is the previous frame's
|
||||
/// choice, kept in `self.last_lod_levels`).
|
||||
///
|
||||
/// Inactive (tombstoned) slots and single-level meshes get level 0 (and their hysteresis state
|
||||
/// resets, so a re-added entity starts fresh). The result has one entry per transform slot.
|
||||
fn compute_lod_levels(
|
||||
&self,
|
||||
scene: &Scene,
|
||||
view: &Mat4,
|
||||
proj: &Mat4,
|
||||
transform_slots: &[TransformSlot],
|
||||
bboxes: &[BBoxSlot],
|
||||
) -> Vec<u32> {
|
||||
let height = self.viewport_height.max(1) as f32;
|
||||
let mut last = self.last_lod_levels.borrow_mut();
|
||||
if last.len() != transform_slots.len() {
|
||||
// Entity slots are append-only, but a resize keeps the old levels for the surviving
|
||||
// slots (their hysteresis is meaningful) and zero-fills the new ones.
|
||||
let keep = last.len().min(transform_slots.len());
|
||||
let tail = last.split_off(keep);
|
||||
last.extend(std::iter::repeat(0).take(transform_slots.len() - keep));
|
||||
drop(tail);
|
||||
}
|
||||
let mut levels = vec![0u32; transform_slots.len()];
|
||||
for (i, t) in transform_slots.iter().enumerate() {
|
||||
if !t.is_active() {
|
||||
last[i] = 0;
|
||||
continue; // level 0 (zeroed vec); reset hysteresis for the tombstone
|
||||
}
|
||||
let mesh_idx = t.flags[0] as usize;
|
||||
let mesh = scene.mesh_by_index(mesh_idx as u32);
|
||||
let max_level = (mesh.num_lod_levels() as u32)
|
||||
.saturating_sub(1)
|
||||
.min(MAX_LOD_LEVELS - 1);
|
||||
if max_level == 0 {
|
||||
last[i] = 0;
|
||||
continue; // single-level mesh: always L0
|
||||
}
|
||||
let b = &bboxes[mesh_idx];
|
||||
let center = Vec3::new(
|
||||
(b.min[0] + b.max[0]) * 0.5,
|
||||
(b.min[1] + b.max[1]) * 0.5,
|
||||
(b.min[2] + b.max[2]) * 0.5,
|
||||
);
|
||||
let half = Vec3::new(
|
||||
(b.max[0] - b.min[0]) * 0.5,
|
||||
(b.max[1] - b.min[1]) * 0.5,
|
||||
(b.max[2] - b.min[2]) * 0.5,
|
||||
);
|
||||
// Mirror the WGSL cull pass exactly (D8): radius = |half-extents| × max(scale).
|
||||
let radius = half.length() * t.scale[0].max(t.scale[1].max(t.scale[2]));
|
||||
let center_world =
|
||||
Vec3::from_array(t.translation) + Quat::from_array(t.rotation) * center;
|
||||
let r_px = projected_radius_px(center_world, radius, *view, *proj, height);
|
||||
let lvl = lod_level(r_px, last[i], max_level, &LOD_THRESHOLDS);
|
||||
last[i] = lvl;
|
||||
levels[i] = lvl;
|
||||
}
|
||||
levels
|
||||
}
|
||||
}
|
||||
|
||||
/// Allocates the depth texture + view backing the render passes' `depth_stencil_attachment`
|
||||
|
||||
@@ -326,6 +326,210 @@ impl Geometry {
|
||||
self.validate()?;
|
||||
Ok(self.to_vertices())
|
||||
}
|
||||
|
||||
/// Number of triangles: index count / 3, or vertex count / 3 for non-indexed geometry.
|
||||
pub fn num_triangles(&self) -> usize {
|
||||
match &self.indices {
|
||||
Some(i) => i.len() / 3,
|
||||
None => self.positions.len() / 3,
|
||||
}
|
||||
}
|
||||
|
||||
/// The source triangles (explicit indices, or implicit for non-indexed input), with
|
||||
/// degenerate (zero-area) triangles discarded. Returns `None` when the input is
|
||||
/// malformed (index/vertex count not a multiple of 3) — callers fall back to a clone.
|
||||
fn non_degenerate_triangles(&self) -> Option<Vec<[u32; 3]>> {
|
||||
let vcount = self.positions.len() as u32;
|
||||
let tris: Vec<[u32; 3]> = if let Some(indices) = &self.indices {
|
||||
if indices.len() % 3 != 0 {
|
||||
return None;
|
||||
}
|
||||
indices
|
||||
.chunks_exact(3)
|
||||
.map(|c| [c[0] as u32, c[1] as u32, c[2] as u32])
|
||||
.collect()
|
||||
} else if vcount % 3 != 0 {
|
||||
return None;
|
||||
} else {
|
||||
(0..vcount).step_by(3).map(|i| [i, i + 1, i + 2]).collect()
|
||||
};
|
||||
Some(
|
||||
tris.into_iter()
|
||||
.filter(|t| Self::triangle_area(&self.positions, t) > 1e-8)
|
||||
.collect(),
|
||||
)
|
||||
}
|
||||
|
||||
/// Area of a triangle given its three vertex indices (half the cross-product magnitude).
|
||||
fn triangle_area(positions: &[[f32; 3]], t: &[u32; 3]) -> f32 {
|
||||
let a = positions[t[0] as usize];
|
||||
let b = positions[t[1] as usize];
|
||||
let c = positions[t[2] as usize];
|
||||
let ux = b[0] - a[0];
|
||||
let uy = b[1] - a[1];
|
||||
let uz = b[2] - a[2];
|
||||
let vx = c[0] - a[0];
|
||||
let vy = c[1] - a[1];
|
||||
let vz = c[2] - a[2];
|
||||
let cx = uy * vz - uz * vy;
|
||||
let cy = uz * vx - ux * vz;
|
||||
let cz = ux * vy - uy * vx;
|
||||
(cx * cx + cy * cy + cz * cz).sqrt() * 0.5
|
||||
}
|
||||
|
||||
/// Greedy decimation (Step 19, D10 — pure function, no GPU): returns an **indexed**
|
||||
/// `Geometry` with ≈ `target_triangles` triangles, preserving the silhouette.
|
||||
///
|
||||
/// Algorithm: degenerate (zero-area) triangles are dropped, then the triangles are
|
||||
/// sorted by **ascending area** (stable, tie-broken by original index — deterministic)
|
||||
/// and the `T - target` smallest are removed. Any subset of a mesh's faces is a valid
|
||||
/// (possibly open) mesh — a removed triangle just leaves boundary edges, so no topology
|
||||
/// repair is needed (full manifold preservation would require edge collapse — out of
|
||||
/// scope, see DRAFT Step 19 "Out of scope"). The kept triangles are rebuilt in original
|
||||
/// triangle order:
|
||||
/// - vertices are **welded by exact position equality** (dedup; non-indexed input is
|
||||
/// welded first), and the output is always indexed;
|
||||
/// - smooth normals are **recomputed** over the kept faces (only when the source had
|
||||
/// normals);
|
||||
/// - UVs / colors take the value of the **first vertex of each welded group** encountered
|
||||
/// during the rebuild (documented trade-off: LOD sacrifices UV precision for the
|
||||
/// silhouette).
|
||||
///
|
||||
/// Fallback (never corrupts): target ≥ triangle count, target = 0, or malformed input
|
||||
/// (empty, or counts not multiples of 3) → `self.clone()`.
|
||||
pub fn decimated(&self, target_triangles: u32) -> Geometry {
|
||||
let Some(tris) = self.non_degenerate_triangles() else {
|
||||
return self.clone();
|
||||
};
|
||||
let t = tris.len() as u32;
|
||||
if t == 0 || target_triangles == 0 || target_triangles >= t {
|
||||
return self.clone();
|
||||
}
|
||||
|
||||
// Ascending area (stable; tie-broken by original index) → drop the smallest first.
|
||||
let mut order: Vec<(f32, u32)> = tris
|
||||
.iter()
|
||||
.enumerate()
|
||||
.map(|(i, tri)| (Self::triangle_area(&self.positions, tri), i as u32))
|
||||
.collect();
|
||||
order.sort_by(|a, b| {
|
||||
a.0.partial_cmp(&b.0)
|
||||
.unwrap_or(std::cmp::Ordering::Equal)
|
||||
.then(a.1.cmp(&b.1))
|
||||
});
|
||||
|
||||
let needed = t - target_triangles;
|
||||
let mut removed = vec![false; tris.len()];
|
||||
for i in 0..needed as usize {
|
||||
removed[order[i].1 as usize] = true;
|
||||
}
|
||||
let kept: Vec<[u32; 3]> = tris
|
||||
.into_iter()
|
||||
.enumerate()
|
||||
.filter(|(i, _)| !removed[*i])
|
||||
.map(|(_, tri)| tri)
|
||||
.collect();
|
||||
|
||||
// Rebuild: weld by exact position, remap indices in original triangle order.
|
||||
let mut new_positions: Vec<[f32; 3]> = Vec::with_capacity(self.positions.len());
|
||||
let mut new_uvs: Option<Vec<[f32; 2]>> = self.uvs.is_some().then(Vec::new);
|
||||
let mut new_colors: Option<Vec<[f32; 4]>> = self.colors.is_some().then(Vec::new);
|
||||
// Key = exact float bits (f32 is not Hash/Eq; `to_bits` preserves exact equality).
|
||||
let mut map: std::collections::HashMap<(u32, u32, u32), u32> =
|
||||
std::collections::HashMap::new();
|
||||
let mut new_indices: Vec<u16> = Vec::with_capacity(kept.len() * 3);
|
||||
|
||||
for tri in &kept {
|
||||
for (_k, &vi) in tri.iter().enumerate() {
|
||||
let pos = self.positions[vi as usize];
|
||||
let key = (pos[0].to_bits(), pos[1].to_bits(), pos[2].to_bits());
|
||||
let ni = if let Some(&ni) = map.get(&key) {
|
||||
ni
|
||||
} else {
|
||||
let ni = new_positions.len() as u32;
|
||||
new_positions.push(pos);
|
||||
if let Some(uvs) = &mut new_uvs {
|
||||
uvs.push(self.uvs.as_ref().unwrap()[vi as usize]);
|
||||
}
|
||||
if let Some(colors) = &mut new_colors {
|
||||
colors.push(self.colors.as_ref().unwrap()[vi as usize]);
|
||||
}
|
||||
map.insert(key, ni);
|
||||
ni
|
||||
};
|
||||
new_indices.push(ni as u16);
|
||||
}
|
||||
}
|
||||
|
||||
// Recompute smooth normals over the kept faces (only when the source had normals).
|
||||
let normals = self.normals.as_ref().map(|_| {
|
||||
let mut acc = vec![[0.0f32; 3]; new_positions.len()];
|
||||
// Accumulate with the REMAPPED (post-weld) indices: `kept` holds original indices,
|
||||
// which may not exist in `new_positions` once welding has merged duplicates.
|
||||
for i in 0..kept.len() {
|
||||
let (ia, ib, ic) = (
|
||||
new_indices[i * 3] as usize,
|
||||
new_indices[i * 3 + 1] as usize,
|
||||
new_indices[i * 3 + 2] as usize,
|
||||
);
|
||||
let a = new_positions[ia];
|
||||
let b = new_positions[ib];
|
||||
let c = new_positions[ic];
|
||||
let n = [
|
||||
(b[1] - a[1]) * (c[2] - a[2]) - (b[2] - a[2]) * (c[1] - a[1]),
|
||||
(b[2] - a[2]) * (c[0] - a[0]) - (b[0] - a[0]) * (c[2] - a[2]),
|
||||
(b[0] - a[0]) * (c[1] - a[1]) - (b[1] - a[1]) * (c[0] - a[0]),
|
||||
];
|
||||
for vi in [ia, ib, ic] {
|
||||
let v = &mut acc[vi];
|
||||
v[0] += n[0];
|
||||
v[1] += n[1];
|
||||
v[2] += n[2];
|
||||
}
|
||||
}
|
||||
acc.iter()
|
||||
.map(|v| {
|
||||
let len = (v[0] * v[0] + v[1] * v[1] + v[2] * v[2]).sqrt();
|
||||
if len < 1e-12 {
|
||||
[0.0, 0.0, 1.0]
|
||||
} else {
|
||||
[v[0] / len, v[1] / len, v[2] / len]
|
||||
}
|
||||
})
|
||||
.collect()
|
||||
});
|
||||
|
||||
Geometry {
|
||||
positions: new_positions,
|
||||
normals,
|
||||
uvs: new_uvs,
|
||||
colors: new_colors,
|
||||
indices: Some(new_indices),
|
||||
}
|
||||
}
|
||||
|
||||
/// Generates `levels` (clamped to `1..=MAX_LOD_LEVELS`) LOD geometries from `self`
|
||||
/// (Step 19, D10): level 0 is `self` (**byte-exact**), level k = `decimated(T × 0.5^k)`
|
||||
/// with the target clamped to `[1, T]` (a mesh never decimates below one triangle).
|
||||
/// Cost is setup-only (a few ms for thousands of triangles), never per frame.
|
||||
pub fn generate_lod_levels(&self, levels: u8) -> Vec<Geometry> {
|
||||
let levels = levels.clamp(1, crate::utils::conf::MAX_LOD_LEVELS as u8);
|
||||
let mut out = vec![self.clone()];
|
||||
let Some(tris) = self.non_degenerate_triangles() else {
|
||||
return out;
|
||||
};
|
||||
let t = tris.len() as u32;
|
||||
if t == 0 {
|
||||
return out;
|
||||
}
|
||||
for k in 1..levels {
|
||||
let target = ((t as f32) * 0.5f32.powi(k as i32))
|
||||
.round()
|
||||
.clamp(1.0, t as f32) as u32;
|
||||
out.push(self.decimated(target));
|
||||
}
|
||||
out
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
@@ -465,4 +669,224 @@ mod tests {
|
||||
let geo = quad();
|
||||
assert_eq!(geo.indices(), Some(&[0, 1, 2, 0, 2, 3][..]));
|
||||
}
|
||||
|
||||
// ========================================================================
|
||||
// decimated / generate_lod_levels (Step 19, D10)
|
||||
// ========================================================================
|
||||
|
||||
/// Fan of three triangles sharing vertex 0, with distinct areas 0.5 / 4 / 16.
|
||||
fn tri_fan() -> Geometry {
|
||||
Geometry::new(vec![
|
||||
[0.0, 0.0, 0.0], // 0 (shared)
|
||||
[1.0, 0.0, 0.0], // 1
|
||||
[0.0, 1.0, 0.0], // 2
|
||||
[2.0, 0.0, 0.0], // 3
|
||||
[0.0, 4.0, 0.0], // 4
|
||||
[4.0, 0.0, 0.0], // 5
|
||||
[0.0, 8.0, 0.0], // 6
|
||||
])
|
||||
.with_normals(vec![[0.0, 0.0, 1.0]; 7])
|
||||
.with_indices(vec![0, 1, 2, 0, 3, 4, 0, 5, 6])
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_removes_smallest_triangles_first() {
|
||||
// Areas 0.5 / 4 / 16 → target 2 keeps the two largest.
|
||||
let out = tri_fan().decimated(2);
|
||||
assert_eq!(out.num_triangles(), 2);
|
||||
// The smallest triangle's unique vertices (1, 2) are gone.
|
||||
assert!(!out.positions.contains(&[1.0, 0.0, 0.0]));
|
||||
assert!(!out.positions.contains(&[0.0, 1.0, 0.0]));
|
||||
assert!(out.positions.contains(&[2.0, 0.0, 0.0]));
|
||||
assert!(out.positions.contains(&[0.0, 4.0, 0.0]));
|
||||
out.validate().expect("decimated output validates");
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_icosahedron_20_to_10() {
|
||||
use crate::math::primitives;
|
||||
let ico = primitives::icosphere(1.0, 0); // 12 vertices / 20 faces, equal area
|
||||
assert_eq!(ico.num_triangles(), 20);
|
||||
let out = ico.decimated(10);
|
||||
assert_eq!(out.num_triangles(), 10);
|
||||
out.validate().expect("output validates");
|
||||
// All output positions come from the input; the bbox stays inside the original.
|
||||
for p in &out.positions {
|
||||
assert!(ico.positions.contains(p));
|
||||
}
|
||||
let in_bb = ico.bbox().expect("input bbox");
|
||||
let out_bb = out.bbox().expect("output bbox");
|
||||
for i in 0..3 {
|
||||
assert!(
|
||||
(out_bb.min[i] - in_bb.min[i]).abs() < 1e-5
|
||||
&& out_bb.min[i] >= in_bb.min[i] - 1e-5
|
||||
&& out_bb.max[i] <= in_bb.max[i] + 1e-5
|
||||
);
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_welds_non_indexed_input() {
|
||||
// Four triangles stored NON-indexed (duplicated corners). A and B share the
|
||||
// (0,0,0)-(4,0,0) edge; C and D are smaller and get removed at target 2.
|
||||
let geo = Geometry::new(vec![
|
||||
[0.0, 0.0, 0.0],
|
||||
[4.0, 0.0, 0.0],
|
||||
[0.0, 4.0, 0.0], // tri A (big)
|
||||
[0.0, 0.0, 0.0], // tri B (big; shared edge duplicated)
|
||||
[4.0, 0.0, 0.0],
|
||||
[0.0, -4.0, 0.0],
|
||||
[0.0, 4.0, 0.0],
|
||||
[1.0, 1.0, 0.0],
|
||||
[0.0, -4.0, 0.0], // tri C (small, removed)
|
||||
[1.0, 1.0, 0.0],
|
||||
[2.0, -1.0, 0.0],
|
||||
[0.0, -4.0, 0.0], // tri D (smallest, removed)
|
||||
]);
|
||||
assert!(geo.indices().is_none());
|
||||
// Target 2 keeps the two big triangles; welding merges A and B's duplicated edge corners.
|
||||
let out = geo.decimated(2);
|
||||
assert!(out.indices().is_some(), "output is always indexed");
|
||||
assert_eq!(out.num_triangles(), 2);
|
||||
assert_eq!(out.positions.len(), 4, "12 input vertices weld to 4");
|
||||
assert_eq!(out.indices().unwrap(), &[0, 1, 2, 0, 1, 3]);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_is_deterministic() {
|
||||
use crate::math::primitives;
|
||||
let ico = primitives::icosphere(1.0, 1); // 80 equal-area faces
|
||||
let a = ico.decimated(30);
|
||||
let b = ico.decimated(30);
|
||||
assert_eq!(a.indices(), b.indices());
|
||||
assert_eq!(a.positions, b.positions);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_fallback_target_at_or_above_count() {
|
||||
let geo = tri_fan(); // 3 triangles
|
||||
assert_eq!(
|
||||
geo.decimated(3).indices(),
|
||||
geo.indices(),
|
||||
"target == T → clone"
|
||||
);
|
||||
assert_eq!(
|
||||
geo.decimated(99).indices(),
|
||||
geo.indices(),
|
||||
"target > T → clone"
|
||||
);
|
||||
assert_eq!(
|
||||
geo.decimated(0).indices(),
|
||||
geo.indices(),
|
||||
"target 0 → clone"
|
||||
);
|
||||
// Malformed non-indexed input (vertex count not a multiple of 3) → clone.
|
||||
let bad = Geometry::new(vec![[0.0, 0.0, 0.0], [1.0, 0.0, 0.0]]);
|
||||
assert_eq!(bad.decimated(1).indices(), bad.indices());
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_drops_degenerate_triangles() {
|
||||
// 2 real triangles + 1 degenerate (collinear) one.
|
||||
let geo = Geometry::new(vec![
|
||||
[0.0, 0.0, 0.0],
|
||||
[2.0, 0.0, 0.0],
|
||||
[0.0, 2.0, 0.0], // tri 0: area 2
|
||||
[0.0, 2.0, 0.0], // tri 1: area 2
|
||||
[2.0, 2.0, 0.0],
|
||||
[2.0, 0.0, 0.0],
|
||||
[3.0, 0.0, 0.0], // tri 2: degenerate (collinear)
|
||||
[4.0, 0.0, 0.0],
|
||||
[5.0, 0.0, 0.0],
|
||||
])
|
||||
.with_indices(vec![0, 1, 2, 3, 4, 5, 6, 7, 8]);
|
||||
// T (non-degenerate) = 2; target 1 → one real triangle kept, degenerate dropped.
|
||||
let out = geo.decimated(1);
|
||||
assert_eq!(out.num_triangles(), 1);
|
||||
assert!(
|
||||
!out.positions.contains(&[3.0, 0.0, 0.0]),
|
||||
"degenerate triangle dropped"
|
||||
);
|
||||
out.validate().expect("output validates");
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_recomputes_unit_normals() {
|
||||
// Garbage (unnormalized) input normals; output normals must be unit length.
|
||||
let geo = tri_fan().with_normals(vec![[9.0, 9.0, 9.0]; 7]);
|
||||
let out = geo.decimated(2);
|
||||
let normals = out.normals.expect("normals recomputed");
|
||||
for n in &normals {
|
||||
let len = (n[0] * n[0] + n[1] * n[1] + n[2] * n[2]).sqrt();
|
||||
assert!((len - 1.0).abs() < 1e-5, "normal not unit: {n:?}");
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn decimated_uv_takes_first_encountered() {
|
||||
// Big triangle (area 2, good UVs) + small one (area 0.5, seam UVs) → the small
|
||||
// one is removed; welded UVs come from the kept triangle's vertices.
|
||||
// A and B are the two big triangles (kept); B reuses v0's position as v4 with a
|
||||
// DIFFERENT (seam) UV. C and D are small and removed at target 2. Welding v4 into v0
|
||||
// must keep the FIRST encountered UV (0,0) and drop the seam (9,9).
|
||||
let geo = Geometry::new(vec![
|
||||
[0.0, 0.0, 0.0],
|
||||
[3.0, 0.0, 0.0],
|
||||
[3.0, 3.0, 0.0],
|
||||
[0.0, 3.0, 0.0],
|
||||
[0.0, 0.0, 0.0], // seam copy of v0
|
||||
[1.5, 1.5, 0.0],
|
||||
])
|
||||
.with_uvs(vec![
|
||||
[0.0, 0.0],
|
||||
[1.0, 0.0],
|
||||
[1.0, 1.0],
|
||||
[0.0, 1.0],
|
||||
[9.0, 9.0], // seam UVs on the duplicate corner
|
||||
[8.0, 8.0],
|
||||
])
|
||||
.with_indices(vec![0, 1, 2, 4, 2, 3, 3, 4, 5, 5, 1, 2]);
|
||||
let out = geo.decimated(2); // keeps the two big triangles
|
||||
assert_eq!(out.num_triangles(), 2);
|
||||
assert_eq!(
|
||||
out.positions.len(),
|
||||
4,
|
||||
"seam corner welds into its first occurrence"
|
||||
);
|
||||
assert_eq!(
|
||||
out.uvs.as_deref(),
|
||||
Some(&[[0.0, 0.0], [1.0, 0.0], [1.0, 1.0], [0.0, 1.0]][..])
|
||||
);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn generate_lod_levels_halving_ratios() {
|
||||
use crate::math::primitives;
|
||||
let geo = primitives::icosphere(1.0, 1); // 80 triangles
|
||||
let levels = geo.generate_lod_levels(3);
|
||||
assert_eq!(levels.len(), 3);
|
||||
assert_eq!(levels[0].num_triangles(), 80, "L0 exact");
|
||||
assert_eq!(levels[0].indices(), geo.indices(), "L0 byte-exact");
|
||||
assert_eq!(levels[1].num_triangles(), 40);
|
||||
assert_eq!(levels[2].num_triangles(), 20);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn generate_lod_levels_one_level() {
|
||||
let geo = tri_fan();
|
||||
let levels = geo.generate_lod_levels(1);
|
||||
assert_eq!(levels.len(), 1);
|
||||
assert_eq!(levels[0].indices(), geo.indices());
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn generate_lod_levels_tiny_mesh_clamps_to_one() {
|
||||
let geo = quad(); // 2 triangles, equal area
|
||||
let levels = geo.generate_lod_levels(4);
|
||||
assert_eq!(levels.len(), 4);
|
||||
assert_eq!(levels[0].num_triangles(), 2);
|
||||
assert_eq!(levels[1].num_triangles(), 1, "2 × 0.5");
|
||||
assert_eq!(levels[2].num_triangles(), 1, "2 × 0.25 → clamped to 1");
|
||||
assert_eq!(levels[3].num_triangles(), 1, "2 × 0.125 → clamped to 1");
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,248 @@
|
||||
//! # LOD — Per-frame Level Selection (pure, testable without a GPU)
|
||||
//!
|
||||
//! The pure functions behind the LOD feature (Step 19, D1/D4/D8). Each frame the **CPU**
|
||||
//! decides which detail level every entity draws; these functions do that math:
|
||||
//!
|
||||
//! - [`projected_radius_px`]: the entity's *perceived size* — its bounding-sphere radius in
|
||||
//! screen pixels (the **same sphere** the GPU frustum culling uses, D8);
|
||||
//! - [`lod_level`]: the level decision with **asymmetric hysteresis** (D4) — the core
|
||||
//! anti-flicker mechanism.
|
||||
//!
|
||||
//! Both are pure (no GPU, no state beyond the caller-supplied `last` level) → unit-testable.
|
||||
//! The Renderer calls them per slot each frame and uploads the resulting levels to the GPU,
|
||||
//! which only maps level → draw args (the packed-buffer offsets live in the per-mesh LOD
|
||||
//! table — see `resources::uniform::LodTable`).
|
||||
|
||||
use glam::{Mat4, Vec3, Vec4};
|
||||
|
||||
/// Projected radius (in **pixels**) of a bounding sphere, given the camera's view/projection.
|
||||
///
|
||||
/// The sphere center (world space) is transformed into view space; a sphere at depth `d` with
|
||||
/// radius `r` subtends `r / d` in view space, which the projection's vertical scale
|
||||
/// (`proj.y.y = 1 / tan(fov / 2)`) maps to NDC — multiplied by `height_px / 2` (half the
|
||||
/// viewport height in pixels) gives pixels.
|
||||
///
|
||||
/// A sphere whose center is inside/behind the near plane (`depth <= 1e-4`) returns
|
||||
/// `f32::INFINITY` — the entity dominates the screen, so the finest level (0) is chosen.
|
||||
pub fn projected_radius_px(
|
||||
center_world: Vec3,
|
||||
radius: f32,
|
||||
view: Mat4,
|
||||
proj: Mat4,
|
||||
height_px: f32,
|
||||
) -> f32 {
|
||||
let v = view * Vec4::new(center_world.x, center_world.y, center_world.z, 1.0);
|
||||
let depth = -v.z; // view space: the camera looks along -Z (glam `look_at_mat4`)
|
||||
if depth <= 1e-4 {
|
||||
return f32::INFINITY;
|
||||
}
|
||||
(radius / depth) * proj.y_axis.y * (height_px * 0.5)
|
||||
}
|
||||
|
||||
/// Level decision with **asymmetric hysteresis** (Step 19, D4).
|
||||
///
|
||||
/// `thresholds` is a **descending** pixel radius: `thresholds[k]` is the radius *above which*
|
||||
/// level k+1 is required (i.e. level k is sufficient up to that bound; level 0 has no bound).
|
||||
/// Levels beyond the threshold count share the last bound (clamped) — e.g. with `[48, 12]`
|
||||
/// only the first three levels are distinct.
|
||||
///
|
||||
/// Hysteresis (dead band):
|
||||
/// - to a **finer** level: immediate, as soon as `radius_px` exceeds the current level's bound;
|
||||
/// - to a **coarser** level: only if `radius_px <= bound(k) * 0.8` (20 % dead band), stepped
|
||||
/// incrementally (each intermediate bound × 0.8 must hold).
|
||||
///
|
||||
/// The "detail loss" pop (going coarser) is therefore delayed; the "detail regain" pop (going
|
||||
/// finer) is immediate — standard engine practice. `f32::INFINITY` (object at the camera)
|
||||
/// always returns 0. The result is always within `0..=max_level`.
|
||||
pub fn lod_level(radius_px: f32, last: u32, max_level: u32, thresholds: &[f32]) -> u32 {
|
||||
if radius_px.is_infinite() || max_level == 0 || thresholds.is_empty() {
|
||||
return 0;
|
||||
}
|
||||
// Bound for level k+1: the k-th threshold, clamped for levels beyond the threshold count.
|
||||
let bound = |k: u32| thresholds[(k as usize).min(thresholds.len() - 1)];
|
||||
let last = (last as usize).min(max_level as usize) as u32;
|
||||
|
||||
// Target without hysteresis: the coarsest level whose bound is still satisfied.
|
||||
let mut target = 0u32;
|
||||
let mut k = 0u32;
|
||||
while k < max_level {
|
||||
if radius_px <= bound(k) {
|
||||
target = k + 1;
|
||||
k += 1;
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if target <= last {
|
||||
// Finer or equal: immediate (no dead band on the way to more detail).
|
||||
target
|
||||
} else {
|
||||
// Coarser: 20 % dead band per step, incremental.
|
||||
let mut lvl = last;
|
||||
while lvl < target {
|
||||
if radius_px <= bound(lvl) * 0.8 {
|
||||
lvl += 1;
|
||||
} else {
|
||||
break;
|
||||
}
|
||||
}
|
||||
lvl
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
use glam::Mat4;
|
||||
use glam::Vec3;
|
||||
|
||||
/// A camera at `(0, 0, dist)` looking at the origin, up `+Y`, with vertical `fov`.
|
||||
fn camera(dist: f32, fov: f32) -> (Mat4, Mat4) {
|
||||
let view =
|
||||
glam::camera::rh::view::look_at_mat4(Vec3::new(0.0, 0.0, dist), Vec3::ZERO, Vec3::Y);
|
||||
let proj = glam::Mat4::perspective_rh_gl(fov, 1.0, 0.1, 100.0);
|
||||
(view, proj)
|
||||
}
|
||||
|
||||
// ========================================================================
|
||||
// projected_radius_px
|
||||
// ========================================================================
|
||||
|
||||
#[test]
|
||||
fn projected_radius_analytic() {
|
||||
// Sphere of radius 1 at the origin; camera 5 units away; fov = 90°
|
||||
// (proj vertical scale = 1/tan(45°) = 1); viewport 1000 px tall.
|
||||
// Expected: (1 / 5) * 1 * 500 = 100 px.
|
||||
let (view, proj) = camera(5.0, std::f32::consts::PI / 2.0);
|
||||
let r = projected_radius_px(Vec3::ZERO, 1.0, view, proj, 1000.0);
|
||||
assert!((r - 100.0).abs() < 1e-3, "expected 100 px, got {r}");
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn projected_radius_scale_invariance() {
|
||||
// 10x bigger object 10x further away → same projected radius (similarity).
|
||||
let (view1, proj1) = camera(5.0, std::f32::consts::PI / 2.0);
|
||||
let (view2, proj2) = camera(50.0, std::f32::consts::PI / 2.0);
|
||||
let r1 = projected_radius_px(Vec3::ZERO, 1.0, view1, proj1, 1000.0);
|
||||
let r2 = projected_radius_px(Vec3::ZERO, 10.0, view2, proj2, 1000.0);
|
||||
assert!((r1 - r2).abs() < 1e-2, "expected equal, got {r1} vs {r2}");
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn projected_radius_at_camera_is_infinite() {
|
||||
// Center at the camera position → depth 0 → INFINITY (finest level).
|
||||
let (view, proj) = camera(5.0, std::f32::consts::PI / 2.0);
|
||||
let r = projected_radius_px(Vec3::new(0.0, 0.0, 5.0), 1.0, view, proj, 1000.0);
|
||||
assert!(r.is_infinite());
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn projected_radius_behind_camera_is_infinite() {
|
||||
// Center behind the camera → negative depth → INFINITY.
|
||||
let (view, proj) = camera(5.0, std::f32::consts::PI / 2.0);
|
||||
let r = projected_radius_px(Vec3::new(0.0, 0.0, 20.0), 1.0, view, proj, 1000.0);
|
||||
assert!(r.is_infinite());
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn projected_radius_narrower_fov_larger_pixels() {
|
||||
// Narrower FOV (zoomed in) → LARGER vertical projection scale (1/tan(fov/2)) →
|
||||
// more pixels for the same sphere at the same distance.
|
||||
let fov_narrow = std::f32::consts::PI / 3.0; // 60°
|
||||
let fov_wide = std::f32::consts::PI / 2.0; // 90°
|
||||
let (v1, p1) = camera(5.0, fov_narrow);
|
||||
let (v2, p2) = camera(5.0, fov_wide);
|
||||
let r1 = projected_radius_px(Vec3::ZERO, 1.0, v1, p1, 1000.0);
|
||||
let r2 = projected_radius_px(Vec3::ZERO, 1.0, v2, p2, 1000.0);
|
||||
assert!(
|
||||
r1 > r2,
|
||||
"narrower FOV should give more pixels: {r1} vs {r2}"
|
||||
);
|
||||
}
|
||||
|
||||
// ========================================================================
|
||||
// lod_level
|
||||
// ========================================================================
|
||||
|
||||
#[test]
|
||||
fn lod_level_simple_thresholds() {
|
||||
let t = [48.0f32, 12.0];
|
||||
// r > 48 → level 0 (too big for any coarser level).
|
||||
assert_eq!(lod_level(100.0, 0, 2, &t), 0);
|
||||
assert_eq!(lod_level(52.0, 0, 2, &t), 0);
|
||||
// 48 >= r > 38.4 (0.8·48): target is L1, but the dead band holds it at L0.
|
||||
assert_eq!(lod_level(44.0, 0, 2, &t), 0);
|
||||
// r <= 38.4 → L1.
|
||||
assert_eq!(lod_level(38.4, 0, 2, &t), 1);
|
||||
assert_eq!(lod_level(30.0, 0, 2, &t), 1);
|
||||
// 12 > r > 9.6 (0.8·12): target L2, dead band holds at L1.
|
||||
assert_eq!(lod_level(10.0, 0, 2, &t), 1);
|
||||
// r <= 9.6 → L2 (both steps pass the band).
|
||||
assert_eq!(lod_level(9.6, 0, 2, &t), 2);
|
||||
assert_eq!(lod_level(9.0, 0, 2, &t), 2);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn lod_level_finer_is_immediate() {
|
||||
let t = [48.0f32, 12.0];
|
||||
// Already coarse (L2); radius grows past 48 → immediately back to L0.
|
||||
assert_eq!(lod_level(100.0, 2, 2, &t), 0);
|
||||
// L2, radius between the bounds → immediately to L1.
|
||||
assert_eq!(lod_level(30.0, 2, 2, &t), 1);
|
||||
// L1, radius past 48 → immediately to L0.
|
||||
assert_eq!(lod_level(52.0, 1, 2, &t), 0);
|
||||
// L1, radius below 12 → target L2 but dead band (10 > 9.6) holds at L1.
|
||||
assert_eq!(lod_level(10.0, 1, 2, &t), 1);
|
||||
// L1, radius below 9.6 → L2.
|
||||
assert_eq!(lod_level(9.0, 1, 2, &t), 2);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn lod_level_oscillation_is_stable() {
|
||||
// Anti-flicker (D4): a radius oscillating ±10 % around threshold 48 (43.2..52.8)
|
||||
// must not make the level flip back and forth.
|
||||
let t = [48.0f32];
|
||||
let mut level = 0u32;
|
||||
for _ in 0..100 {
|
||||
for r in [43.2f32, 52.8, 43.2, 52.8] {
|
||||
level = lod_level(r, level, 2, &t);
|
||||
}
|
||||
}
|
||||
// Whatever level it settled on, it must not have changed on the last pass.
|
||||
let before = level;
|
||||
for r in [43.2f32, 52.8, 43.2, 52.8] {
|
||||
level = lod_level(r, level, 2, &t);
|
||||
}
|
||||
assert_eq!(before, level, "level flickered around the threshold");
|
||||
// From L0 the oscillation never leaves L0 (coarser needs r ≤ 38.4).
|
||||
assert_eq!(lod_level(43.2, 0, 2, &t), 0);
|
||||
assert_eq!(lod_level(52.8, 0, 2, &t), 0);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn lod_level_clamped_thresholds_for_extra_levels() {
|
||||
// 4 levels but only 2 thresholds: levels 2 and 3 share the last bound (12).
|
||||
let t = [48.0f32, 12.0];
|
||||
// r = 9 passes both bands (38.4, 9.6) AND the clamped third bound (0.8·12) → L3.
|
||||
assert_eq!(lod_level(9.0, 0, 3, &t), 3);
|
||||
// r = 10 passes the first two targets but the clamped band holds at L2.
|
||||
assert_eq!(lod_level(10.0, 0, 3, &t), 1);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn lod_level_infinite_returns_zero() {
|
||||
let t = [48.0f32, 12.0];
|
||||
assert_eq!(lod_level(f32::INFINITY, 2, 2, &t), 0);
|
||||
assert_eq!(lod_level(f32::INFINITY, 0, 2, &t), 0);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn lod_level_degenerate_inputs() {
|
||||
let t = [48.0f32];
|
||||
assert_eq!(lod_level(1.0, 5, 0, &t), 0); // max_level 0
|
||||
assert_eq!(lod_level(1.0, 0, 2, &[]), 0); // no thresholds
|
||||
// Stale `last` beyond max_level is clamped, not a panic.
|
||||
assert_eq!(lod_level(100.0, 9, 2, &t), 0);
|
||||
}
|
||||
}
|
||||
@@ -15,14 +15,17 @@
|
||||
//! - `transform.rs`: Defines the `Transform` struct and its conversion to matrix form
|
||||
//! - `geometry.rs`: Defines the `Geometry` struct for mesh data storage
|
||||
//! - `primitives.rs`: Procedural mesh generators (cube, sphere, cylinder, cone, torus…) returning `Geometry`
|
||||
//! - `lod.rs`: Pure LOD level-selection functions (projected radius + hysteresis, Step 19)
|
||||
|
||||
pub mod frustum;
|
||||
pub mod geometry;
|
||||
pub mod lod;
|
||||
pub mod primitives;
|
||||
pub mod transform;
|
||||
|
||||
// Re-exports
|
||||
pub use frustum::Frustum;
|
||||
pub use geometry::{BBox, Geometry, GeometryError};
|
||||
pub use lod::{lod_level, projected_radius_px};
|
||||
pub use primitives::{cone, cube, cylinder, icosphere, plane, torus, uv_sphere};
|
||||
pub use transform::Transform;
|
||||
|
||||
+307
-44
@@ -14,84 +14,209 @@
|
||||
//! - **CPU+GPU retention (DRAFT Step 8, D5)**: `geometry` (CPU) and the vertex/index buffers (GPU) coexist.
|
||||
//! The GPU buffers are uploaded once at creation; the `Arc<Geometry>` is kept for CPU-side computations
|
||||
//! without re-uploading per frame.
|
||||
//! - **LOD packing (Step 19, D7)**: a multi-level mesh packs ALL its levels into **one** vertex buffer and
|
||||
//! **one** index buffer (level k lives at a byte offset), because WebGPU forbids dynamic offsets on
|
||||
//! vertex/index bindings — only the draw ARGS move per frame. The per-level offsets live in the
|
||||
//! `LodRow`s (uploaded to the GPU LOD table); the shadow/main passes always bind level 0.
|
||||
//!
|
||||
//! ## Construction (DRAFT Step 8, D4)
|
||||
//! The single canonical constructor is [`Mesh::from_geometry`]. The former `Mesh::new`/`Mesh::with_material`
|
||||
//! ## Construction (DRAFT Step 8, D4; Step 19, D6/D7)
|
||||
//! The single canonical constructor is [`Mesh::from_geometry_lod`] (levels + mode); [`Mesh::from_geometry`]
|
||||
//! is its one-level convenience wrapper. The former `Mesh::new`/`Mesh::with_material`
|
||||
//! (which took raw `&[Vertex]`) were removed in Step 8: the `Scene` declares meshes from a `Geometry`, and
|
||||
//! `Mesh` derives its interleaved vertices internally via `Geometry::to_vertices()`.
|
||||
|
||||
use crate::math::Geometry;
|
||||
use crate::resources::Material;
|
||||
use crate::resources::Vertex;
|
||||
use crate::resources::uniform::LodRow;
|
||||
use std::sync::Arc;
|
||||
use wgpu::util::DeviceExt;
|
||||
|
||||
/// How a mesh's LOD levels were produced (Step 19, D6).
|
||||
#[derive(Debug, Clone, Copy, PartialEq, Eq, Default)]
|
||||
pub enum LodMode {
|
||||
/// No LOD: a single level (the plain `create_mesh` path).
|
||||
#[default]
|
||||
Off,
|
||||
/// Levels 1.. were auto-generated by greedy decimation (`Geometry::decimated`, D10).
|
||||
Auto,
|
||||
/// Levels 1.. were supplied explicitly (`Scene::add_mesh_lod`).
|
||||
Explicit,
|
||||
}
|
||||
|
||||
/// Persistent GPU geometry: vertex positions, optional indices, draw call counters, and the retained
|
||||
/// CPU `Geometry` (Step 8). Created once via `Mesh::from_geometry()` during scene setup; referenced by
|
||||
/// Renderer for every frame. A Mesh optionally references the `Material` used to render it (`Option<Arc<Material>>`).
|
||||
/// When `material()` is `None`, the `Scene` supplies its default material at draw time (DRAFT Step 7.3.5).
|
||||
///
|
||||
/// Since Step 19 a Mesh may carry several LOD levels: they are packed into the single vertex/index
|
||||
/// buffers (D7) and described by `lod_rows` (uploaded per mesh as a [`crate::resources::uniform::LodTable`]).
|
||||
pub struct Mesh {
|
||||
/// Shared CPU geometry this mesh was built from (Step 8, D5). Retained for CPU-side computation
|
||||
/// Shared CPU geometry of level 0 (the full mesh, Step 8, D5). Retained for CPU-side computation
|
||||
/// (bounding boxes, UV access, normal queries) and shared across meshes with identical geometry.
|
||||
geometry: Arc<Geometry>,
|
||||
/// GPU buffer containing vertex attribute data (position, UV, color).
|
||||
/// GPU buffer containing the packed vertex data of ALL levels (one `Vertex` per position, levels
|
||||
/// concatenated in level order; level 0 first).
|
||||
pub vertex_buffer: wgpu::Buffer,
|
||||
/// Optional GPU buffer for indexed drawing. Present when the mesh uses index-based rendering instead of simple vertex iteration.
|
||||
/// Optional GPU buffer with the packed indices of all indexed levels (rebased onto the packed
|
||||
/// vertex layout). `None` when no level is indexed.
|
||||
pub index_buffer: Option<wgpu::Buffer>,
|
||||
/// Number of vertices in the mesh. Used as `0..num_vertices` for non-indexed draws.
|
||||
/// Number of vertices of level 0. Used as `0..num_vertices` for non-indexed draws.
|
||||
pub num_vertices: u32,
|
||||
/// Number of indices in the index buffer. Used as `0..num_indices` for indexed draws.
|
||||
/// Number of indices of level 0. Used as `0..num_indices` for indexed draws.
|
||||
pub num_indices: u32,
|
||||
/// The Material used to render this mesh. `None` until assigned; the Renderer falls back to the
|
||||
/// Scene's default material when absent (DRAFT Step 7.3.5).
|
||||
material: Option<Arc<Material>>,
|
||||
/// How the levels were produced (Step 19, D6).
|
||||
lod_mode: LodMode,
|
||||
/// The CPU geometry of every level, level 0 first (all retained; `geometry` is `lod_levels[0]`).
|
||||
lod_levels: Vec<Arc<Geometry>>,
|
||||
/// The packed-buffer offsets per level (mirrors the uploaded per-mesh LOD table).
|
||||
lod_rows: Vec<LodRow>,
|
||||
}
|
||||
|
||||
/// Error returned by [`pack_levels`] when the packed vertex total exceeds the u16 index range.
|
||||
#[derive(Debug, Clone, Copy, PartialEq, Eq)]
|
||||
pub struct PackError {
|
||||
/// Total packed vertices (all levels) — must be < 65536 for u16 rebased indices.
|
||||
pub total_vertices: u32,
|
||||
}
|
||||
|
||||
impl std::fmt::Display for PackError {
|
||||
fn fmt(&self, f: &mut std::fmt::Formatter<'_>) -> std::fmt::Result {
|
||||
write!(
|
||||
f,
|
||||
"packed LOD vertex total {} exceeds the 65535 u16 index limit",
|
||||
self.total_vertices
|
||||
)
|
||||
}
|
||||
}
|
||||
|
||||
impl std::error::Error for PackError {}
|
||||
|
||||
/// Pure packing of LOD levels (Step 19, D7) — no GPU, unit-testable.
|
||||
///
|
||||
/// Concatenates every level's interleaved vertices into one flat slice (level 0 first) and rebases
|
||||
/// every indexed level's indices onto the packed vertex layout. Returns the packed vertices, the
|
||||
/// packed indices (`None` when no level is indexed), and one [`LodRow`] per level carrying its
|
||||
/// **element** offsets + counts. Offsets are in ELEMENT units (vertex indices / index elements),
|
||||
/// not bytes: the GPU cull pass copies them straight into the WebGPU indirect draw args, whose
|
||||
/// `first_vertex`/`first_index` fields are element indices — the packed buffers are bound in full
|
||||
/// (offset 0) and only the args' first_* fields move per level.
|
||||
///
|
||||
/// Fails with [`PackError`] when the **total** packed vertex count reaches 65536 (u16 rebased
|
||||
/// indices cannot address it). Callers (the Scene APIs) validate the total beforehand and report
|
||||
/// the error to the user; a single-level mesh can never fail (its indices are already u16).
|
||||
pub(crate) fn pack_levels(
|
||||
levels: &[Arc<Geometry>],
|
||||
) -> Result<(Vec<Vertex>, Option<Vec<u16>>, Vec<LodRow>), PackError> {
|
||||
let total: u32 = levels.iter().map(|l| l.positions.len() as u32).sum();
|
||||
if total >= 65536 {
|
||||
return Err(PackError {
|
||||
total_vertices: total,
|
||||
});
|
||||
}
|
||||
|
||||
let mut vertices: Vec<Vertex> = Vec::new();
|
||||
let mut indices: Vec<u16> = Vec::new();
|
||||
let mut any_indexed = false;
|
||||
let mut rows: Vec<LodRow> = Vec::with_capacity(levels.len());
|
||||
|
||||
let mut vertex_base = 0u32; // element (vertex index) base of the level in the packed buffer
|
||||
let mut index_base = 0u32; // element (index element) base of the level in the packed buffer
|
||||
for level in levels {
|
||||
let level_vertices = level.to_vertices();
|
||||
// Element units (NOT bytes): the row's offsets feed the WebGPU indirect draw args
|
||||
// (first_vertex = vertex index, first_index = index element) — see the doc above.
|
||||
let level_vertex_offset = vertex_base;
|
||||
let level_vertex_count = level_vertices.len() as u32;
|
||||
|
||||
let (level_index_offset, level_index_count) = match level.indices() {
|
||||
Some(data) => {
|
||||
any_indexed = true;
|
||||
for &idx in data {
|
||||
// Safe: the total-vertex check above guarantees no u16 overflow.
|
||||
indices.push(idx as u32 as u16 + vertex_base as u16);
|
||||
}
|
||||
(index_base, data.len() as u32)
|
||||
}
|
||||
None => (0, 0),
|
||||
};
|
||||
|
||||
rows.push(LodRow::new(
|
||||
level_vertex_offset,
|
||||
level_vertex_count,
|
||||
level_index_offset,
|
||||
level_index_count,
|
||||
));
|
||||
vertex_base += level_vertex_count;
|
||||
index_base += level_index_count;
|
||||
vertices.extend(level_vertices);
|
||||
}
|
||||
|
||||
Ok((vertices, any_indexed.then_some(indices), rows))
|
||||
}
|
||||
|
||||
impl Mesh {
|
||||
/// Canonical constructor (Step 8, D4): builds GPU buffers from a shared CPU `Geometry`.
|
||||
/// Canonical constructor (Step 8, D4, now D6/D7): builds the packed GPU buffers from a list of
|
||||
/// LOD levels (level 0 = the full mesh, always present).
|
||||
///
|
||||
/// Inputs: device (GPU command source for buffer creation), geometry (shared CPU vertex data to
|
||||
/// upload), material (optional appearance; `None` falls back to the Scene default at draw time).
|
||||
///
|
||||
/// Internal steps: 1) derive interleaved `Vertex` array via `geometry.to_vertices()`; 2) create the
|
||||
/// vertex buffer (one `Vertex` per position); 3) if the geometry has indices, create the index buffer
|
||||
/// and set `num_indices`, else leave it `None`.
|
||||
///
|
||||
/// The provided `geometry` is retained on the mesh (`geometry` accessor) alongside the uploaded GPU
|
||||
/// buffers, so the CPU data remains readable for later phases without re-uploading each frame (DRAFT Step 8, D5).
|
||||
/// All levels are packed into ONE vertex buffer and ONE index buffer (D7 — WebGPU forbids dynamic
|
||||
/// offsets on vertex/index bindings; only the draw args move). `num_vertices`/`num_indices`
|
||||
/// describe **level 0** (the shadow and main passes always bind level 0); the per-level draw
|
||||
/// arguments are emitted by the GPU cull pass from the uploaded LOD table.
|
||||
pub fn from_geometry_lod(
|
||||
device: &wgpu::Device,
|
||||
levels: Vec<Arc<Geometry>>,
|
||||
material: Option<Arc<Material>>,
|
||||
mode: LodMode,
|
||||
) -> Self {
|
||||
assert!(!levels.is_empty(), "a mesh needs at least level 0");
|
||||
// Packing can only fail when the packed vertex total reaches 65536; the Scene APIs
|
||||
// validate that beforehand. A single level (from_geometry) can never fail.
|
||||
let (vertices, indices, rows) =
|
||||
pack_levels(&levels).expect("packed LOD vertex total exceeds the u16 limit");
|
||||
|
||||
let vertex_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("Mesh Vertex Buffer (packed LOD)"),
|
||||
contents: bytemuck::cast_slice(&vertices),
|
||||
usage: wgpu::BufferUsages::VERTEX | wgpu::BufferUsages::COPY_SRC,
|
||||
});
|
||||
|
||||
let index_buffer = indices.map(|data| {
|
||||
device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("Mesh Index Buffer (packed LOD)"),
|
||||
usage: wgpu::BufferUsages::INDEX | wgpu::BufferUsages::COPY_SRC,
|
||||
contents: bytemuck::cast_slice(&data),
|
||||
})
|
||||
});
|
||||
|
||||
let l0 = &levels[0];
|
||||
Self {
|
||||
geometry: Arc::clone(l0),
|
||||
vertex_buffer,
|
||||
index_buffer,
|
||||
num_vertices: l0.positions.len() as u32,
|
||||
num_indices: l0.indices().map(|i| i.len() as u32).unwrap_or(0),
|
||||
material,
|
||||
lod_mode: mode,
|
||||
lod_levels: levels,
|
||||
lod_rows: rows,
|
||||
}
|
||||
}
|
||||
|
||||
/// One-level convenience constructor (Step 8, D4): builds GPU buffers from a shared CPU
|
||||
/// `Geometry` (no LOD — `LodMode::Off`, `lod_rows` has a single row).
|
||||
pub fn from_geometry(
|
||||
device: &wgpu::Device,
|
||||
geometry: Arc<Geometry>,
|
||||
material: Option<Arc<Material>>,
|
||||
) -> Self {
|
||||
let vertices = geometry.to_vertices();
|
||||
let vertex_buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("Mesh Vertex Buffer"),
|
||||
contents: bytemuck::cast_slice(&vertices),
|
||||
usage: wgpu::BufferUsages::VERTEX,
|
||||
});
|
||||
|
||||
let (index_buffer, num_indices) = if let Some(data) = geometry.indices() {
|
||||
let buffer = device.create_buffer_init(&wgpu::util::BufferInitDescriptor {
|
||||
label: Some("Mesh Index Buffer"),
|
||||
usage: wgpu::BufferUsages::INDEX,
|
||||
contents: bytemuck::cast_slice(data),
|
||||
});
|
||||
(Some(buffer), data.len() as u32)
|
||||
} else {
|
||||
(None, 0)
|
||||
};
|
||||
|
||||
Self {
|
||||
geometry,
|
||||
vertex_buffer,
|
||||
index_buffer,
|
||||
num_vertices: vertices.len() as u32,
|
||||
num_indices,
|
||||
material,
|
||||
}
|
||||
Self::from_geometry_lod(device, vec![geometry], material, LodMode::Off)
|
||||
}
|
||||
|
||||
/// Returns a reference to the shared CPU geometry this mesh was built from (Step 8, D5).
|
||||
/// Returns a reference to the level-0 CPU geometry this mesh was built from (Step 8, D5).
|
||||
/// Read-only accessor for CPU-side queries (bounding boxes, UVs, normals).
|
||||
pub fn geometry(&self) -> &Arc<Geometry> {
|
||||
&self.geometry
|
||||
@@ -108,4 +233,142 @@ impl Mesh {
|
||||
pub fn set_material(&mut self, material: Arc<Material>) {
|
||||
self.material = Some(material);
|
||||
}
|
||||
|
||||
/// How this mesh's LOD levels were produced (Step 19, D6).
|
||||
pub fn lod_mode(&self) -> LodMode {
|
||||
self.lod_mode
|
||||
}
|
||||
|
||||
/// Number of LOD levels (1 for a plain mesh).
|
||||
pub fn num_lod_levels(&self) -> usize {
|
||||
self.lod_rows.len()
|
||||
}
|
||||
|
||||
/// The per-level packed-buffer rows (level 0 first).
|
||||
pub fn lod_rows(&self) -> &[LodRow] {
|
||||
&self.lod_rows
|
||||
}
|
||||
|
||||
/// The CPU geometry of level k (0 = full mesh). Used by `Scene::add_mesh_lod` when
|
||||
/// reconstructing a mesh with a modified level list.
|
||||
pub fn lod_levels_arc(&self, k: usize) -> Arc<Geometry> {
|
||||
Arc::clone(&self.lod_levels[k])
|
||||
}
|
||||
|
||||
/// The per-mesh GPU LOD table (uploaded once by the Renderer; read by the GPU cull pass).
|
||||
pub fn lod_table(&self) -> crate::resources::uniform::LodTable {
|
||||
crate::resources::uniform::LodTable::from_rows(&self.lod_rows)
|
||||
}
|
||||
|
||||
/// Whether **level 0** is indexed (the shadow/main passes always bind level 0).
|
||||
pub fn l0_indexed(&self) -> bool {
|
||||
self.lod_rows
|
||||
.first()
|
||||
.map(|r| r.index_count > 0)
|
||||
.unwrap_or(false)
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
use crate::math::primitives;
|
||||
|
||||
#[test]
|
||||
fn pack_levels_offsets_and_rebasing() {
|
||||
// Two levels: L0 = icosahedron (12 verts / 60 indices), L1 = decimated to 10 (welded).
|
||||
let l0 = Arc::new(primitives::icosphere(1.0, 0));
|
||||
let l1 = Arc::new(l0.decimated(10));
|
||||
let (vertices, indices, rows) =
|
||||
pack_levels(&[Arc::clone(&l0), Arc::clone(&l1)]).expect("small mesh packs");
|
||||
|
||||
// Packed vertices = L0 + L1 concatenated.
|
||||
assert_eq!(vertices.len(), l0.positions.len() + l1.positions.len());
|
||||
// Packed indices = 60 + L1's (rebased by L0's vertex count).
|
||||
let packed_indices = indices.expect("both levels indexed");
|
||||
assert_eq!(packed_indices.len(), 60 + l1.indices().unwrap().len());
|
||||
// L0 indices unchanged (rebase base 0); L1 indices rebased by 12.
|
||||
for (i, idx) in l0.indices().unwrap().iter().enumerate() {
|
||||
assert_eq!(packed_indices[i], *idx);
|
||||
}
|
||||
for (j, idx) in l1.indices().unwrap().iter().enumerate() {
|
||||
assert_eq!(packed_indices[60 + j], *idx + 12);
|
||||
}
|
||||
|
||||
// Rows carry the ELEMENT offsets (vertex indices / index elements, not bytes).
|
||||
assert_eq!(rows.len(), 2);
|
||||
assert_eq!(rows[0].vertex_offset, 0);
|
||||
assert_eq!(rows[0].vertex_count, 12);
|
||||
assert_eq!(rows[0].index_offset, 0);
|
||||
assert_eq!(rows[0].index_count, 60);
|
||||
assert_eq!(rows[1].vertex_offset, 12);
|
||||
assert_eq!(rows[1].vertex_count, l1.positions.len() as u32);
|
||||
assert_eq!(rows[1].index_offset, 60);
|
||||
assert_eq!(rows[1].index_count, l1.indices().unwrap().len() as u32);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn pack_levels_mixed_indexedness() {
|
||||
// Non-indexed L0 (6 verts) + indexed L1 (welded) — the Auto-mode case from DRAFT D7.
|
||||
let l0 = Arc::new(
|
||||
Geometry::new(vec![
|
||||
[0.0, 0.0, 0.0],
|
||||
[2.0, 0.0, 0.0],
|
||||
[0.0, 2.0, 0.0],
|
||||
[0.0, 0.0, 0.0],
|
||||
[2.0, 0.0, 0.0],
|
||||
[0.0, -2.0, 0.0],
|
||||
])
|
||||
.with_indices(vec![0, 1, 2, 3, 4, 5]),
|
||||
);
|
||||
// Force L0 non-indexed: strip the indices.
|
||||
let l0_nonidx = Arc::new(Geometry::new(l0.positions.clone()));
|
||||
let l1 = Arc::new(l0.decimated(1)); // welded + indexed
|
||||
let (vertices, indices, rows) =
|
||||
pack_levels(&[Arc::clone(&l0_nonidx), Arc::clone(&l1)]).expect("small mesh packs");
|
||||
assert_eq!(vertices.len(), 6 + l1.positions.len());
|
||||
let packed = indices.expect("an indexed level exists");
|
||||
// L0 contributes no indices; L1's are rebased by 6.
|
||||
assert_eq!(packed.len(), l1.indices().unwrap().len());
|
||||
assert_eq!(rows[0].index_count, 0, "non-indexed L0 row");
|
||||
assert_eq!(rows[1].index_offset, 0, "L1 is the first indexed level");
|
||||
assert_eq!(rows[1].vertex_offset, 6, "element units, not bytes");
|
||||
for (j, idx) in l1.indices().unwrap().iter().enumerate() {
|
||||
assert_eq!(packed[j], *idx + 6);
|
||||
}
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn pack_levels_all_non_indexed() {
|
||||
let l0 = Arc::new(Geometry::new(vec![
|
||||
[0.0, 0.0, 0.0],
|
||||
[1.0, 0.0, 0.0],
|
||||
[0.0, 1.0, 0.0],
|
||||
]));
|
||||
let (vertices, indices, rows) = pack_levels(&[l0.clone()]).expect("small mesh packs");
|
||||
assert_eq!(vertices.len(), 3);
|
||||
assert!(indices.is_none());
|
||||
assert_eq!(rows[0].index_count, 0);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn lod_table_from_rows() {
|
||||
let l0 = Arc::new(primitives::icosphere(1.0, 0));
|
||||
let l1 = Arc::new(l0.decimated(10));
|
||||
let (_, _, rows) = pack_levels(&[l0, l1]).expect("small mesh packs");
|
||||
let table = crate::resources::uniform::LodTable::from_rows(&rows);
|
||||
assert_eq!(table.count, 2);
|
||||
assert_eq!(table.rows[0].index_count, 60);
|
||||
assert_eq!(table.rows[1].vertex_count, rows[1].vertex_count);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn pack_levels_rejects_u16_overflow() {
|
||||
// A level with 70k vertices (indexed) cannot be packed with u16 rebased indices.
|
||||
let big = Arc::new(
|
||||
Geometry::new(vec![[0.0, 0.0, 0.0]; 70_000]).with_indices(vec![0u16; 70_000 / 3 * 3]),
|
||||
);
|
||||
let err = pack_levels(&[big]).unwrap_err();
|
||||
assert_eq!(err.total_vertices, 70_000);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -24,13 +24,13 @@ pub mod vertex;
|
||||
pub use camera::{Camera, CameraController, PITCH_LIMIT};
|
||||
pub use lights::Lights;
|
||||
pub use material::Material;
|
||||
pub use mesh::Mesh;
|
||||
pub use mesh::{LodMode, Mesh, PackError};
|
||||
pub use texture::{Texture, TextureError};
|
||||
pub use uniform::{
|
||||
BBOX_SLOT_SIZE, BBoxSlot, CULL_UNIFORMS_SIZE, CullUniforms, DRAW_SLOT_SIZE, DrawSlot,
|
||||
FRAME_UNIFORMS_SIZE, FrameUniforms, Light, LightType, MAT_SLOT_SIZE, MAX_LIGHTS, MatSlot,
|
||||
OBJECT_UNIFORM_SIZE, ObjectUniform, SHADOW_UNIFORM_SIZE, ShadowUniform, TRANSFORM_SLOT_SIZE,
|
||||
TransformSlot,
|
||||
FRAME_UNIFORMS_SIZE, FrameUniforms, LOD_ROW_SIZE, LOD_TABLE_SIZE, Light, LightType, LodRow,
|
||||
LodTable, MAT_SLOT_SIZE, MAX_LIGHTS, MatSlot, OBJECT_UNIFORM_SIZE, ObjectUniform,
|
||||
SHADOW_UNIFORM_SIZE, ShadowUniform, TRANSFORM_SLOT_SIZE, TransformSlot,
|
||||
};
|
||||
pub use vertex::Vertex;
|
||||
|
||||
|
||||
@@ -383,6 +383,79 @@ impl CullUniforms {
|
||||
}
|
||||
}
|
||||
|
||||
/// Size in bytes of one [`LodRow`] (16 B = 4 u32), matching WGSL `LodRow`.
|
||||
pub const LOD_ROW_SIZE: u64 = 16;
|
||||
|
||||
/// Size in bytes of one [`LodTable`] (80 B = 5 × 16 B), matching WGSL `LodTable`.
|
||||
pub const LOD_TABLE_SIZE: u64 = 80;
|
||||
|
||||
/// One LOD level of a mesh's **packed** vertex/index buffers (Step 19, D7) — the per-level
|
||||
/// draw offsets the GPU cull pass needs to emit the level's indirect draw args.
|
||||
/// Mirrors the WGSL `LodRow` (16 bytes: `vec4<u32>`).
|
||||
#[repr(C)]
|
||||
#[derive(Debug, Clone, Copy, PartialEq, Pod, Zeroable, Default)]
|
||||
pub struct LodRow {
|
||||
/// **Element** offset of this level's vertices within the mesh's packed vertex buffer
|
||||
/// (a vertex index, not a byte offset — the WebGPU `drawIndirectNonIndexed` first_vertex
|
||||
/// is a vertex index; it goes straight into the indirect args' `first_vertex`).
|
||||
pub vertex_offset: u32,
|
||||
/// Number of vertices of this level.
|
||||
pub vertex_count: u32,
|
||||
/// **Element** offset of this level's indices within the mesh's packed index buffer
|
||||
/// (an index element, not a byte offset — the `drawIndirectIndexed` first_index is an
|
||||
/// index element; 0 when the level is non-indexed or no earlier level has indices).
|
||||
pub index_offset: u32,
|
||||
/// Number of indices of this level (0 when non-indexed); the draw count.
|
||||
pub index_count: u32,
|
||||
}
|
||||
|
||||
impl LodRow {
|
||||
/// Builds a row from the packed-buffer offsets the packer computed.
|
||||
/// `index_offset`/`index_count` are 0 for a non-indexed level.
|
||||
pub fn new(vertex_offset: u32, vertex_count: u32, index_offset: u32, index_count: u32) -> Self {
|
||||
Self {
|
||||
vertex_offset,
|
||||
vertex_count,
|
||||
index_offset,
|
||||
index_count,
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
/// Per-mesh LOD table (Step 19, D7): one [`LodRow`] per level, uploaded once per mesh and
|
||||
/// read by the GPU cull pass to map a CPU-decided level to indirect draw args. Level 0
|
||||
/// **always** exists and is byte-exact with the full mesh. Mirrors the WGSL `LodTable`
|
||||
/// (80 bytes = count @0 + 4 × 16-byte rows @16..80, each row a `vec4<u32>`).
|
||||
#[repr(C)]
|
||||
#[derive(Debug, Clone, Copy, PartialEq, Pod, Zeroable, Default)]
|
||||
pub struct LodTable {
|
||||
/// Number of valid levels (1 = no LOD, single level).
|
||||
pub count: u32,
|
||||
/// Padding so the rows start at byte 16 (WGSL `LodTable`: `count : u32` @0, a 12-byte pad,
|
||||
/// then the `array<LodRow, 4>` @16 — the WGSL pad is an `array<u32, 3>` (alignment 4, NOT
|
||||
/// `vec3<u32>` which would align to 16 and grow the struct to 96 bytes).
|
||||
pub _pad: [u32; 3],
|
||||
/// One 16-byte row per level (`LOD_ROW_SIZE`-spaced), zeroed beyond `count`.
|
||||
pub rows: [LodRow; crate::utils::conf::MAX_LOD_LEVELS as usize],
|
||||
}
|
||||
|
||||
impl LodTable {
|
||||
/// Builds a table from up to `MAX_LOD_LEVELS` rows (row 0 = level 0 = full mesh).
|
||||
pub fn from_rows(rows: &[LodRow]) -> Self {
|
||||
let mut table = Self::default();
|
||||
table._pad = [0; 3];
|
||||
table.count = rows.len() as u32;
|
||||
for (i, row) in rows
|
||||
.iter()
|
||||
.enumerate()
|
||||
.take(crate::utils::conf::MAX_LOD_LEVELS as usize)
|
||||
{
|
||||
table.rows[i] = *row;
|
||||
}
|
||||
table
|
||||
}
|
||||
}
|
||||
|
||||
#[cfg(test)]
|
||||
mod tests {
|
||||
use super::*;
|
||||
@@ -479,4 +552,39 @@ mod tests {
|
||||
assert_eq!(offset_of!(CullUniforms, num_slots), 96);
|
||||
assert_eq!(offset_of!(CullUniforms, culling), 100);
|
||||
}
|
||||
|
||||
#[test]
|
||||
fn lod_layouts_match_wgsl() {
|
||||
// LodRow: four u32 -> 16 (a vec4<u32> in WGSL).
|
||||
assert_eq!(size_of::<LodRow>(), 16);
|
||||
assert_eq!(LOD_ROW_SIZE, 16);
|
||||
assert_eq!(offset_of!(LodRow, vertex_offset), 0);
|
||||
assert_eq!(offset_of!(LodRow, vertex_count), 4);
|
||||
assert_eq!(offset_of!(LodRow, index_offset), 8);
|
||||
assert_eq!(offset_of!(LodRow, index_count), 12);
|
||||
|
||||
// LodTable: count @0 + 4 rows @16 -> 80.
|
||||
assert_eq!(size_of::<LodTable>(), 80);
|
||||
assert_eq!(LOD_TABLE_SIZE, 80);
|
||||
assert_eq!(offset_of!(LodTable, count), 0);
|
||||
assert_eq!(offset_of!(LodTable, rows), 16);
|
||||
|
||||
// from_rows: count + rows filled, rest zeroed.
|
||||
let table = LodTable::from_rows(&[LodRow {
|
||||
vertex_offset: 0,
|
||||
vertex_count: 100,
|
||||
index_offset: 0,
|
||||
index_count: 300,
|
||||
}]);
|
||||
assert_eq!(table.count, 1);
|
||||
assert_eq!(table.rows[0].vertex_count, 100);
|
||||
assert_eq!(table.rows[1], LodRow::default());
|
||||
|
||||
// MAX_LOD_LEVELS rows fit exactly.
|
||||
let table = LodTable::from_rows(&vec![
|
||||
LodRow::default();
|
||||
crate::utils::conf::MAX_LOD_LEVELS as usize
|
||||
]);
|
||||
assert_eq!(table.count, crate::utils::conf::MAX_LOD_LEVELS);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -277,6 +277,148 @@ impl Scene {
|
||||
Ok(id.to_string())
|
||||
}
|
||||
|
||||
/// Builds, (optionally) links to a Material, and registers a **multi-level** Mesh in one
|
||||
/// declarative call (Step 19, D6/D7). Level 0 is `geometry` (byte-exact); levels 1.. are
|
||||
/// auto-generated by greedy decimation (`Geometry::generate_lod_levels`, D10) at halving
|
||||
/// targets. All levels are packed into the mesh's single vertex/index buffers (D7), and the
|
||||
/// per-level offsets are uploaded per frame as the mesh's LOD table for the GPU cull pass.
|
||||
///
|
||||
/// Inputs: id (unique mesh id), geometry (level 0 — the full mesh), material (optional
|
||||
/// material id, same rule as [`create_mesh`]), levels (2..=MAX_LOD_LEVELS).
|
||||
/// Returns `Err` if the id exists, the material is unknown, `levels` is out of range, or the
|
||||
/// packed vertex total exceeds the u16 index limit (65535).
|
||||
pub fn create_mesh_with_lod(
|
||||
&mut self,
|
||||
id: &str,
|
||||
geometry: Geometry,
|
||||
material: Option<&str>,
|
||||
levels: u8,
|
||||
) -> Result<String, String> {
|
||||
use crate::utils::conf::MAX_LOD_LEVELS;
|
||||
if self.meshes.contains_key(id) {
|
||||
return Err(format!("Mesh ID '{}' already exists.", id));
|
||||
}
|
||||
if levels < 2 || u32::from(levels) > MAX_LOD_LEVELS {
|
||||
return Err(format!(
|
||||
"LOD levels must be in 2..={MAX_LOD_LEVELS} (got {levels})."
|
||||
));
|
||||
}
|
||||
let lod_levels: Vec<Arc<Geometry>> = geometry
|
||||
.generate_lod_levels(levels)
|
||||
.into_iter()
|
||||
.map(Arc::new)
|
||||
.collect();
|
||||
let total: u32 = lod_levels.iter().map(|l| l.positions.len() as u32).sum();
|
||||
if total >= 65536 {
|
||||
return Err(format!(
|
||||
"Packed LOD vertex total ({total}) exceeds the 65535 u16 index limit; use fewer levels or a smaller mesh."
|
||||
));
|
||||
}
|
||||
let mut mesh = Mesh::from_geometry_lod(
|
||||
self.device(),
|
||||
lod_levels,
|
||||
None,
|
||||
crate::resources::LodMode::Auto,
|
||||
);
|
||||
if let Some(name) = material {
|
||||
let mat = self
|
||||
.materials
|
||||
.get(name)
|
||||
.ok_or_else(|| format!("Material '{}' does not exist.", name))?
|
||||
.clone();
|
||||
mesh.set_material(mat);
|
||||
}
|
||||
self.meshes.insert(id.to_string(), Arc::new(mesh));
|
||||
self.mesh_order.push(id.to_string());
|
||||
Ok(id.to_string())
|
||||
}
|
||||
|
||||
/// Adds (or replaces) an **explicitly provided** LOD level on an existing mesh (Step 19, D6).
|
||||
/// The level is packed into the mesh's buffers alongside the others (D7) and the per-mesh LOD
|
||||
/// table is updated (it is re-uploaded every frame, so the change takes effect next frame).
|
||||
///
|
||||
/// Inputs: id (existing mesh id), level (index ≥ 1; must be ≤ the current level count —
|
||||
/// append at the end or replace in place), geometry (the level's geometry).
|
||||
/// Validation: the level must validate, carry the **same attribute set and indexedness** as
|
||||
/// level 0, keep the packed vertex total under 65536, and stay within `MAX_LOD_LEVELS`.
|
||||
pub fn add_mesh_lod(&mut self, id: &str, level: u8, geometry: Geometry) -> Result<(), String> {
|
||||
use crate::utils::conf::MAX_LOD_LEVELS;
|
||||
let current = self
|
||||
.meshes
|
||||
.get(id)
|
||||
.ok_or_else(|| format!("Mesh '{}' does not exist.", id))?;
|
||||
let levels_count = current.num_lod_levels();
|
||||
if level < 1 || u32::from(level) > MAX_LOD_LEVELS {
|
||||
return Err(format!(
|
||||
"LOD level must be in 1..={MAX_LOD_LEVELS} (got {level})."
|
||||
));
|
||||
}
|
||||
if level as usize > levels_count {
|
||||
return Err(format!(
|
||||
"Mesh '{}' has {levels_count} level(s); level {level} does not exist and the next free level is {levels_count}.",
|
||||
id
|
||||
));
|
||||
}
|
||||
let l0 = current.geometry();
|
||||
let attr = |g: &Geometry| (g.normals.is_some(), g.uvs.is_some(), g.colors.is_some());
|
||||
if attr(&geometry) != attr(l0) {
|
||||
return Err(format!(
|
||||
"LOD level {level} of mesh '{}' must have the same attribute set (normals/UVs/colors) as level 0.",
|
||||
id
|
||||
));
|
||||
}
|
||||
if geometry.indices().is_some() != l0.indices().is_some() {
|
||||
return Err(format!(
|
||||
"LOD level {level} of mesh '{}' must have the same indexedness as level 0.",
|
||||
id
|
||||
));
|
||||
}
|
||||
geometry
|
||||
.validate()
|
||||
.map_err(|e| format!("LOD level {level} of mesh '{}': {e}", id))?;
|
||||
|
||||
let mut new_levels: Vec<Arc<Geometry>> = (0..levels_count)
|
||||
.map(|i| current.lod_levels_arc(i))
|
||||
.collect();
|
||||
if level as usize == levels_count {
|
||||
if levels_count >= MAX_LOD_LEVELS as usize {
|
||||
return Err(format!(
|
||||
"Mesh '{}' already has the maximum of {MAX_LOD_LEVELS} LOD levels.",
|
||||
id
|
||||
));
|
||||
}
|
||||
new_levels.push(Arc::new(geometry)); // append the next level (L_k at vec index k)
|
||||
} else {
|
||||
new_levels[level as usize] = Arc::new(geometry); // replace L_k in place (vec index k)
|
||||
}
|
||||
let total: u32 = new_levels.iter().map(|l| l.positions.len() as u32).sum();
|
||||
if total >= 65536 {
|
||||
return Err(format!(
|
||||
"Packed LOD vertex total ({total}) exceeds the 65535 u16 index limit."
|
||||
));
|
||||
}
|
||||
|
||||
let material = current.material().cloned();
|
||||
let mesh = Mesh::from_geometry_lod(
|
||||
self.device(),
|
||||
new_levels,
|
||||
material,
|
||||
crate::resources::LodMode::Explicit,
|
||||
);
|
||||
// Entities reference the mesh by (stable) index, not by Arc — swapping the Arc is safe.
|
||||
self.meshes.insert(id.to_string(), Arc::new(mesh));
|
||||
Ok(())
|
||||
}
|
||||
|
||||
/// The per-mesh LOD tables in `mesh_order` order (one 80-byte table per mesh, level 0 first)
|
||||
/// — the payload of the GPU `lod_tables` buffer, uploaded every frame (Step 19, D7).
|
||||
pub fn mesh_lod_tables(&self) -> Vec<crate::resources::LodTable> {
|
||||
self.mesh_order
|
||||
.iter()
|
||||
.map(|name| self.meshes[name].lod_table())
|
||||
.collect()
|
||||
}
|
||||
|
||||
/// Returns the Scene's default material: the `standard` shader pipeline, built lazily on first
|
||||
/// call and cached afterwards. Used by `Renderer::render_scene` for meshes that carry no material.
|
||||
/// Note: the flat (unlit) look is *not* a property of this material — it is driven by the
|
||||
|
||||
@@ -9,11 +9,17 @@
|
||||
// = no-op) instead of a CPU-side per-entity loop.
|
||||
//
|
||||
// All buffers are fixed-capacity (MAX_ENTITIES = 256, see ARCHI_CPU_GPU.md D12) and are
|
||||
// allocated once. Each frame the CPU rewrites the transform slots and cull uniforms;
|
||||
// everything else is GPU-driven.
|
||||
// allocated once. Each frame the CPU rewrites the transform slots, cull uniforms, LOD levels
|
||||
// and LOD tables; everything else is GPU-driven.
|
||||
//
|
||||
// GPU buffer layouts mirror the bytemuck structs in `resources::uniform` (byte-for-byte):
|
||||
// TransformSlot (64B), MatSlot (256B), BBoxSlot (32B), DrawSlot (80B), CullUniforms (112B).
|
||||
// TransformSlot (64B), MatSlot (256B), BBoxSlot (32B), DrawSlot (80B), CullUniforms (112B),
|
||||
// LodRow (16B), LodTable (80B).
|
||||
//
|
||||
// Step 19 (LOD): the CPU decides each entity's level (screen-space size + hysteresis, D8);
|
||||
// the cull pass maps it to the packed level's draw args through the per-mesh LOD table
|
||||
// (binding 4) and per-slot level array (binding 3). With LOD disabled the CPU writes level 0
|
||||
// everywhere and the args are byte-identical to the pre-LOD behavior.
|
||||
//
|
||||
// GOTCHA — WGSL `select` argument order: `select(reject, accept, cond)` returns the SECOND
|
||||
// argument when `cond` is true and the FIRST when false (the reverse of HLSL's
|
||||
@@ -63,6 +69,22 @@ struct CullUniforms {
|
||||
_pad : vec2u,
|
||||
};
|
||||
|
||||
// 16 bytes: one LOD level's draw offsets (Step 19, D7). ELEMENT units, not bytes: x is a
|
||||
// vertex index (drawIndirectNonIndexed first_vertex) and z an index element (drawIndirectIndexed
|
||||
// first_index) — the packed vertex/index buffers are bound in full (offset 0), only these
|
||||
// first_* values move per level. w = 0 marks a non-indexed level.
|
||||
struct LodRow {
|
||||
o : vec4u, // x = first_vertex, y = vertex_count, z = first_index, w = index_count
|
||||
};
|
||||
|
||||
// 80 bytes: the per-mesh LOD table (Step 19, D7). count @0, 12-byte pad (an array<u32,3> —
|
||||
// NOT vec3u, which would align to 16 and grow the struct to 96 B), rows @16..80.
|
||||
struct LodTable {
|
||||
count : u32, // number of valid levels (1 = no LOD)
|
||||
_pad : array<u32, 3>,
|
||||
rows : array<LodRow, 4>, // MAX_LOD_LEVELS = 4; zeroed beyond count
|
||||
};
|
||||
|
||||
// ---- Bind groups (Step 15.5) ----
|
||||
// Group 0 (transforms) is shared by both entry points; group 1 (matrices) by `compute_matrices`;
|
||||
// group 2 (cull uniforms + bboxes + draw args) by `cull`. Each pipeline infers the subset it uses.
|
||||
@@ -71,6 +93,10 @@ struct CullUniforms {
|
||||
@group(2) @binding(0) var<uniform> cull_u : CullUniforms;
|
||||
@group(2) @binding(1) var<storage, read> bboxes : array<BBoxSlot>;
|
||||
@group(2) @binding(2) var<storage, read_write> draw_args : array<DrawSlot>;
|
||||
// Step 19: per-slot CPU-decided LOD level (one u32 per entity slot) and the per-mesh LOD tables
|
||||
// (one 80-byte LodTable per mesh, mesh_order order — same indexing as `bboxes`).
|
||||
@group(2) @binding(3) var<storage, read> lod_levels : array<u32>;
|
||||
@group(2) @binding(4) var<storage, read> lod_tables : array<LodTable>;
|
||||
|
||||
// ---- Shared helpers ----
|
||||
|
||||
@@ -119,10 +145,23 @@ fn identity_mat() -> mat4x4f {
|
||||
|
||||
// Fills a draw slot with a non-zero (visible) count of `count`, or zero (culled / inactive).
|
||||
// The count lands in `.a.x`; `.a.y` (instance count) is the constant 1; the rest stays zero.
|
||||
// (A level-0 draw produced by `write_level_args` is byte-identical to this: first_* = 0.)
|
||||
fn set_draw_count(i : u32, count : u32) {
|
||||
draw_args[i].a = vec4u(count, 1u, 0u, 0u);
|
||||
}
|
||||
|
||||
// Fills the slot's indirect args with the draw command of LOD level `row` (Step 19, D7):
|
||||
// indexed levels use the 5-field layout (index_count, instances, first_index, base_vertex,
|
||||
// base_instance) and non-indexed levels the 4-field layout (vertex_count, instances,
|
||||
// first_vertex, base_instance). The element-unit offsets in the row become the first_* fields.
|
||||
fn write_level_args(i : u32, row : LodRow) {
|
||||
if (row.o.w > 0u) {
|
||||
draw_args[i].a = vec4u(row.o.w, 1u, row.o.z, 0u);
|
||||
} else {
|
||||
draw_args[i].a = vec4u(row.o.y, 1u, row.o.x, 0u);
|
||||
}
|
||||
}
|
||||
|
||||
// ---- Pass 1: derive world matrices (Step 15.5) ----
|
||||
// Dispatched for MAX_ENTITIES; inactive slots get the identity matrix (a harmless stale read).
|
||||
@compute
|
||||
@@ -137,9 +176,11 @@ fn compute_matrices(@builtin(global_invocation_id) gid : vec3u) {
|
||||
}
|
||||
}
|
||||
|
||||
// ---- Pass 2: cull + fill indirect draw args (Step 15.6) ----
|
||||
// ---- Pass 2: LOD level → draw args + cull (Steps 15.6 + 19) ----
|
||||
// Dispatched for MAX_ENTITIES. Slots at or beyond `num_slots` (and inactive slots) are zeroed so
|
||||
// the indirect render passes skip them; visible slots keep their packed count.
|
||||
// the indirect render passes skip them. For visible slots the CPU-decided LOD level (binding 3)
|
||||
// is mapped through the mesh's LOD table (binding 4) to the level's draw command; culling
|
||||
// (when enabled) zeroes it against the frustum planes.
|
||||
@compute
|
||||
@workgroup_size(64)
|
||||
fn cull(@builtin(global_invocation_id) gid : vec3u) {
|
||||
@@ -158,32 +199,41 @@ fn cull(@builtin(global_invocation_id) gid : vec3u) {
|
||||
return;
|
||||
}
|
||||
|
||||
// Culling disabled: every active entity is visible, with its packed count.
|
||||
if (cull_u.culling == 0u) {
|
||||
set_draw_count(i, u32(t.flags.z));
|
||||
return;
|
||||
// LOD (Step 19, D8): the CPU decides the level per slot (projected size + hysteresis);
|
||||
// the GPU only maps it to the packed level's draw args. Clamped to the mesh's level count
|
||||
// (the CPU clamps too — defense in depth against a stale level after an add_mesh_lod).
|
||||
let table = lod_tables[u32(t.flags.x)];
|
||||
var level = lod_levels[i];
|
||||
if (level >= table.count) {
|
||||
level = table.count - 1u;
|
||||
}
|
||||
let row = table.rows[level].o;
|
||||
|
||||
// Culling enabled: test the entity's world bounding sphere against the frustum planes.
|
||||
let b = bboxes[u32(t.flags.x)];
|
||||
let center_local = (b.min + b.max) * 0.5;
|
||||
// World center = translation + rotation * local center (no scale; the radius carries the scale).
|
||||
let center_world = t.translation + rotate_by_quat(center_local, t.rotation);
|
||||
let half_extents = (b.max - b.min) * 0.5;
|
||||
let radius = length(half_extents) * max(t.scale.x, max(t.scale.y, t.scale.z));
|
||||
|
||||
// Culling: test the entity's world bounding sphere against the frustum planes.
|
||||
var visible = true;
|
||||
for (var p = 0u; p < 6u; p = p + 1u) {
|
||||
let plane = cull_u.planes[p];
|
||||
let dist = dot(plane.xyz, center_world) + plane.w;
|
||||
if (dist < -radius) {
|
||||
visible = false;
|
||||
break;
|
||||
if (cull_u.culling == 1u) {
|
||||
let b = bboxes[u32(t.flags.x)];
|
||||
let center_local = (b.min + b.max) * 0.5;
|
||||
// World center = translation + rotation * local center (no scale; the radius carries it).
|
||||
let center_world = t.translation + rotate_by_quat(center_local, t.rotation);
|
||||
let half_extents = (b.max - b.min) * 0.5;
|
||||
let radius = length(half_extents) * max(t.scale.x, max(t.scale.y, t.scale.z));
|
||||
|
||||
for (var p = 0u; p < 6u; p = p + 1u) {
|
||||
let plane = cull_u.planes[p];
|
||||
let dist = dot(plane.xyz, center_world) + plane.w;
|
||||
if (dist < -radius) {
|
||||
visible = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// NOTE: WGSL `select(reject, accept, cond)` — the accept value is the SECOND argument.
|
||||
// visible => full count; culled => 0.
|
||||
let count = select(0u, u32(t.flags.z), visible);
|
||||
set_draw_count(i, count);
|
||||
if (!visible) {
|
||||
set_draw_count(i, 0u);
|
||||
return;
|
||||
}
|
||||
|
||||
// Visible: the chosen level's draw command (level 0 is byte-identical to the pre-LOD args).
|
||||
write_level_args(i, LodRow(row));
|
||||
}
|
||||
|
||||
@@ -57,6 +57,17 @@ pub const GPU_DRIVEN_SHADER: &str = include_str!("../shaders/gpu_driven.wgsl");
|
||||
/// 256 is a multiple of the 64-wide workgroup size, giving a whole number of workgroups.
|
||||
pub const MAX_ENTITIES: u32 = 256;
|
||||
|
||||
/// Maximum number of LOD levels per mesh (Étape 19, D8). The per-mesh LOD table
|
||||
/// (`LodTable`, 80 bytes) carries one 16-byte row per level — 4 rows + the count header.
|
||||
pub const MAX_LOD_LEVELS: u32 = 4;
|
||||
|
||||
/// Default LOD thresholds in **pixels** of projected bounding-sphere radius (Étape 19, D4/D8):
|
||||
/// `thresholds[k]` is the radius *above which* level k+1 is required (descending). With these
|
||||
/// values: `r > 48` → L0, `12 < r ≤ 48` → L1, `r ≤ 12` → L2+ (clamped). Constants in v1
|
||||
/// (per-scene/mesh configurability is a follow-up); the hysteresis dead band (×0.8 to go
|
||||
/// coarser) lives in `math::lod::lod_level`.
|
||||
pub const LOD_THRESHOLDS: [f32; 2] = [48.0, 12.0];
|
||||
|
||||
/// Workgroup size of the GPU-driven compute shaders (matches the `@workgroup_size` in
|
||||
/// `gpu_driven.wgsl`). The compute dispatch is `MAX_ENTITIES / WORKGROUP_SIZE` workgroups.
|
||||
pub const GPU_WORKGROUP_SIZE: u32 = 64;
|
||||
|
||||
Reference in New Issue
Block a user