archi particules
This commit is contained in:
@@ -0,0 +1,959 @@
|
||||
# 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 × 64 bytes) │ │
|
||||
│ │ • 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 (64 bytes)
|
||||
|
||||
```rust
|
||||
/// Miroir du struct WGSL `Particle`.
|
||||
/// 64 bytes, `#[repr(C)]`, `Pod + Zeroable`.
|
||||
#[repr(C)]
|
||||
#[derive(Copy, Clone, Pod, Zeroable)]
|
||||
pub struct Particle {
|
||||
pub pos: [f32; 3], // offset 0 — position monde (xyz)
|
||||
pub _pad0: f32, // offset 12
|
||||
pub vel: [f32; 3], // offset 16 — vélocité (xyz)
|
||||
pub _pad1: f32, // offset 28
|
||||
pub life: f32, // offset 32 — vie restante (seconds)
|
||||
pub max_life: f32, // offset 36 — vie initiale (pour fade normalisé)
|
||||
pub size: f32, // offset 40 — taille courante (world units)
|
||||
pub size_growth: f32, // offset 44 — croissance par seconde (+ = grandir, - = rétrécir)
|
||||
pub angle: f32, // offset 48 — rotation 2D courante (radians)
|
||||
pub angular_vel: f32, // offset 52 — vitesse angulaire (rad/s)
|
||||
pub color: [f32; 4], // offset 56 — RGBA (alpha modulée par le driver)
|
||||
// Total : 64 bytes
|
||||
}
|
||||
```
|
||||
|
||||
### 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 |
|
||||
|
||||
> 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 |
|
||||
|
||||
---
|
||||
|
||||
## 3. Le Pool (ressource GPU)
|
||||
|
||||
### 3.1 Création
|
||||
|
||||
Le pool est créé **une fois** au setup. Il alloue :
|
||||
- Le buffer storage (N × 64 bytes)
|
||||
- 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<String>,
|
||||
/// 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 × 64 bytes). 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 (buffer + texture + sampler + camera uniform ref).
|
||||
pub(crate) bind_group: wgpu::BindGroup,
|
||||
/// Sampler (Linear, ClampToEdge).
|
||||
pub(crate) sampler: wgpu::Sampler,
|
||||
/// Buffer uniform 4 bytes : alive_count (écrit par le driver, lu par le vertex shader).
|
||||
pub(crate) count_buffer: wgpu::Buffer,
|
||||
/// Taille max du pool.
|
||||
pub max_count: u32,
|
||||
/// Driver actuellement attaché (None = pool inactif).
|
||||
pub(crate) driver: Option<Box<dyn ParticleDriver>>,
|
||||
/// 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. Écrire alive_count estimé dans count_buffer
|
||||
2. Compute dispatch (le pool fait le dispatch) :
|
||||
- workgroups = ceil(max_count / 64)
|
||||
- Le shader intègre TOUS les slots alive
|
||||
- Le shader écrit alive_count exact dans count_buffer (atomic ou 2e pass)
|
||||
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 RGBA (alpha initiale = 1.0, modulée par life/max_life).
|
||||
pub color: [f32; 4],
|
||||
}
|
||||
```
|
||||
|
||||
**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. Écrire alive_count dans count_buffer
|
||||
2. Pas de compute dispatch
|
||||
3. post_compute : (rien)
|
||||
4. Draw (si alive_count > 0)
|
||||
```
|
||||
|
||||
**Config** : identique à `GpuEmitterConfig` (mêmes paramètres de simulation).
|
||||
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<Arc<dyn Fn(&Particle, &Scene) -> [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 | blending |
|
||||
|--------|-------|----------|----------|-------|---------|------|------|------|--------|---------|-------|----------|
|
||||
| **Explosion** | Point | Sphere | — | 2-8 | (0,-9.8,0) | 0.1 | 0.5-1.5 | 0.05-0.15 | -0.05 | ±5 | [1,0.8,0.3,1] | 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.6,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 | [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.5,0.5,0.5,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.5,0,1] | 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 | [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,1,0.8,1] | Additive |
|
||||
| **Débris** | Point | Sphere | — | 1-6 | (0,-9.8,0) | 0.02 | 1-3 | 0.03-0.08 | 0 | ±10 | [0.6,0.4,0.2,1] | Alpha |
|
||||
|
||||
Disponibles comme constructors : `GpuEmitterConfig::fire()`, `::smoke()`, `::rain()`, etc.
|
||||
|
||||
---
|
||||
|
||||
## 6. Buffers GPU
|
||||
|
||||
| Buffer | Type | Taille | Écrit par | Lu par |
|
||||
|--------|------|--------|-----------|--------|
|
||||
| `particle_data` | Storage (RW) | N × 64 B | Driver (spawn/update) | Compute + Render |
|
||||
| `count_buffer` | Uniform (RW) | 4 B | Driver / Compute | Vertex shader (instance_count) |
|
||||
| `emitter_params` | Uniform (RO) | 64 B | Driver GPU (par frame) | Compute shader |
|
||||
| `camera_params` | Uniform (RO) | 160 B | Renderer (existant) | Vertex shader (view, proj) |
|
||||
|
||||
### `emitter_params` (uniform, 64 bytes)
|
||||
|
||||
```wgsl
|
||||
struct EmitterParams {
|
||||
dt: f32, // delta time de la frame
|
||||
gravity: vec3<f32>, // accélération
|
||||
drag: f32, // friction
|
||||
alive_count: f32, // count courant (écrit par le compute via atomic)
|
||||
pool_size: f32, // taille max du pool
|
||||
_pad: vec2<f32>, // alignment
|
||||
}
|
||||
// Total : 4+12+4+4+4+8 = 36 → pad à 48 (align 16)
|
||||
```
|
||||
|
||||
### 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) | VERTEX |
|
||||
| 0 | 1 | Storage RO | `particle_data` | VERTEX |
|
||||
| 0 | 2 | Uniform RO | `count_buffer` (alive_count) | VERTEX |
|
||||
| 0 | 3 | Sampler | Sampler | FRAGMENT |
|
||||
| 0 | 4 | Texture | Texture particule | FRAGMENT |
|
||||
|
||||
> **Note** : le `count_buffer` est lu par le vertex shader pour déterminer
|
||||
> `instance_count` (via `@builtin(instance_index)` et un early-out si `ii >= count`).
|
||||
|
||||
---
|
||||
|
||||
## 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 + storage RO + count + sampler + texture
|
||||
},
|
||||
vertex: vs_main (billboard),
|
||||
fragment: fs_main (texture × color),
|
||||
primitive: TriangleList,
|
||||
vertex_buffer_layouts: [], // PAS de vertex buffer ! (quad généré en shader)
|
||||
multisample: sample_count du contexte,
|
||||
color_states: [blending mode du pool],
|
||||
depth_stencil: Some(LessEqual, ALWAYS), // depth test oui, depth write non (transparence)
|
||||
}
|
||||
```
|
||||
|
||||
> **Pas de vertex buffer** : le quad billboard est généré dans le vertex shader
|
||||
> via `@builtin(vertex_index)` (4 vertices) — même pattern que le fullscreen triangle
|
||||
> du tone mapping, mais avec 4 sommets au lieu de 3.
|
||||
|
||||
### 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 {
|
||||
pos: vec3<f32>,
|
||||
pad0: f32,
|
||||
vel: vec3<f32>,
|
||||
pad1: f32,
|
||||
life: f32,
|
||||
max_life: f32,
|
||||
size: f32,
|
||||
size_growth: f32,
|
||||
angle: f32,
|
||||
angular_vel: f32,
|
||||
color: vec4<f32>,
|
||||
}
|
||||
|
||||
struct EmitterParams {
|
||||
dt: f32,
|
||||
gravity: vec3<f32>,
|
||||
drag: f32,
|
||||
alive_count: f32,
|
||||
pool_size: f32,
|
||||
pad: vec2<f32>,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<storage, read_write> particles: array<Particle>;
|
||||
@group(0) @binding(1) var<uniform> 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
|
||||
p.color.a = 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;
|
||||
}
|
||||
```
|
||||
|
||||
> **alive_count** : 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 {
|
||||
pos: vec3<f32>, pad0: f32,
|
||||
vel: vec3<f32>, pad1: f32,
|
||||
life: f32, max_life: f32,
|
||||
size: f32, size_growth: f32,
|
||||
angle: f32, angular_vel: f32,
|
||||
color: vec4<f32>,
|
||||
}
|
||||
|
||||
struct CameraParams {
|
||||
view: mat4x4<f32>,
|
||||
proj: mat4x4<f32>,
|
||||
}
|
||||
|
||||
struct VsOut {
|
||||
@builtin(position) clip: vec4<f32>,
|
||||
@location(0) frag_color: vec4<f32>,
|
||||
@location(1) uv: vec2<f32>,
|
||||
}
|
||||
|
||||
@group(0) @binding(0) var<uniform> camera: CameraParams;
|
||||
@group(0) @binding(1) var<storage, read> particles: array<Particle>;
|
||||
@group(0) @binding(2) var<uniform> count_buf: f32;
|
||||
|
||||
const QUAD: array<vec2<f32>, 4> = array<vec2<f32>, 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;
|
||||
|
||||
// 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<f32>,
|
||||
@location(0) frag_color: vec4<f32>,
|
||||
@location(1) uv: vec2<f32>,
|
||||
}
|
||||
|
||||
@group(0) @binding(3) var samp: sampler;
|
||||
@group(0) @binding(4) var tex: texture_2d<f32>;
|
||||
|
||||
@fragment
|
||||
fn fs_main(in: FsIn) -> @location(0) vec4<f32> {
|
||||
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(4, pool.max_count) // early-out en shader
|
||||
│ │
|
||||
│ └─ (fin pools)
|
||||
│
|
||||
├─ 4. Post-process (bloom → DoF → fog → tone mapping)
|
||||
│
|
||||
└─ 5. Present
|
||||
```
|
||||
|
||||
> **Note sur `draw(4, max_count)`** : 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],
|
||||
..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<ManualPoolHandle<'_>>;
|
||||
}
|
||||
```
|
||||
|
||||
---
|
||||
|
||||
## 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 × 64 B). Pas de compute, pas d'upload, pas de draw. |
|
||||
| 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 | Simple, pas de buffer indirect | Indirect draw (V2) |
|
||||
| Spawn CPU uniquement | Le random au spawn est CPU | GPU spawn (curated random) |
|
||||
| Pas de texture sheet (UV animation) | 1 UV par particule | UV offset + animation |
|
||||
| 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 | v1 : early-out en shader. v2 : indirect. | Complexité du pipeline |
|
||||
| 2 | alive_count : CPU count vs atomic GPU | v1 : CPU (driver compte). Simple. | Précision du count |
|
||||
| 3 | Multi-pools même texture : partager le pipeline ? | Oui (Arc<RenderPipeline> 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 | **64 bytes/particule** | Couvre pos/vel/life/size/angle/color. Align 16. |
|
||||
| D5 | **Billboard camera-facing** (pas world-facing) | Standard pour les VFX. Plus simple. |
|
||||
| D6 | **Quad via vertex_index** (pas de vertex buffer) | Cohérent avec TM/bloom. 0 buffer vertices. |
|
||||
| D7 | **`draw(4, max_count)` + early-out** (v1) | Pas de buffer indirect. GPU skip les mortes. |
|
||||
| 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. |
|
||||
Reference in New Issue
Block a user