# Architecture du Système de Particules (GPU) > Document de référence pour l'implémentation. > Statut : **DESIGN FINAL** — base pour les DRAFT d'implémentation. > Principe directeur : **Pool ≠ Driver**. La ressource GPU (pool) est séparée du mécanisme de simulation (driver). --- ## 1. Vue d'ensemble ``` ┌─────────────────────────────────────────────────────────────────────────┐ │ SCÈNE │ │ │ │ ┌───────────────────────────────────────────────────────────────────┐ │ │ │ PARTICLE POOL (ressource GPU, créée une fois) │ │ │ │ │ │ │ │ • Buffer storage (N × 80 bytes [D15/D19]) │ │ │ │ • Pipeline render (billboard instancé) │ │ │ │ • Texture + Sampler + Blending mode │ │ │ │ • Draw call (instanced, 1 par frame) │ │ │ │ │ │ │ │ Rôle : STOCKER + AFFICHER │ │ │ ├───────────────────────────────────────────────────────────────────┤ │ │ │ DRIVER (mécanisme de mise à jour, attaché au pool) │ │ │ │ │ │ │ │ ┌───────────────┐ ┌───────────────┐ ┌───────────────────┐ │ │ │ │ │ GPU Emitter │ │ CPU Emitter │ │ Manual │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ Compute pass │ │ write_buffer │ │ write_range │ │ │ │ │ │ (simulation │ │ (full upload │ │ (user writes │ │ │ │ │ │ en shader) │ │ par frame) │ │ directement) │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ │ Rôle : │ │ Rôle : │ │ Rôle : │ │ │ │ │ │ SIMULER (GPU) │ │ SIMULER (CPU) │ │ CONTRÔLE TOTAL │ │ │ │ │ └───────────────┘ └───────────────┘ └───────────────────┘ │ │ │ │ │ │ │ │ Un pool a UN SEUL driver actif à la fois. │ │ │ │ Changer de driver = swap (pas de reallocation GPU). │ │ │ └───────────────────────────────────────────────────────────────────┘ │ │ │ └─────────────────────────────────────────────────────────────────────────┘ ``` ### Séparation des responsabilités | Couche | Sait | Ne sait pas | |--------|------|-------------| | **Pool** | Comment afficher (pipeline, texture, blending, draw) | Comment les particules bougent | | **Driver** | Comment mettre à jour le buffer (compute / upload / écriture directe) | Comment afficher | | **User** | Quelle intention (feu, pluie, explosion, custom) | Les détails wgpu | --- ## 2. État par particule (80 bytes : `uv_rect` [D15], sans padding [D19]) ```rust /// Miroir du struct WGSL `Particle` (§8.1/§8.2 — un seul struct partagé compute + render, [D19]). /// 80 bytes avec `uv_rect` [D15] : espace **storage** → vec3/vec4 align 4 → **pas de padding** [D19]. #[repr(C)] #[derive(Copy, Clone, Pod, Zeroable)] pub struct Particle { pub pos: [f32; 3], // offset 0 — position monde (xyz) pub vel: [f32; 3], // offset 12 — vélocité (xyz) pub life: f32, // offset 24 — vie restante (seconds) pub max_life: f32, // offset 28 — vie initiale (pour fade normalisé) pub size: f32, // offset 32 — taille courante (world units) pub size_growth: f32, // offset 36 — croissance par seconde (+ = grandir, - = rétrécir) pub angle: f32, // offset 40 — rotation 2D courante (radians) pub angular_vel: f32, // offset 44 — vitesse angulaire (rad/s) pub color: [f32; 4], // offset 48 — RGBA (alpha : `alpha_scale` [D18]) pub uv_rect: [f32; 4], // offset 64 — zone UV (ox, oy, sx, sy) [D15] // Total : 80 bytes (layout identique Rust = WGSL storage, align 4) [D19] } ``` ### Transform par particule Chaque particule porte un **transform complet** : | Composante | Stockage | Mise à jour (GPU compute) | Effet visuel | |-----------|----------|---------------------------|--------------| | Translation 3D | `pos` | `pos += vel * dt` | Déplacement | | Trajectoire | `vel` | `vel += gravity * dt; vel *= (1 - drag * dt)` | Arc, chute, flottement | | Scale 2D | `size` + `size_growth` | `size += size_growth * dt` | Croissance (fumée) / rétrécissement (feu) | | Rotation 2D | `angle` + `angular_vel` | `angle += angular_vel * dt` | Spin (turbulence, pétale) | | Opacité | `color.a` | `color.a = life / max_life` | Fade out progressif | | Zone UV | `uv_rect` | fixe (set au spawn) [D15] | Multi-motifs depuis une texture (atlas) | > Le RGB de `color` est fixé au spawn (CPU). Seule l'alpha est modulée par le driver. > **V2** : color ramp (`mix(color_start, color_end, 1.0 - life/max_life)`). ### Taille du pool | `max_count` | Buffer size | Upload CPU (si driver CPU) @60fps | Verdict | |-------------|-------------|-----------------------------------|---------| | 1 024 | 64 KB | 3.8 MB/s | Trivial | | 10 000 | 640 KB | 38 MB/s | OK | | 50 000 | 3.2 MB | 190 MB/s | Correct (PC) | | 65 536 | 4 MB | 245 MB/s | Max recommandé pour CPU | | 100 000 | 6.4 MB | 380 MB/s | → Préférer driver GPU | > [D15] Avec `uv_rect` (80 B) : tous les chiffres × 1.25 (10K → 800 KB ≈ 48 MB/s, 100K → 8 MB ≈ 480 MB/s). Verdicts inchangés. --- ## 3. Le Pool (ressource GPU) ### 3.1 Création Le pool est créé **une fois** au setup. Il alloue : - Le buffer storage (N × 80 bytes [D15]) - Le buffer d'index compact + args indirect ([D17], cf §18) - Le pipeline render (billboard instancé) - Le bind group layout - La texture + sampler (si fournie) - Le blending state ```rust pub struct ParticlePoolConfig { /// Nombre maximum de particules dans le pool. pub max_count: u32, /// Texture de la particule (ID dans scene.textures). /// `None` = texture disque doux intégrée (16×16, alpha radiale). pub texture: Option, /// Mode de blending. pub blending: BlendingMode, } pub enum BlendingMode { /// `src × 1 + dst × 1` — feu, sparkles, magie. Additive, /// `src × src.a + dst × (1 - src.a)` — fumée, neige, vapeur. Alpha, } ``` ### 3.2 Composantes internes du pool ```rust pub struct ParticlePool { /// Buffer storage (N × 80 B [D15/D19]). COPY_DST | STORAGE. pub(crate) buffer: wgpu::Buffer, /// Pipeline render (billboard instancé). pub(crate) pipeline: wgpu::RenderPipeline, /// Pipeline layout du render. pub(crate) layout: wgpu::BindGroupLayout, /// Bind group (particles + compact_index + camera_params + texture + sampler — tout est /// possédé par le pool, [D17/D19]). pub(crate) bind_group: wgpu::BindGroup, /// Sampler (Linear, ClampToEdge). pub(crate) sampler: wgpu::Sampler, /// Index compact : 1 u32 par slot [D17/D19], lu en VS par storage (binding 2). /// STORAGE | COPY_DST. pub(crate) compact_index: wgpu::Buffer, /// Args indirect draw (16 B : 4 × u32, `drawIndirect`) [D17/D19]. STORAGE | COPY_DST. pub(crate) indirect_args: wgpu::Buffer, /// Camera params (128 B : view + proj), possédée par le pool, écrite par le renderer chaque frame. pub(crate) camera_params: wgpu::Buffer, /// Taille max du pool. pub max_count: u32, /// Driver actuellement attaché (None = pool inactif). pub(crate) driver: Option>, /// Mode de blending (déterminé au build du pipeline). pub blending: BlendingMode, } ``` ### 3.3 Comportement | État du pool | Compute | Upload CPU | Draw | |-------------|---------|------------|------| | Pas de driver | ❌ | ❌ | ❌ (count = 0) | | Driver GPU actif | ✅ (1 dispatch) | Spawns seulement | ✅ | | Driver CPU actif | ❌ | Full range alive | ✅ | | Driver Manual actif | ❌ | Ce que l'utilisateur envoie | ✅ | | Driver actif mais `paused = true` | ❌ | ❌ | ❌ (count = 0) | **Coût zéro** : un pool sans driver (ou pausé) n'apparaît pas dans le render pass. --- ## 4. Les Drivers ### 4.1 Trait commun ```rust /// Trait implémenté par chaque type de driver. /// Le pool appelle ces méthodes dans l'ordre à chaque frame. pub trait ParticleDriver: Send { /// Appelé AVANT le compute (si driver GPU) ou avant le draw. /// Le driver peut écrire dans le pool buffer (spawns, updates). fn pre_compute(&mut self, queue: &wgpu::Queue, pool: &mut ParticlePool, dt: f32); /// Appelé APRÈS le compute (driver GPU uniquement). /// Permet de mettre à jour des uniforms post-simulation. fn post_compute(&mut self, queue: &wgpu::Queue, pool: &mut ParticlePool); /// Le driver veut-il qu'on fasse un compute dispatch cette frame ? fn needs_compute(&self) -> bool; /// Le driver veut qu'on drawe le pool ? fn needs_draw(&self) -> bool; } ``` ### 4.2 Driver GPU (GpuEmitter) **Rôle** : la simulation est faite par un compute shader. Le CPU ne fait que les spawns. ``` Frame loop (driver GPU) : 1. pre_compute : a. Accumulateur : acc += rate * dt b. n_spawn = floor(acc); acc -= n_spawn c. Trouver n_spawn slots morts dans le pool d. Pour chaque slot : tirer pos/vel/life/size/angle/color (random CPU) e. queue.write_buffer(&pool.buffer, offset_morts, new_particles) f. [D17] rien ici — les args indirect sortent de la compaction (fused dans le compute, §18) 2. Compute dispatch (le pool fait le dispatch) : - workgroups = ceil(max_count / 64) - Le shader intègre TOUS les slots alive - Le shader écrit `compact_index` + `indirect_args` (compaction fused, [D17]) 3. post_compute : (optionnel, ex: sync alive_count) 4. Draw (si alive_count > 0) ``` **Config** : ```rust pub struct GpuEmitterConfig { /// Taux d'émission (particules / seconde). pub rate: f32, /// Forme de spawn (où naît la particule). pub shape: SpawnShape, /// Pattern de vitesse (dans quelle direction). pub velocity: VelocityPattern, /// Plage de vitesse (pour Fixed/Cone/Sphere). Ignoré pour Box. pub speed_range: (f32, f32), /// Gravité (m/s²) appliquée par le compute. pub gravity: [f32; 3], /// Friction linéaire (0.0 = aucune, 1.0 = stop immédiat). pub drag: f32, /// Plage de vie (seconds). pub life_range: (f32, f32), /// Plage de taille initiale. pub size_range: (f32, f32), /// Croissance de taille par seconde (négatif = rétrécir). pub size_growth: f32, /// Plage de vitesse angulaire initiale (rad/s). pub angular_vel_range: (f32, f32), /// Couleur de base : [min, max] par canal **RGB** (le canal alpha est ignoré → `alpha_scale`). /// Chaque canal est tiré uniformément dans [min, max] au spawn [D18]. pub color_range: ([f32; 4], [f32; 4]), /// Palette de zones UV : chaque spawn tire un rect uniformément au hasard [D15/D18]. /// `uv_rects_count = 1` = toute la texture (comportement d'avant D15). pub uv_rects: [[f32; 4]; UV_RECTS_MAX], // UV_RECTS_MAX = 8 pub uv_rects_count: u8, // ≤ UV_RECTS_MAX /// Alpha cible (0..1, défaut 1.0) : l'intégration écrit `color.a = alpha_scale × life/max_life` [D18]. pub alpha_scale: f32, } ``` **Position de l'émetteur** : mise à jour par l'utilisateur à tout moment : ```rust scene.set_emitter_position(pool_id, [x, y, z]); ``` Le driver lit cette position à chaque spawn. ### 4.3 Driver CPU (CpuEmitter) **Rôle** : la simulation est faite en Rust (CPU). Le résultat est uploadé dans le pool à chaque frame. Pas de compute pass. ``` Frame loop (driver CPU) : 1. pre_compute (ici : simulation + upload) : a. Pour chaque slot alive : intégrer (pos += vel*dt, vel += grav*dt, life -= dt, ...) b. Tuer les particules mortes (life <= 0 → alive = 0) c. Accumulateur de spawn : acc += rate * dt d. Pour chaque nouveau spawn : écrire dans un slot mort e. queue.write_buffer(&pool.buffer, 0, &all_alive_particles) f. [D17] Construire `compact_index` + `indirect_args` au CPU (compaction exacte) 2. Pas de compute dispatch 3. post_compute : (rien) 4. Draw (si alive_count > 0) ``` **Config** : identique à `GpuEmitterConfig` (mêmes paramètres de simulation, y compris `color_range`/`uv_rects`/`alpha_scale` [D18] — les tirages se font en Rust au spawn). La différence est **où** l'intégration se fait. **Avantage du CPU** : possibilité d'interactions (collision avec objets de la scène, champs de force locaux, attracteurs). L'utilisateur peut implémenter une closure `fn custom_force(&self, particle: &Particle, scene: &Scene) -> Vec3`. ```rust pub struct CpuEmitterConfig { // ... même champs que GpuEmitterConfig ... /// Force custom (appelée par particule par frame). None = aucune. /// Permet : collision, attracteur, champ de vent, etc. pub custom_force: Option [f32; 3] + Send + Sync>>, } ``` **Limite** : bandwidth. Au-delà de ~10K particules, l'upload full buffer devient coûteux. Pour 50K+, préférer le driver GPU. ### 4.4 Driver Manual **Rôle** : l'utilisateur contrôle **directement** le contenu du buffer. Le pool ne fait que draw. ``` Frame loop (driver Manual) : 1. L'utilisateur appelle pool.upload(queue) quand il veut (ou le driver fait un upload automatique de la "dirty range") 2. Pas de compute dispatch 3. Draw (si alive_count > 0) ``` **API** : ```rust /// Handle exposé à l'utilisateur pour le driver manual. pub struct ManualPoolHandle<'a> { pool: &'a mut ParticlePool, queue: &'a wgpu::Queue, } impl ManualPoolHandle<'_> { /// Lire une particule (retourne une copie). pub fn get(&self, index: u32) -> Particle; /// Écrire une particule (marque le pool dirty). pub fn set(&mut self, index: u32, particle: Particle); /// Écrire un range de particules. pub fn write_range(&mut self, start: u32, count: u32, particles: &[Particle]); /// Définir manuellement le alive_count. pub fn set_count(&mut self, count: u32); /// Uploader le buffer dans le GPU (à appeler avant le draw). pub fn upload(&mut self); } ``` **Cas d'usage** : - Animation de particules le long d'une courbe (splines) - Transfert de particules entre pools - Import d'une animation de particules pré-calculée - Debug (visualiser l'état interne) - Effets très custom (morphing, shader-like en CPU) --- ## 5. Modèles d'émission ### 5.1 Shape de spawn ```rust pub enum SpawnShape { /// Position exacte de l'émetteur (point). Point, /// Boîte centrée sur l'émetteur. `extent` = demi-extents xyz. Box { extent: [f32; 3] }, /// Sphère centrée sur l'émetteur. Sphere { radius: f32 }, /// Plan perpendiculaire à `normal`, centré sur l'émetteur. /// `size` = demi-largeur x/y dans le plan local. Plane { normal: [f32; 3], size: [f32; 2] }, } ``` ### 5.2 Pattern de vitesse ```rust pub enum VelocityPattern { /// Vélocité fixe : `base_dir * speed`. (balles, rayons) Fixed, /// Cône : direction aléatoire dans un cône de demi-angle `spread_deg` autour de `base_dir`. Cone { spread_deg: f32 }, /// Sphère : direction aléatoire uniforme sur la sphère. (explosion) Sphere, /// Boîte : chaque composante tirée indépendamment dans [min_i, max_i]. (pluie) Box { min: [f32; 3], max: [f32; 3] }, } ``` ### 5.3 Direction base La direction base est implicitement dérivée : - Pour `Cone` / `Fixed` : l'utilisateur donne un `base_dir: [f32; 3]` dans le config - Pour `Sphere` : pas de direction base (isotrope) - Pour `Box` : pas de direction base (les min/max définissent tout) ```rust // Dans GpuEmitterConfig / CpuEmitterConfig : pub base_dir: [f32; 3], // utilisé par Fixed et Cone, ignoré par Sphere et Box ``` ### 5.4 Table des presets | Preset | Shape | Velocity | base_dir | speed | gravity | drag | life | size | growth | ang_vel | color_range (RGB min→max) [D18] | α_scale [D18] | blending | |--------|-------|----------|----------|-------|---------|------|------|------|--------|---------|---------------------------|---------|----------| | **Explosion** | Point | Sphere | — | 2-8 | (0,-9.8,0) | 0.1 | 0.5-1.5 | 0.05-0.15 | -0.05 | ±5 | [0.8,0.3,0]→[1,0.8,0.3] | 1.0 | Additive | | **Jet d'eau** | Point | Cone 10° | (0,1,0) | 3-5 | (0,-9.8,0) | 0.01 | 1.0-2.0 | 0.03-0.05 | 0 | 0 | [0.3,0.5,0.8]→[0.5,0.8,1] | 0.8 | Alpha | | **Pluie** | Plan (0,1,0) 20×20 | Box | — | vy:-5, vx/z:±0.3 | (0,0,0) | 0 | 3.0-5.0 | 0.01-0.02 | 0 | 0 | [0.7,0.8,1]→[1,1,1] | 0.3 | Alpha | | **Fumée** | Box [0.15,0,0.15] | Cone 45° | (0,1,0) | 0.3-1.0 | (0,0.3,0) | 0.3 | 2.0-4.0 | 0.2-0.4 | +0.3 | ±2 | [0.3,0.3,0.3]→[0.7,0.7,0.7] | 0.4 | Alpha | | **Feu** | Point | Cone 20° | (0,1,0) | 1-3 | (0,0.5,0) | 0.2 | 0.3-0.8 | 0.15-0.3 | -0.2 | ±8 | [1,0.3,0]→[1,0.9,0.2] | 1.0 | Additive | | **Neige** | Plan (0,1,0) 10×10 | Box | — | vy:-0.5, vx/z:±0.2 | (0,0,0) | 0 | 5-10 | 0.02-0.04 | 0 | ±3 | [0.9,0.9,1]→[1,1,1] | 0.8 | Alpha | | **Sparkles** | Sphere r=0.3 | Sphere | — | 0.1-0.5 | (0,0,0) | 0.5 | 1-3 | 0.02-0.05 | -0.01 | 0 | [1,0.8,0.4]→[1,1,1] | 1.0 | Additive | | **Débris** | Point | Sphere | — | 1-6 | (0,-9.8,0) | 0.02 | 1-3 | 0.03-0.08 | 0 | ±10 | [0.4,0.3,0.15]→[0.8,0.6,0.3] | 1.0 | Alpha | Par défaut `uv_rects = [(0,0,1,1)]` (toute la texture) [D15]. Exemple d'atlas : `Neige` avec 3 flocons → `uv_rects: [(0,0,⅓,1), (⅓,0,⅓,1), (⅔,0,⅓,1)]` — chaque spawn tire un flocon [D18]. Disponibles comme constructors : `GpuEmitterConfig::fire()`, `::smoke()`, `::rain()`, etc. --- ## 6. Buffers GPU | Buffer | Type | Taille | Écrit par | Lu par | |--------|------|--------|-----------|--------| | `particle_data` | Storage (RW) | N × 80 B [D15/D19] | Driver (spawn/update) | Compute + Render | | `compact_index` | Storage (RO) | N × u32 [D17/D19] | Compute (compaction) / Driver | Vertex shader (slot via `compact_index[ii]`) | | `indirect_args` | Storage (RO) | 16 B [D17/D19] | Compute (compaction) / Driver | `drawIndirect` | | `emitter_params` | Uniform (RO) | 64 B | Driver GPU (par frame) | Compute shader | | `camera_params` | Uniform (RO) | 128 B (view + proj) | Renderer (par frame) | Vertex shader | ### `emitter_params` (uniform, 64 bytes) ```wgsl struct EmitterParams { dt: f32, // delta time de la frame gravity: vec3, // accélération drag: f32, // friction alive_count: f32, // count courant (écrit par le compute via atomic) pool_size: f32, // taille max du pool alpha_scale: f32, // [D18] alpha cible (modulée par life/max_life dans le compute) _pad: vec2, // alignment } // Total : 48 B (espace uniform, vec3 align 16) [D18] ``` ### Bind groups **Compute pass (driver GPU uniquement)** : | Group | Binding | Type | Contenu | Visibility | |-------|---------|------|---------|------------| | 0 | 0 | Storage RW | `particle_data` | COMPUTE | | 0 | 1 | Uniform RO | `emitter_params` | COMPUTE | **Render pass (toujours)** : | Group | Binding | Type | Contenu | Visibility | |-------|---------|------|---------|------------| | 0 | 0 | Uniform RO | `camera_params` (view + proj, 128 B) | VERTEX | | 0 | 1 | Storage RO | `particle_data` | VERTEX | | 0 | 2 | Storage RO | `compact_index` [D17] | VERTEX | | 0 | 3 | Sampler | Sampler | FRAGMENT | | 0 | 4 | Texture | Texture particule | FRAGMENT | > **Note [D17/D19]** : le slot de l'instance arrive par **storage** (`compact_index[ii]`), pas par > attribut vertex (layout vide, pattern TM). Le count exact vient des args indirect — plus d'early-out. --- ## 7. Pipelines ### 7.1 Pipeline Compute (driver GPU) ```rust wgpu::ComputePipeline { layout: ComputePipelineLayout { bind_group_layouts: [bgl_compute], // storage RW + params }, // entry point : "cs_update" } ``` Dispatch : `workgroups = ceil(max_count / 64)`, un workgroup de 64 threads. ### 7.2 Pipeline Render (billboard) ```rust wgpu::RenderPipeline { layout: RenderPipelineLayout { bind_group_layouts: [bgl_render], // camera + particles + compact_index + sampler + texture (5 bindings) [D17] }, vertex: vs_main (billboard), fragment: fs_main (texture × color), primitive: TriangleList, vertex_buffer_layouts: [], // layout VIDE — quad via QUAD[vi], slot via storage binding 2 [D17/D19] multisample: sample_count du contexte, color_states: [blending mode du pool], depth_stencil: Some(LessEqual, ALWAYS), // depth test oui, depth write non (transparence) } ``` > **Vertex layout vide** (comme le fullscreen du tone mapping) : le quad est généré en shader > (`QUAD[vi]`), et le slot de l'instance arrive par **storage** (`compact_index[ii]`) — pas > d'attribut vertex [D17/D19]. ### 7.3 Blend states | Mode | color_ops | Formula | |------|-----------|---------| | **Additive** | src: One, dst: One | `output = src × 1 + dst × 1` | | **Alpha** | src: SrcAlpha, dst: OneMinusSrcAlpha | `output = src × a + dst × (1-a)` | Le blending est **figé au build du pipeline** (2 pipelines par pool si on veut les 2 modes — mais v1 : 1 mode par pool). ### 7.4 Depth - **Depth test** : `CompareFunction::LessEqual` (les particules derrière les objets opaques sont masquées) - **Depth write** : `STENCIL_WRITE_ONLY` → en pratique `depth_write_enabled: false` (les particules ne créent pas d'ombre depth entre elles) - **Stencil** : `Always` (pas de stencil) --- ## 8. Shaders (WGSL) ### 8.1 Compute shader (`particle_update.wgsl`) ```wgsl // particle_update.wgsl // Compute pass : intègre toutes les particules alive. struct Particle { // 80 B — espace storage, vec3/vec4 align 4, pas de padding [D19] pos: vec3, vel: vec3, life: f32, max_life: f32, size: f32, size_growth: f32, angle: f32, angular_vel: f32, color: vec4, uv_rect: vec4, // [D15] } struct EmitterParams { // espace uniform : vec3 align 16 → padding conservé [D19] dt: f32, gravity: vec3, drag: f32, alive_count: f32, pool_size: f32, alpha_scale: f32, // [D18] _pad: vec2, // → total 48 B } @group(0) @binding(0) var particles: array; @group(0) @binding(1) var params: EmitterParams; @compute @workgroup_size(64) fn cs_update() { let idx = u32(InvocationIndex); if idx >= u32(params.pool_size) { return; } var p = particles[idx]; // Skip les particules mortes if p.life <= 0.0 { return; } // Intégration (semi-implicit Euler) p.vel += params.gravity * params.dt; p.vel *= max(0.0, 1.0 - params.drag * params.dt); p.pos += p.vel * params.dt; // Vie p.life -= params.dt; // Size p.size = max(0.0, p.size + p.size_growth * params.dt); // Rotation p.angle += p.angular_vel * params.dt; // Fade out [D18] p.color.a = params.alpha_scale * clamp(p.life / p.max_life, 0.0, 1.0); // Kill if p.life <= 0.0 { p.life = 0.0; p.color.a = 0.0; } particles[idx] = p; } ``` > ⚠️ **[D17] Supplanté** : le count est exact, produit par la compaction fused dans ce compute (cf §18). > > **alive_count** (version initiale) : pour la v1, le count est géré par le CPU (le driver compte les > spawns/kills). Le compute ne fait que l'intégration. Un atomic dans le compute > serait plus "pur GPU" mais ajoute de la complexité sans gain visible en v1. ### 8.2 Vertex shader (billboard) ```wgsl // particle_billboard.wgsl (vertex) struct Particle { // 80 B — espace storage, vec3/vec4 align 4, pas de padding [D19] pos: vec3, vel: vec3, life: f32, max_life: f32, size: f32, size_growth: f32, angle: f32, angular_vel: f32, color: vec4, uv_rect: vec4, // [D15] } struct CameraParams { view: mat4x4, proj: mat4x4, } struct VsOut { @builtin(position) clip: vec4, @location(0) frag_color: vec4, @location(1) uv: vec2, } @group(0) @binding(0) var camera: CameraParams; @group(0) @binding(1) var particles: array; @group(0) @binding(2) var count_buf: f32; const QUAD: array, 4> = array, 4>( vec2(-0.5, -0.5), vec2( 0.5, -0.5), vec2( 0.5, 0.5), vec2(-0.5, 0.5), ); @vertex fn vs_main( @builtin(vertex_index) vi: u32, @builtin(instance_index) ii: u32, ) -> VsOut { var out: VsOut; // ⚠️ [D17] version initiale : early-out + draw(4, max_count). // → remplacé par index buffer compact + drawIndirect, cf §18. // Early-out si au-delà du count alive if f32(ii) >= count_buf { out.clip = vec4(0.0, 0.0, -1.0, 1.0); // hors écran out.frag_color = vec4(0.0); out.uv = vec2(0.0); return out; } let p = particles[ii]; // Quad unitaire let q = QUAD[vi]; // Rotation 2D let c = cos(p.angle); let s = sin(p.angle); let rot = vec2(q.x * c - q.y * s, q.x * s + q.y * c) * p.size; // Billboard : axes caméra (right, up) extraits de la view matrix let right = vec3(camera.view[0][0], camera.view[1][0], camera.view[2][0]); let up = vec3(camera.view[0][1], camera.view[1][1], camera.view[2][1]); let world = p.pos + right * rot.x + up * rot.y; out.clip = camera.proj * camera.view * vec4(world, 1.0); out.frag_color = p.color; out.uv = q + vec2(0.5); // 0..1 return out; } ``` ### 8.3 Fragment shader ```wgsl // particle_billboard.wgsl (fragment) struct FsIn { @builtin(position) clip: vec4, @location(0) frag_color: vec4, @location(1) uv: vec2, } @group(0) @binding(3) var samp: sampler; @group(0) @binding(4) var tex: texture_2d; @fragment fn fs_main(in: FsIn) -> @location(0) vec4 { let t = textureSample(tex, samp, in.uv); return in.frag_color * t; } ``` > Pour le blending **Additive**, le hardware fait `src × 1 + dst × 1` → le résultat > s'additive naturally. Pour **Alpha**, `src × src.a + dst × (1-src.a)`. > Le shader est identique dans les 2 cas — le blending est dans le pipeline. --- ## 9. Frame loop (ordre exact) ``` App::render_scene(frame) : │ ├─ 1. Shadow pass (si ombres) │ ├─ 2. Opaque pass (entities, indirect draws) │ ├─ 3. PARTICLES (nouveau) │ │ │ ├─ Pour chaque pool actif, dans l'ordre : │ │ │ │ a. driver.pre_compute(queue, pool, dt) │ │ → GPU : écrit spawns + params │ │ → CPU : simule + upload full │ │ → Manual : upload si dirty │ │ │ │ b. Si driver.needs_compute() : │ │ → encoder.dispatch_workgroups(ceil(N/64)) │ │ │ │ c. driver.post_compute(queue, pool) │ │ │ │ d. Si driver.needs_draw() : │ │ → render_pass.set_pipeline(pool.pipeline) │ │ → render_pass.set_bind_group(0, pool.bind_group) │ │ → render_pass.draw_indirect(pool.indirect_args, 0) // [D17] args écrits par le compute/driver │ │ │ └─ (fin pools) │ ├─ 4. Post-process (bloom → DoF → fog → tone mapping) │ └─ 5. Present ``` > ⚠️ **[D17] NOTE SUPPLANTÉE** : voir §18 — compaction + `drawIndirect`, plus de early-out. > > **Note sur `draw(4, max_count)`** (version initiale) : on drawe TOUS les slots du pool (max_count > instances). Le vertex shader fait un early-out (`if ii >= count → clip hors écran`). > C'est moins élégant qu'un indirect draw mais : > - Évite un buffer indirect + un compute de réduction > - Le GPU skip les instances mortes très tôt (vertex shader) > - Pour 50K instances "mortes", le coût est ~0 (vertex shader trivial) > > **V2** : indirect draw avec `alive_count` exact (nécessite un buffer indirect > écrit par le compute ou un 2e dispatch de réduction). --- ## 10. API utilisateur (résumé) ### 10.1 Setup ```rust impl AppHandler for MyScene { fn setup(&mut self, app: &mut wsg_lib::App) { // 1. Créer la texture (optionnel — None = disque blanc) app.scene.add_texture("fire_tex", fire_texture); // 2. Créer le pool app.scene.create_particle_pool("fire_pool", ParticlePoolConfig { max_count: 10_000, texture: Some("fire_tex"), blending: BlendingMode::Additive, }).unwrap(); // 3. Attacher un driver app.scene.attach_gpu_emitter("fire_pool", GpuEmitterConfig::fire()).unwrap(); app.scene.set_emitter_position("fire_pool", [0.0, 0.5, 0.0]); // Autre pool avec driver CPU app.scene.create_particle_pool("debris_pool", ParticlePoolConfig { max_count: 5_000, texture: None, blending: BlendingMode::Alpha, }).unwrap(); app.scene.attach_cpu_emitter("debris_pool", CpuEmitterConfig { // ... custom_force: Some(Arc::new(|p, scene| { // Attraction vers le centre -p.pos.map(|c| c * 2.0) })), ..CpuEmitterConfig::explosion() }).unwrap(); // Pool manual app.scene.create_particle_pool("custom_pool", ParticlePoolConfig { max_count: 1_000, texture: None, blending: BlendingMode::Alpha, }).unwrap(); app.scene.attach_manual("custom_pool").unwrap(); } } ``` ### 10.2 Per-frame ```rust impl AppHandler for MyScene { fn update(&mut self, app: &mut wsg_lib::App) { // GPU emitter : rien à faire (le driver gère tout) // CPU emitter : rien à faire non plus (le driver simule) // Manual : l'utilisateur écrit if let Some(handle) = app.scene.manual_handle("custom_pool") { for i in 0..100 { let t = (i as f32 / 100.0) * std::f32::consts::TAU; handle.set(i, Particle { pos: [t.cos() * 2.0, 0.0, t.sin() * 2.0], vel: [0.0; 3], life: 1.0, max_life: 1.0, size: 0.1, size_growth: 0.0, angle: t, angular_vel: 1.0, color: [1.0, 1.0, 1.0, 1.0], uv_rect: [0.0, 0.0, 1.0, 1.0], // [D15] ..Particle::ZERO }); } handle.set_count(100); handle.upload(); } } } ``` ### 10.3 Runtime control ```rust // Pauser/reprendre un pool scene.set_pool_active("fire_pool", false); scene.set_pool_active("fire_pool", true); // Repositionner un émetteur scene.set_emitter_position("fire_pool", [1.0, 2.0, 3.0]); // Burst (émettre N particules d'un coup, en plus du rate continu) scene.emitter_burst("fire_pool", 200); // Changer le driver (swap, pas de reallocation) scene.detach_driver("fire_pool"); scene.attach_cpu_emitter("fire_pool", CpuEmitterConfig::smoke()).unwrap(); ``` --- ## 11. Méthodes Scene (nouveaux) ```rust impl Scene { /// Créer un pool de particules (alloue le buffer + pipeline). /// Le pool est inactif jusqu'à l'attachement d'un driver. pub fn create_particle_pool(&mut self, id: &str, config: ParticlePoolConfig) -> Result<(), String>; /// Attacher un driver GPU (compute emitter) au pool. /// Remplace tout driver précédemment attaché. pub fn attach_gpu_emitter(&mut self, pool_id: &str, config: GpuEmitterConfig) -> Result<(), String>; /// Attacher un driver CPU (simulation Rust + upload) au pool. pub fn attach_cpu_emitter(&mut self, pool_id: &str, config: CpuEmitterConfig) -> Result<(), String>; /// Attacher le mode manual (l'utilisateur écrit le buffer directement). pub fn attach_manual(&mut self, pool_id: &str) -> Result<(), String>; /// Détacher le driver (le pool devient inactif). pub fn detach_driver(&mut self, pool_id: &str); /// Définir la position d'un émetteur (driver GPU/CPU). pub fn set_emitter_position(&mut self, pool_id: &str, pos: [f32; 3]); /// Pauser/reprendre un pool. pub fn set_pool_active(&mut self, pool_id: &str, active: bool); /// Burst : émettre N particules immédiatement (en plus du rate). pub fn emitter_burst(&mut self, pool_id: &str, count: u32); /// Obtenir un handle pour le driver manual (None si pas en mode manual). pub fn manual_handle(&mut self, pool_id: &str) -> Option>; } ``` --- ## 12. Non-régression et opt-in | Condition | Coût | |-----------|------| | Aucun pool créé | **Zéro**. Pas de buffer, pas de pipeline, pas de draw. | | Pool créé, pas de driver | Buffer alloué (N × 80 B). Pas de compute, pas d'upload, pas de draw (args = 0 → no-op). | | Pool + driver, `active = false` | Idem ci-dessus. | | Pool + driver GPU actif | 1 compute dispatch + 1 draw. CPU : spawns seulement. | | Pool + driver CPU actif | 1 write_buffer + 1 draw. CPU : O(N) simulation. | | Pool + driver Manual actif | 1 write_buffer (si dirty) + 1 draw. | **Zéro impact** sur les scènes existantes sans particules. --- ## 13. Texture par défaut (si `texture: None`) Texture **16×16** générée au build : ``` . . . . . . . . . + + + + + + . . + # # # # + . . + # # # # + . . + # # # # + . . + # # # # + . . + + + + + + . . . . . . . . . ``` Disque doux (alpha radiale : 1.0 au centre → 0.0 au bord). RGBA blanc. Utilisée par défaut si l'utilisateur ne fournit pas de texture. --- ## 14. Limitations v1 | Limitation | Justification | V2 | |------------|---------------|-----| | 1 driver par pool | Simplifie la gestion du buffer | Multi-drivers (rangs du pool) | | 1 blending par pool | Pipeline figé | Multi-pipelines par pool | | Fade alpha linéaire | Simple, couvre 90% | Color ramp (mix start→end) | | Pas de collision particule↔objet | O(n×m) | Driver CPU + custom_force | | Pas de collision particule↔particule | O(n²) | Spatial hash (V3) | | ~~`draw(4, max_count)` + early-out~~ → **[D17] compaction + indirect en v1** (§18) | Supplanté | — | | Spawn CPU uniquement | Le random au spawn est CPU | GPU spawn (curated random) | | Pas de texture sheet **animée** ([D15] donne déjà la zone UV statique par particule en v1) | Zone UV statique par vie | UV animation f(t) (V2) | | Pas de sorting (alpha) | Additive = commutatif. Alpha = approximation | Back-to-front sort (V2) | --- ## 15. Questions ouvertes (à trancher par DRAFT) | # | Question | Tendance | Impact | |---|----------|----------|--------| | 1 | `draw(4, max_count)` vs indirect draw | **Tranché → D17** (compaction + indirect en v1, §18) | — | | 2 | alive_count : CPU count vs atomic GPU | **Tranché → D17** (count exact GPU via compaction, §18) | — | | 3 | Multi-pools même texture : partager le pipeline ? | Oui (Arc partagé) | Mémoire | | 4 | Le pool est-il `Send` ? (multi-thread spawn) | Oui — le buffer est GPU, le driver est `Send` | Flexibilité | | 5 | `emitter_burst` : overwrite les plus vieilles si pool saturé ? | Oui (ring buffer sémantique) | Robustesse | | 6 | Le driver CPU fait-il un upload **full** ou **dirty range** ? | v1 : full (simple). V2 : dirty range. | Bandwidth | | 7 | Les pools sont-ils rendus dans un seul render pass ou un par pool ? | Un seul pass, `set_pipeline` + `set_bind_group` par pool | Draw calls | | 8 | Faut-il un `ParticlePool::alive_count()` public (pour l'UI) ? | Oui (le driver maintient le count) | Debug | --- ## 16. Estimation de complexité (par DRAFT) | DRAFT | Contenu | Lignes estimées | |-------|---------|-----------------| | **Étape A** | Struct `Particle` + `ParticlePoolConfig` + enums + `ParticlePool` (buffer + pipeline + draw) | ~250 | | **Étape B** | Compute shader + driver GPU (spawn + dispatch) | ~200 | | **Étape C** | Driver CPU (simulation Rust + upload) + `custom_force` | ~150 | | **Étape D** | Driver Manual + `ManualPoolHandle` | ~100 | | **Étape E** | Intégration Renderer (frame loop, order) + `Scene` methods | ~150 | | **Étape F** | Presets + Example `particles.rs` | ~150 | | **Étape G** | Tests (WGSL validate + layout + pool) | ~80 | | **Total** | | **~1 080 lines** | Chaque étape est un DRAFT séparé, testable indépendamment. --- ## 17. Résumé des décisions | # | Décision | Justification | |---|----------|---------------| | D1 | **Pool ≠ Driver** (séparation stricte) | Flexibilité, swappability, zéro waste | | D2 | **3 drivers** : GPU, CPU, Manual | Gradient de contrôle | | D3 | **1 driver par pool** (v1) | Simplicité. Multi-drivers = V2. | | D4 | **80 bytes/particule** (avec `uv_rect` [D15] ; **confirmé sans padding** [D19]) | Couvre pos/vel/life/size/angle/color/uv. Storage align 4 — Rust = WGSL [D19]. | | D5 | **Billboard camera-facing** (pas world-facing) | Standard pour les VFX. Plus simple. | | D6 | **Quad via vertex_index** (vertex layout vide ; le slot arrive par storage `compact_index[ii]` [D17/D19]) | Cohérent avec TM/bloom. | | D7 | ~~**`draw(4, max_count)` + early-out**~~ → **supplanté par D17** (§18) | — | | D8 | **Spawn toujours CPU** (même driver GPU) | Le random + la décision "qui spawn" est CPU. Le GPU intègre. | | D9 | **Blend figé au pipeline** (1 mode par pool) | Pas de switch de pipeline par frame. | | D10 | **Depth test oui, depth write non** | Transparence correcte sans artefacts. | | D11 | **Texture par défaut : disque 16×16** | Zéro config pour un résultat acceptable. | | D12 | **Pool inactif si pas de driver** | Zéro coût. L'allocation buffer est le seul coût. | | D13 | **Preset methods** (`GpuEmitterConfig::fire()`) | UX : 1 ligne pour un effet correct. | | D14 | **`custom_force` (driver CPU)** | Le seul cas où CPU > GPU : interactions. | | D15 | **UV par particule** (`uv_rect` vec4 → struct 80 B) | WebGPU : pas d'indexage dynamique de textures → multi-motifs via UV. §18 | | D16 | **Layout complet figé, pas d'opt-in par attribut** | Gain nul, explosion de variants, coupling pool/émetteur. §18 | | D17 | **Compaction + indirect draw** (fused dans le compute ; détails corrigés par D19) | Scale avec count actif, pas capacity. Count exact GPU. §18 | | D18 | **Randomisation par particule au spawn** (`color_range` + `uv_rects` + `alpha_scale`) | Principe D16 : les plages vivent dans la config (coût GPU zéro), le résultat dans l'état. §18 | | D19 | **Audit layout/détails** : 80 B sans padding, un seul struct partagé, `compact_index` ×1 par storage, args 16 B | L'align uniform space du draft avait fuité en storage space. §18 | --- ## 18. Décisions de la session de design (à appliquer par DRAFT) Ces 5 décisions tranchent les questions §15 (Q1, Q2), supplantent D4/D6/D7 et corrigent l'application initiale de D15/D17 (audit D19). Le reste du doc décrit la version initiale : appliquer les impacts listés ici. ### D15 — UV par particule (`uv_rect`) - `uv_rect: vec4` (ox, oy, sx, sy en UV texture) ajouté à `Particle` → struct 64 → **80 B** (appliqué dans §2, §8.1, §8.2 — D19 a corrigé l'application initiale : padding résiduel de l'espace uniform). - VS : `out.uv = uv_rect.xy + (q + 0.5) * uv_rect.zw` (au lieu de `q + 0.5`). Défaut (0,0,1,1) = comportement d'avant. - **Pourquoi** : WebGPU interdit l'indexage dynamique de textures en uniform (1 texture = 1 bind group = 1 pipeline). La variété de motifs ne peut venir QUE des UV → multi-motifs depuis UNE texture (atlas partagé entre presets). - Statique par vie (fixé au spawn). Animation d'atlas (uv = f(t)) = **v2** (cf §14). - Impact : §2/§3.2/§6/§8.1/§8.2 appliqués ; §5.2/§5.4 appliqués par D18 (la zone unique devient la palette `uv_rects`) ; §10.2 appliqué. ### D16 — Layout COMPLET et figé, PAS d'opt-in par attribut - Le pool a TOUJOURS le layout complet. Pas de variante « pool léger » (sans color/angle/vel). - **Pourquoi** : a) gain nul en pratique : la VRAM est allouée UNE fois (capacity = le réglage utilisateur), pas par frame ; b) opt-in = explosion de variants : 4 attributs optionnels → jusqu'à 16 combos pipeline/shader (attribut absent = layout WGSL différent = shader variant différent) ; c) piège de coupling : le pool est créé UNE FOIS et partagé — impossible de l'élargir après coup pour un émetteur voulant l'attribut manquant. - Le « enable » utile est au niveau **config d'émetteur** (quels champs le driver écrit/randomise au spawn) — coût GPU zéro. - Impact : aucun (le §2 reste le seul layout). Ferme toute question « pool minimal ». ### D17 — Compaction + indirect draw (supplante D7, tranche Q1+Q2) - Emulation d'instancing standard WebGPU : buffer `compact_index` (**1 u32 par slot** [D19], pas 4 copies) lié en **storage** (binding 2) — le VS fait `let slot = compact_index[ii]` puis `particles[slot]`. Le vertex layout reste **vide** (quad généré en shader via `QUAD[vi]`, pattern TM). - La **compaction est FUSIONNÉE dans le compute d'intégration** (~15 lignes WGSL, pas de 3e pass) : chaque slot alive (life > 0 après intégration) fait `atomicAdd(&compact_count)` et écrit son slot **une fois** dans `compact_index` [D19] ; le dernier workgroup (détection par atomic global vs nb de workgroups) écrit les args indirect. - Draw : **`drawIndirect`** (vertexCount = 4, instanceCount = alive → 4 × alive vertices au total) au lieu de `draw(4, max_count)` + early-out. Le VS n'a plus besoin du count → `count_buffer` remplacé par `indirect_args` (**16 B : 4 × u32**, draw non-indexé [D19]). - **Pourquoi** : la compaction scale avec le count **actif** ; le template statique (early-out) taxe la **capacity** à CHAQUE frame — scène idle avec grand pool = 0.3–1 ms/frame de VS inutile. Coût compaction : pire cas ~0.3 ms sur iGPU faible (scène active). - Conséquence : le count est **exact et GPU** (sortie de la compaction) → tranche Q2 (plus de count CPU estimé). - Par driver : **GPU** = compaction fused (ci-dessus). **CPU** = sait exactement les slots alive → build lui-même l'index (4 copies) + args au CPU. **Manual** = `set_count` remplit l'index identité (l'utilisateur pack ses alive en tête de buffer) + args. - Non-régression : pool sans driver → args indirect = 0 → `drawIndirect` no-op (§12 inchangé). - Impact : §3.2/§6/§7.2/§9 appliqués (buffers + bind group + pipeline + frame loop, détails corrigés par D19) ; §8.1 (compaction fused) + §8.2 (supprimer early-out + `count_buf`, lire l'index) + §4.2/§4.3/§4.4 (chaque driver produit index + args) → à appliquer par le DRAFT B (les corps v0 restent, notes ci-dessus). --- ### D18 — Randomisation par particule au spawn : `color_range` + `uv_rects` + `alpha_scale` Demande : « randomiser un peu la couleur, et randomiser un peu le rect UV ». Principe général (extension de D16) : **tout** champ de l'état 80 B randomisable au spawn a sa **plage** dans la config d'émetteur (coût GPU zéro — les tirages sont CPU au spawn [D8]) et son **résultat** dans l'état (par particule). Trois paramètres ajoutés aux configs d'émetteurs : - `color_range: ([f32; 4], [f32; 4])` — [min, max] par canal **RGB** (le canal alpha est ignoré → `alpha_scale`). Spawn : `c = lerp(min, max, rand)` par canal. - `uv_rects: [[f32; 4]; UV_RECTS_MAX]` + `uv_rects_count: u8` (UV_RECTS_MAX = 8) — **palette** de zones UV (généralise la zone unique de D15). Spawn : tirage uniforme d'un rect. `count = 1` → défaut D15 (toute la texture, comportement inchangé). - `alpha_scale: f32` (0..1, défaut 1.0) — alpha cible, modulé par la vie normalisée. **Source unique de l'alpha** : l'intégration (compute GPU ou simulation CPU) réécrit à chaque frame `color.a = alpha_scale × clamp(life/max_life, 0, 1)` (remplace l'« alpha initiale » de `color`). Driver Manual : pas d'intégration → l'utilisateur écrit `color` (avec alpha) directement, `alpha_scale` est inopérant. `EmitterParams` (uniform lu par le compute) gagne `alpha_scale` → struct **48 B** (toujours dans le buffer 64 B déclaré). Impact : §4.2/§4.3 (configs — appliqué), §5.4 (presets : colonne `color` → `color_range` + `α_scale` — appliqué), §6/§8.1 (EmitterParams + formule de fade — appliqué), §10.2 (manual : pas de config, écriture directe — inchangé). **Pas dans l'étape A du DRAFT** (étapes B/C/D). ### D19 — Audit layout/détails (préparation DRAFT) : 80 B sans padding, un seul struct partagé, `compact_index` ×1 Audit fait en écrivant le DRAFT. Trois corrections, toutes dans le même sens : **l'alignement de l'espace uniform du draft avait fuité dans l'espace storage**. 1. **Pas de padding dans les structs particule.** `Particle` vit en **espace storage**, où `vec3` a **align 4** (l'espace uniform a align 16 — d'où venaient `pad0`/`pad1` du draft). Le Rust `[f32; 3]` est aussi align 4 → layout naturellement identique dans les deux langages : pos@0, vel@12, life@24, max_life@28, size@32, size_growth@36, angle@40, angular_vel@44, color@48, uv_rect@64 → **80 B exactement, sans padding**. L'état précédent (avec padding) faisait en réalité **88 B** — en contradiction avec la mention « 80 bytes » (le commentaire d'offset 64 de `uv_rect` incohérent avec color@56+16=72). Appliqué : §2, §8.1, §8.2. 2. **Un seul struct `Particle` partagé** entre les entry points compute et render : le double struct `ParticleAlive` du draft (64 B, pad) était un **préfixe incompatible** du `Particle` padé (life@32 vs life@24) → le cull aurait lu les mauvais offsets. Supprimé : un storage load ne dépend pas de la taille du struct déclaré, la vue « sous-ensemble » n'a aucun gain. (`EmitterParams` **conserve** son padding — espace uniform, vec3 align 16.) 3. **`compact_index` : 1 u32 par slot (pas 4 copies), lu par **storage binding** (pas d'attribut vertex).** Le buffer est lié en storage (binding 2) et le VS fait `compact_index[ii]` — le vertex layout du pipeline reste vide (pattern TM). La compaction écrit **un** u32 par particule alive (pas quatre). `indirect_args` = **16 B** (4 × u32, `drawIndirect` non-indexé) et non 20 B. Usage du buffer : STORAGE | COPY_DST (pas de VERTEX). Appliqué : §3.2, §6, §7.2, §17 (D6/D17), §18 (D17). Règle future : ajouter un champ à `Particle` → alignement storage (vec3/vec4 = align 4, pas de padding) + mettre à jour les deux miroirs (Rust §2 + WGSL §8) + les tests de layout (DRAFT A).