From 0fbcbfd6bd25f4d822bf4597dedab8988f9ee95e Mon Sep 17 00:00:00 2001 From: Thierry Berger Date: Fri, 11 Apr 2025 14:29:52 +0200 Subject: [PATCH 1/2] wip wip adding particles create a bigger buffer than necessary + store a u32 to know the actual number of particles Not working yet. --- examples2d/src/dynamic_add.rs | 108 +++++++++++++++++++++++++ examples2d/src/elastic_cut2.rs | 1 + examples2d/src/elasticity2.rs | 1 + examples2d/src/main.rs | 5 ++ examples2d/src/sand2.rs | 1 + examples3d/src/dynamic_add.rs | 68 ++++++++++++++++ examples3d/src/main.rs | 5 ++ examples3d/src/sand3.rs | 3 + src/grid/sort.wgsl | 14 ++-- src/models/mod.rs | 13 ++- src/pipeline.rs | 8 +- src/solver/g2p.rs | 1 + src/solver/p2g.wgsl | 8 +- src/solver/particle2d.rs | 110 +++++++++++++++++++++++--- src/solver/particle_update.rs | 3 +- src/solver/particle_update.wgsl | 16 ++-- src_testbed/lib.rs | 4 +- src_testbed/prep_vertex_buffer.rs | 3 +- src_testbed/prep_vertex_buffer2d.wgsl | 14 ++-- src_testbed/prep_vertex_buffer3d.wgsl | 14 ++-- src_testbed/step.rs | 2 + 21 files changed, 356 insertions(+), 46 deletions(-) create mode 100644 examples2d/src/dynamic_add.rs create mode 100644 examples3d/src/dynamic_add.rs diff --git a/examples2d/src/dynamic_add.rs b/examples2d/src/dynamic_add.rs new file mode 100644 index 0000000..ae54ad4 --- /dev/null +++ b/examples2d/src/dynamic_add.rs @@ -0,0 +1,108 @@ +use wgsparkl2d::rapier::prelude::{ColliderBuilder, RigidBodyBuilder}; +use wgsparkl_testbed2d::{wgsparkl, Callbacks, RapierData}; + +use bevy::render::renderer::RenderDevice; +use nalgebra::{vector, Vector2}; +use wgsparkl::models::DruckerPrager; +use wgsparkl::{ + models::ElasticCoefficients, + pipeline::MpmData, + solver::{Particle, SimulationParams}, +}; +use wgsparkl2d::solver::ParticleDynamics; +use wgsparkl_testbed2d::{AppState, PhysicsContext}; + +pub fn dynamic_demo( + device: RenderDevice, + app_state: &mut AppState, + callbacks: &mut Callbacks, +) -> PhysicsContext { + let mut rapier_data = RapierData::default(); + let device = device.wgpu_device(); + + let offset_y = 46.0; + let cell_width = 0.2; + let mut particles = vec![]; + let position = vector![0.5, 0.5] * cell_width / 2.0 + Vector2::y() * offset_y; + let density = 1000.0; + let radius = cell_width / 4.0; + particles.push(Particle { + position, + dynamics: ParticleDynamics::with_density(radius, density), + model: ElasticCoefficients::from_young_modulus(10_000_000.0, 0.2), + plasticity: Some(DruckerPrager::new(10_000_000.0, 0.2)), + phase: None, + color: None, + }); + + if !app_state.restarting { + app_state.num_substeps = 10; + app_state.gravity_factor = 1.0; + }; + + let params = SimulationParams { + gravity: vector![0.0, -9.81] * app_state.gravity_factor, + dt: (1.0 / 60.0) / (app_state.num_substeps as f32), + padding: 0.0, + }; + + /* + * Static platforms. + */ + let rb = RigidBodyBuilder::fixed().translation(vector![35.0, -1.0]); + let rb_handle = rapier_data.bodies.insert(rb); + let co = ColliderBuilder::cuboid(42.0, 1.0); + rapier_data + .colliders + .insert_with_parent(co, rb_handle, &mut rapier_data.bodies); + + let rb = RigidBodyBuilder::fixed() + .translation(vector![-25.0, 45.0]) + .rotation(0.5); + let rb_handle = rapier_data.bodies.insert(rb); + let co = ColliderBuilder::cuboid(1.0, 52.0); + rapier_data + .colliders + .insert_with_parent(co, rb_handle, &mut rapier_data.bodies); + + let rb = RigidBodyBuilder::fixed() + .translation(vector![95.0, 45.0]) + .rotation(-0.5); + let rb_handle = rapier_data.bodies.insert(rb); + let co = ColliderBuilder::cuboid(1.0, 52.0); + rapier_data + .colliders + .insert_with_parent(co, rb_handle, &mut rapier_data.bodies); + + callbacks.0.push(Box::new( + move |_, physics: &mut PhysicsContext, _, _, queue| { + physics.data.particles.push( + queue, + &Particle { + position, + dynamics: ParticleDynamics::with_density(radius, density), + model: ElasticCoefficients::from_young_modulus(2_000_000_000.0, 0.2), + plasticity: Some(DruckerPrager::new(2_000_000_000.0, 0.2)), + phase: None, + color: None, + }, + ) + }, + )); + + let data = MpmData::new( + device, + params, + &particles, + 1_000, + &rapier_data.bodies, + &rapier_data.colliders, + cell_width, + 60_000, + ); + PhysicsContext { + data, + rapier_data, + particles, + } +} diff --git a/examples2d/src/elastic_cut2.rs b/examples2d/src/elastic_cut2.rs index 6a32239..97748e9 100644 --- a/examples2d/src/elastic_cut2.rs +++ b/examples2d/src/elastic_cut2.rs @@ -103,6 +103,7 @@ pub fn elastic_cut_demo( device, params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, cell_width, diff --git a/examples2d/src/elasticity2.rs b/examples2d/src/elasticity2.rs index 638f038..e35b96c 100644 --- a/examples2d/src/elasticity2.rs +++ b/examples2d/src/elasticity2.rs @@ -85,6 +85,7 @@ pub fn elastic_demo( device, params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, cell_width, diff --git a/examples2d/src/main.rs b/examples2d/src/main.rs index 3bd46f5..07231f3 100644 --- a/examples2d/src/main.rs +++ b/examples2d/src/main.rs @@ -1,6 +1,7 @@ use bevy::prelude::*; use wgsparkl_testbed2d::{init_testbed, SceneInitFn, SceneInits}; +mod dynamic_add; mod elastic_cut2; mod elasticity2; mod sand2; @@ -20,6 +21,10 @@ pub fn main() { fn register_scenes(world: &mut World) { let scenes: Vec<(String, SceneInitFn)> = vec![ ("sand".to_string(), Box::new(sand2::sand_demo)), + ( + "dynamic add".to_string(), + Box::new(dynamic_add::dynamic_demo), + ), ("elastic".to_string(), Box::new(elasticity2::elastic_demo)), ( "elastic cut".to_string(), diff --git a/examples2d/src/sand2.rs b/examples2d/src/sand2.rs index 0f57d5b..5c1edbe 100644 --- a/examples2d/src/sand2.rs +++ b/examples2d/src/sand2.rs @@ -163,6 +163,7 @@ pub fn sand_demo( device, params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, cell_width, diff --git a/examples3d/src/dynamic_add.rs b/examples3d/src/dynamic_add.rs new file mode 100644 index 0000000..a851e04 --- /dev/null +++ b/examples3d/src/dynamic_add.rs @@ -0,0 +1,68 @@ +use wgsparkl_testbed3d::{wgsparkl, Callbacks, RapierData}; + +use bevy::render::renderer::RenderDevice; +use nalgebra::vector; +use rapier3d::prelude::{ColliderBuilder, RigidBodyBuilder}; +use wgsparkl::models::DruckerPrager; +use wgsparkl::{ + models::ElasticCoefficients, + pipeline::MpmData, + solver::{Particle, ParticleDynamics, SimulationParams}, +}; +use wgsparkl_testbed3d::{AppState, PhysicsContext}; + +use crate::utils::default_scene; + +pub fn dynamic_demo( + device: RenderDevice, + app_state: &mut AppState, + _callbacks: &mut Callbacks, +) -> PhysicsContext { + let mut rapier_data = RapierData::default(); + let device = device.wgpu_device(); + + if !app_state.restarting { + app_state.num_substeps = 20; + app_state.gravity_factor = 1.0; + }; + + let params = SimulationParams { + gravity: vector![0.0, -9.81, 0.0] * app_state.gravity_factor, + dt: (1.0 / 60.0) / (app_state.num_substeps as f32), + }; + + default_scene::spawn_ground_and_walls(&mut rapier_data); + + callbacks.0.push(Box::new( + move |_, physics: &mut PhysicsContext, _, _, queue| { + physics.data.particles.push( + queue, + &Particle { + position, + dynamics: ParticleDynamics::with_density(radius, density), + model: ElasticCoefficients::from_young_modulus(2_000_000_000.0, 0.2), + plasticity: Some(DruckerPrager::new(2_000_000_000.0, 0.2)), + phase: None, + color: None, + }, + ) + }, + )); + + let particles = vec![]; + let cell_width = 1.0; + let data = MpmData::new( + device, + params, + &particles, + &rapier_data.bodies, + &rapier_data.colliders, + cell_width, + 60_000, + ); + PhysicsContext { + data, + rapier_data, + particles, + } +} diff --git a/examples3d/src/main.rs b/examples3d/src/main.rs index 09158a5..742197c 100644 --- a/examples3d/src/main.rs +++ b/examples3d/src/main.rs @@ -4,6 +4,7 @@ use wgsparkl_testbed3d::{init_testbed, SceneInitFn, SceneInits}; pub mod utils; mod banana3; +mod dynamic_add; mod elastic_cut3; mod glb_to_point_cloud_color; mod heightfield3; @@ -24,6 +25,10 @@ pub fn main() { fn register_scenes(world: &mut World) { let scenes: Vec<(String, SceneInitFn)> = vec![ ("sand".to_string(), Box::new(sand3::sand_demo)), + ( + "dynamic addition".to_string(), + Box::new(dynamic_add::dynamic_demo), + ), ( "heightfield".to_string(), Box::new(heightfield3::heightfield_demo), diff --git a/examples3d/src/sand3.rs b/examples3d/src/sand3.rs index 2c7570a..31decbe 100644 --- a/examples3d/src/sand3.rs +++ b/examples3d/src/sand3.rs @@ -1,3 +1,5 @@ +use std::ptr::dangling; + use wgsparkl_testbed3d::{wgsparkl, Callbacks, RapierData}; use bevy::render::renderer::RenderDevice; @@ -88,6 +90,7 @@ pub fn sand_demo( device, params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, cell_width, diff --git a/src/grid/sort.wgsl b/src/grid/sort.wgsl index 979760c..327c6d4 100644 --- a/src/grid/sort.wgsl +++ b/src/grid/sort.wgsl @@ -5,18 +5,20 @@ @group(1) @binding(0) -var particles_pos: array; +var num_particles u32; @group(1) @binding(1) -var scan_values: array>; // This has to be atomic for finalize_particles_sort. Should it be a different buffer? +var particles_pos: array; @group(1) @binding(2) -var sorted_particle_ids: array; +var scan_values: array>; // This has to be atomic for finalize_particles_sort. Should it be a different buffer? @group(1) @binding(3) -var particle_node_linked_lists: array; +var sorted_particle_ids: array; @group(1) @binding(4) -var rigid_particles_pos: array; +var particle_node_linked_lists: array; @group(1) @binding(5) -var rigid_particle_node_linked_lists: array; +var rigid_particles_pos: array; @group(1) @binding(6) +var rigid_particle_node_linked_lists: array; +@group(1) @binding(7) var rigid_particle_needs_block: array>; // Disable this kernel on macos because of the underlying compareExchangeMap which is diff --git a/src/models/mod.rs b/src/models/mod.rs index c5cb6f9..0e8640b 100644 --- a/src/models/mod.rs +++ b/src/models/mod.rs @@ -17,15 +17,18 @@ pub struct GpuModels { } impl GpuModels { - pub fn from_particles(device: &Device, particles: &[Particle]) -> Self { + pub fn from_particles(device: &Device, particles: &[Particle], maximum_size: usize) -> Self { + let additional_unused_particles = maximum_size - particles.len(); let models: Vec<_> = particles.iter().map(|p| p.model).collect(); let plasticity: Vec<_> = particles .iter() .map(|p| p.plasticity.unwrap_or(DruckerPrager::new(-1.0, -1.0))) + .chain(vec![DruckerPrager::new(-1.0, -1.0); additional_unused_particles]) .collect(); let plastic_states: Vec<_> = particles .iter() .map(|_| DruckerPragerPlasticState::default()) + .chain(vec![DruckerPragerPlasticState::default(); additional_unused_particles]) .collect(); let phases: Vec<_> = particles .iter() @@ -35,6 +38,10 @@ impl GpuModels { max_stretch: -1.0, }) }) + .chain(vec![ParticlePhase { + phase: 0.0, + max_stretch: -1.0, + }; additional_unused_particles]) .collect(); Self { linear_elasticity: GpuVector::init(device, &models, BufferUsages::STORAGE), @@ -47,6 +54,10 @@ impl GpuModels { phases: GpuVector::init(device, &phases, BufferUsages::STORAGE), } } + + pub fn push() { + + } } fn lame_lambda_mu(young_modulus: f32, poisson_ratio: f32) -> (f32, f32) { diff --git a/src/pipeline.rs b/src/pipeline.rs index 87309c3..f0bced0 100644 --- a/src/pipeline.rs +++ b/src/pipeline.rs @@ -99,6 +99,7 @@ impl MpmData { device: &Device, params: SimulationParams, particles: &[Particle], + particles_max_count: usize, bodies: &RigidBodySet, colliders: &ColliderSet, cell_width: f32, @@ -119,6 +120,7 @@ impl MpmData { device, params, particles, + particles_max_count, bodies, colliders, coupling, @@ -131,6 +133,7 @@ impl MpmData { device: &Device, params: SimulationParams, particles: &[Particle], + particles_max_count: usize, bodies: &RigidBodySet, colliders: &ColliderSet, coupling: Vec, @@ -140,8 +143,8 @@ impl MpmData { let sampling_step = cell_width; // TODO: * 1.5 ? let bodies = GpuBodySet::from_rapier(device, bodies, colliders, &coupling); let sim_params = GpuSimulationParams::new(device, params); - let models = GpuModels::from_particles(device, particles); - let particles = GpuParticles::from_particles(device, particles); + let models = GpuModels::from_particles(device, particles, particles_max_count); + let particles = GpuParticles::from_particles(device, particles, particles_max_count); let rigid_particles = GpuRigidParticles::from_rapier(device, colliders, &bodies, &coupling, sampling_step); let grid = GpuGrid::with_capacity(device, grid_capacity, cell_width); @@ -325,6 +328,7 @@ mod test { gpu.device(), params, &cpu_particles, + 100_000, &RigidBodySet::default(), &ColliderSet::default(), cell_width, diff --git a/src/solver/g2p.rs b/src/solver/g2p.rs index eee6a7f..28ff140 100644 --- a/src/solver/g2p.rs +++ b/src/solver/g2p.rs @@ -43,6 +43,7 @@ impl WgG2P { 1, [ particles.positions.buffer(), + particles.current_size.buffer(), particles.dynamics.buffer(), particles.sorted_ids.buffer(), sim_params.params.buffer(), diff --git a/src/solver/p2g.wgsl b/src/solver/p2g.wgsl index b56250b..87554fe 100644 --- a/src/solver/p2g.wgsl +++ b/src/solver/p2g.wgsl @@ -8,13 +8,13 @@ #import wgrapier::body as Body; -@group(1) @binding(0) -var particles_pos: array; @group(1) @binding(1) -var particles_dyn: array; +var particles_pos: array; @group(1) @binding(2) -var nodes_linked_lists: array; +var particles_dyn: array; @group(1) @binding(3) +var nodes_linked_lists: array; +@group(1) @binding(4) var particle_node_linked_lists: array; @group(2) @binding(0) diff --git a/src/solver/particle2d.rs b/src/solver/particle2d.rs index 32b9e22..5ae49a8 100644 --- a/src/solver/particle2d.rs +++ b/src/solver/particle2d.rs @@ -4,7 +4,7 @@ use crate::solver::ParticlePhase; use encase::ShaderType; use nalgebra::{vector, Matrix2, Point2, Vector2}; use rapier::geometry::{ColliderSet, Polyline, Segment}; -use wgcore::tensor::GpuVector; +use wgcore::tensor::{GpuScalar, GpuVector}; use wgcore::Shader; use wgparry::shape::ShapeBuffers; use wgpu::{BufferUsages, Device}; @@ -148,36 +148,122 @@ pub struct GpuParticles { pub dynamics: GpuVector, pub sorted_ids: GpuVector, pub node_linked_lists: GpuVector, + /// The number of particles in use. + /// + /// As GPU buffers have to be re-created from the ground up to be resized, + /// it's an expensive operation, so we can allocate big buffers and keep track + /// of the used size. + pub current_size: GpuScalar, + /// As current_size is a uniform buffer, we can only change it from CPU, so it's easy to cache. + /// + /// If we don't cache it, we should retrieve it from current_size which is stored on GPU. + pub current_size_cached: u32, + /// Current actual size of the buffer on the GPU. + pub maximum_size: u32, } impl GpuParticles { pub fn is_empty(&self) -> bool { - self.positions.is_empty() + self.current_size_cached == 0 } pub fn len(&self) -> usize { - self.positions.len() as usize + self.current_size_cached as usize } - pub fn from_particles(device: &Device, particles: &[Particle]) -> Self { - let positions: Vec<_> = particles.iter().map(|p| p.position).collect(); - let dynamics: Vec<_> = particles.iter().map(|p| p.dynamics).collect(); - + /// Creates a new GPU particle buffer. + /// + /// `maximum_size` should be equal or greater than `particles.len()`. + /// Setting a greater value will allow to add more particles later. + pub fn from_particles(device: &Device, particles: &[Particle], maximum_size: usize) -> Self { + let particles_len = particles.len(); + assert!( + maximum_size >= particles_len, + "`maximum_size` should be equal or greater than `particles.len()`." + ); + let additional_unused_particles = maximum_size - particles_len; + let positions: Vec<_> = particles + .iter() + .map(|p| p.position) + .chain(vec![vector![0.0, 0.0]; additional_unused_particles]) + .collect::>(); + let dynamics: Vec<_> = particles + .iter() + .map(|p| p.dynamics) + .chain(vec![ + ParticleDynamics::with_density(0.0, 0.0); + additional_unused_particles + ]) + .collect(); Self { positions: GpuVector::init( device, &positions, - BufferUsages::STORAGE | BufferUsages::COPY_SRC, + BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST, + ), + dynamics: GpuVector::encase( + device, + &dynamics, + BufferUsages::STORAGE | BufferUsages::COPY_DST, + ), + sorted_ids: GpuVector::uninit( + device, + maximum_size as u32, + BufferUsages::STORAGE | BufferUsages::COPY_DST, ), - dynamics: GpuVector::encase(device, &dynamics, BufferUsages::STORAGE), - sorted_ids: GpuVector::uninit(device, particles.len() as u32, BufferUsages::STORAGE), node_linked_lists: GpuVector::uninit( device, - particles.len() as u32, - BufferUsages::STORAGE, + maximum_size as u32, + BufferUsages::STORAGE | BufferUsages::COPY_DST, ), + current_size: GpuScalar::init( + device, + particles_len as u32, + BufferUsages::UNIFORM | BufferUsages::COPY_DST, + ), + current_size_cached: particles_len as u32, + maximum_size: maximum_size as u32, } } + + // TODO: push multiple particles at once. + /// This is designed to work in conjunction with [`crate::models::GpuModels::push`] + pub fn push(&mut self, queue: &wgpu::Queue, particle: &Particle) { + if self.maximum_size == self.current_size_cached { + eprintln!("Particle buffer is full."); + return; + } + println!( + "ParticleDynamics min_size: {} bytes", + ParticleDynamics::min_size().get() + ); + println!("Cdf min_size: {} bytes", Cdf::min_size().get()); + println!("pushed a new particle!"); + // TODO: provide a helper in Tensor to push values? + + // Push position + let offset = self.current_size_cached as u64 * Vector2::::min_size().get(); + let position = particle.position; + // Turn value into &[u8] (raw bytes) + let bytes = bytemuck::bytes_of(&position); + queue.write_buffer(&self.positions.buffer(), offset, bytes); + + // Push dynamics. + let offset = self.current_size_cached as u64 * ParticleDynamics::min_size().get(); + let dynamics = particle.dynamics; + // Turn value into &[u8] (raw bytes) + let mut buffer = encase::StorageBuffer::new(Vec::::new()); + buffer.write(&dynamics).unwrap(); + let bytes = buffer.as_ref(); + queue.write_buffer(&self.dynamics.buffer(), offset, bytes); + + self.current_size_cached += 1; + queue.write_buffer( + &self.current_size.buffer(), + 0, + bytemuck::bytes_of(&self.current_size_cached), + ); + } } #[derive(Shader)] diff --git a/src/solver/particle_update.rs b/src/solver/particle_update.rs index b6a4801..f9b609c 100644 --- a/src/solver/particle_update.rs +++ b/src/solver/particle_update.rs @@ -54,6 +54,7 @@ impl WgParticleUpdate { .bind( 1, [ + particles.current_size.buffer(), particles.positions.buffer(), particles.dynamics.buffer(), models.linear_elasticity.buffer(), @@ -64,7 +65,7 @@ impl WgParticleUpdate { ], ) // .bind(2, [bodies.shapes().buffer(), bodies.poses().buffer()]) - .queue(particles.positions.len().div_ceil(64) as u32); + .queue(particles.len().div_ceil(64) as u32); } } diff --git a/src/solver/particle_update.wgsl b/src/solver/particle_update.wgsl index 972045a..1993e22 100644 --- a/src/solver/particle_update.wgsl +++ b/src/solver/particle_update.wgsl @@ -18,18 +18,22 @@ #import wgparry::cuboid as Cuboid; @group(1) @binding(0) -var particles_pos: array; +var num_particles: u32; @group(1) @binding(1) -var particles_dyn: array; +var particles_pos: array; @group(1) @binding(2) -var constitutive_model: array; +var particles_pos: array; @group(1) @binding(3) -var plasticity: array; +var particles_dyn: array; @group(1) @binding(4) -var plastic_state: array; +var constitutive_model: array; @group(1) @binding(5) -var phases: array; +var plasticity: array; @group(1) @binding(6) +var plastic_state: array; +@group(1) @binding(7) +var phases: array; +@group(1) @binding(8) var params: Params::SimulationParams; @group(2) @binding(0) diff --git a/src_testbed/lib.rs b/src_testbed/lib.rs index 6ce3b3c..5130fed 100644 --- a/src_testbed/lib.rs +++ b/src_testbed/lib.rs @@ -6,7 +6,7 @@ pub extern crate wgsparkl2d as wgsparkl; #[cfg(feature = "dim3")] pub extern crate wgsparkl3d as wgsparkl; -use bevy::render::renderer::RenderDevice; +use bevy::render::renderer::{RenderDevice, RenderQueue}; #[cfg(feature = "dim2")] pub use instancing2d as instancing; #[cfg(feature = "dim3")] @@ -114,7 +114,7 @@ pub struct PhysicsContext { pub struct Callbacks(pub Vec); pub type Callback = Box< - dyn FnMut(Option<&mut RenderContext>, &mut PhysicsContext, &Timestamps, &AppState) + dyn FnMut(Option<&mut RenderContext>, &mut PhysicsContext, &Timestamps, &AppState, &RenderQueue) + Send + Sync, >; diff --git a/src_testbed/prep_vertex_buffer.rs b/src_testbed/prep_vertex_buffer.rs index 44a3509..ba36a7a 100644 --- a/src_testbed/prep_vertex_buffer.rs +++ b/src_testbed/prep_vertex_buffer.rs @@ -94,13 +94,14 @@ impl WgPrepVertexBuffer { KernelInvocationBuilder::new(queue, &self.main) .bind0([ vertex_buffer, + particles.current_size.buffer(), particles.positions.buffer(), particles.dynamics.buffer(), grid.meta.buffer(), params.params.buffer(), config.buffer.buffer(), ]) - .queue(particles.positions.len().div_ceil(64) as u32); + .queue(particles.current_size_cached.div_ceil(64) as u32); if let Some(vertex_buffer) = vertex_buffer_rigid_particles { KernelInvocationBuilder::new(queue, &self.main_rigid_particles) diff --git a/src_testbed/prep_vertex_buffer2d.wgsl b/src_testbed/prep_vertex_buffer2d.wgsl index b36e035..40cace0 100644 --- a/src_testbed/prep_vertex_buffer2d.wgsl +++ b/src_testbed/prep_vertex_buffer2d.wgsl @@ -8,14 +8,16 @@ @group(0) @binding(0) var instances: array; @group(0) @binding(1) -var particles_pos: array; +var num_particles: u32; @group(0) @binding(2) -var particles_dyn: array; +var particles_pos: array; @group(0) @binding(3) -var grid: Grid::Grid; +var particles_dyn: array; @group(0) @binding(4) -var params: Params::SimulationParams; +var grid: Grid::Grid; @group(0) @binding(5) +var params: Params::SimulationParams; +@group(0) @binding(6) var config: RenderConfig; struct RenderConfig { @@ -42,7 +44,7 @@ fn main( ) { let particle_id = tid.x; - if particle_id < arrayLength(&instances) { + if particle_id < num_particles { let def_grad = particles_dyn[particle_id].def_grad; instances[particle_id].deformation = mat3x3(vec3(def_grad.x, 0.0), vec3(def_grad.y, 0.0), vec3(0.0, 0.0, 1.0)); instances[particle_id].position = vec3(particles_pos[particle_id].pt, 0.0); @@ -96,7 +98,7 @@ fn main_rigid_particles( ) { let particle_id = tid.x; - if particle_id < arrayLength(&instances) { + if particle_id < num_particles { instances[particle_id].deformation = mat3x3f( 0.4, 0.0, 0.0, 0.0, 0.4, 0.0, diff --git a/src_testbed/prep_vertex_buffer3d.wgsl b/src_testbed/prep_vertex_buffer3d.wgsl index 4dca049..5a1ece8 100644 --- a/src_testbed/prep_vertex_buffer3d.wgsl +++ b/src_testbed/prep_vertex_buffer3d.wgsl @@ -7,15 +7,19 @@ @group(0) @binding(0) var instances: array; +/// Stores the number of particles in use, because we cannot rely on `arrayLength(particles_pos)`, +/// as the array may be bigger to allow more efficient dynamic addition or removal of particles. @group(0) @binding(1) -var particles_pos: array; +var num_particles: u32; @group(0) @binding(2) -var particles_dyn: array; +var particles_pos: array; @group(0) @binding(3) -var grid: Grid::Grid; +var particles_dyn: array; @group(0) @binding(4) -var params: Params::SimulationParams; +var grid: Grid::Grid; @group(0) @binding(5) +var params: Params::SimulationParams; +@group(0) @binding(6) var config: RenderConfig; struct RenderConfig { @@ -43,7 +47,7 @@ fn main( ) { let particle_id = tid.x; - if particle_id < arrayLength(&instances) { + if particle_id < num_particles { let def_grad = particles_dyn[particle_id].def_grad; instances[particle_id].deformation = def_grad; instances[particle_id].position = particles_pos[particle_id].pt; diff --git a/src_testbed/step.rs b/src_testbed/step.rs index f252e8f..c29557a 100644 --- a/src_testbed/step.rs +++ b/src_testbed/step.rs @@ -20,6 +20,7 @@ pub struct TimestampChannel { } pub fn callbacks( + mut render_queue: ResMut, mut render: ResMut, mut physics: ResMut, app_state: ResMut, @@ -32,6 +33,7 @@ pub fn callbacks( &mut physics, timings.as_ref(), &app_state, + &*render_queue, ); } } From e6b16b708ae0eac129ca59994d83bdc6f16f8f46 Mon Sep 17 00:00:00 2001 From: Thierry Berger Date: Mon, 14 Apr 2025 18:06:56 +0200 Subject: [PATCH 2/2] working example to add particles --- examples2d/src/dynamic_add.rs | 36 ++++--- examples2d/src/elastic_cut2.rs | 1 + examples2d/src/elasticity2.rs | 1 + examples2d/src/sand2.rs | 1 + examples3d/src/banana3.rs | 3 +- examples3d/src/dynamic_add.rs | 9 +- examples3d/src/elastic_cut3.rs | 1 + examples3d/src/glb_to_point_cloud_color.rs | 1 + examples3d/src/heightfield3.rs | 1 + examples3d/src/sand3.rs | 2 - src/grid/grid.rs | 27 ++++-- src/grid/sort.rs | 4 +- src/grid/sort.wgsl | 8 +- src/grid/touch_particle_blocks2d.wgsl | 8 +- src/grid/touch_particle_blocks3d.wgsl | 8 +- src/models/mod.rs | 82 +++++++++++++--- src/pipeline.rs | 11 +++ src/solver/g2p.rs | 1 - src/solver/p2g.wgsl | 8 +- src/solver/particle2d.rs | 7 +- src/solver/particle3d.rs | 107 +++++++++++++++++++-- src/solver/particle_update.wgsl | 18 ++-- src_testbed/lib.rs | 16 +++ src_testbed/prep_vertex_buffer2d.wgsl | 4 +- src_testbed/prep_vertex_buffer3d.wgsl | 2 +- src_testbed/startup.rs | 23 +++++ src_testbed/step.rs | 2 +- 27 files changed, 305 insertions(+), 87 deletions(-) diff --git a/examples2d/src/dynamic_add.rs b/examples2d/src/dynamic_add.rs index ae54ad4..bbbc65f 100644 --- a/examples2d/src/dynamic_add.rs +++ b/examples2d/src/dynamic_add.rs @@ -23,17 +23,19 @@ pub fn dynamic_demo( let offset_y = 46.0; let cell_width = 0.2; let mut particles = vec![]; - let position = vector![0.5, 0.5] * cell_width / 2.0 + Vector2::y() * offset_y; + let position = + vector![-20f32, 0.0] + vector![0.5, 0.5] * cell_width / 2.0 + Vector2::y() * offset_y; let density = 1000.0; - let radius = cell_width / 4.0; - particles.push(Particle { + let radius = cell_width / 4f32; + let particle_proto = Particle { position, dynamics: ParticleDynamics::with_density(radius, density), model: ElasticCoefficients::from_young_modulus(10_000_000.0, 0.2), plasticity: Some(DruckerPrager::new(10_000_000.0, 0.2)), phase: None, color: None, - }); + }; + particles.push(particle_proto); if !app_state.restarting { app_state.num_substeps = 10; @@ -74,19 +76,20 @@ pub fn dynamic_demo( .colliders .insert_with_parent(co, rb_handle, &mut rapier_data.bodies); + let mut tick = 0; callbacks.0.push(Box::new( move |_, physics: &mut PhysicsContext, _, _, queue| { - physics.data.particles.push( - queue, - &Particle { - position, - dynamics: ParticleDynamics::with_density(radius, density), - model: ElasticCoefficients::from_young_modulus(2_000_000_000.0, 0.2), - plasticity: Some(DruckerPrager::new(2_000_000_000.0, 0.2)), - phase: None, - color: None, - }, - ) + tick += 1; + if tick % 2 != 0 { + // Do not add particles each frame, to avoid clumping particles together. + return; + } + for i in 0..500 { + let mut particle = particle_proto.clone(); + particle.position += vector![1.0, 0.0] * (i as f32) * cell_width * 1.1; + physics.push_particle(queue, &particle); + } + physics.reset_graphics = true; }, )); @@ -94,7 +97,7 @@ pub fn dynamic_demo( device, params, &particles, - 1_000, + 100_000, &rapier_data.bodies, &rapier_data.colliders, cell_width, @@ -104,5 +107,6 @@ pub fn dynamic_demo( data, rapier_data, particles, + reset_graphics: false, } } diff --git a/examples2d/src/elastic_cut2.rs b/examples2d/src/elastic_cut2.rs index 97748e9..3db3d4c 100644 --- a/examples2d/src/elastic_cut2.rs +++ b/examples2d/src/elastic_cut2.rs @@ -113,5 +113,6 @@ pub fn elastic_cut_demo( data, rapier_data, particles, + reset_graphics: false, } } diff --git a/examples2d/src/elasticity2.rs b/examples2d/src/elasticity2.rs index e35b96c..b16b4b1 100644 --- a/examples2d/src/elasticity2.rs +++ b/examples2d/src/elasticity2.rs @@ -95,5 +95,6 @@ pub fn elastic_demo( data, rapier_data, particles, + reset_graphics: false, } } diff --git a/examples2d/src/sand2.rs b/examples2d/src/sand2.rs index 5c1edbe..674bdf3 100644 --- a/examples2d/src/sand2.rs +++ b/examples2d/src/sand2.rs @@ -173,5 +173,6 @@ pub fn sand_demo( data, rapier_data, particles, + reset_graphics: false, } } diff --git a/examples3d/src/banana3.rs b/examples3d/src/banana3.rs index d9f7592..5509807 100644 --- a/examples3d/src/banana3.rs +++ b/examples3d/src/banana3.rs @@ -114,6 +114,7 @@ pub fn demo( device.wgpu_device(), params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, 1f32 / default_scene::SAMPLE_PER_UNIT, @@ -134,7 +135,7 @@ pub fn demo( fn move_knife_function(body_handle: RigidBodyHandle) -> Callback { let mut ticks = 0; Box::new( - move |_render, physics: &mut PhysicsContext, _timestamps, _app_state: &AppState| { + move |_render, physics: &mut PhysicsContext, _timestamps, _app_state: &AppState, _| { let t = ticks as f32 * physics.rapier_data.integration_parameters.dt; ticks += 1; diff --git a/examples3d/src/dynamic_add.rs b/examples3d/src/dynamic_add.rs index a851e04..ad7083c 100644 --- a/examples3d/src/dynamic_add.rs +++ b/examples3d/src/dynamic_add.rs @@ -2,7 +2,6 @@ use wgsparkl_testbed3d::{wgsparkl, Callbacks, RapierData}; use bevy::render::renderer::RenderDevice; use nalgebra::vector; -use rapier3d::prelude::{ColliderBuilder, RigidBodyBuilder}; use wgsparkl::models::DruckerPrager; use wgsparkl::{ models::ElasticCoefficients, @@ -16,7 +15,7 @@ use crate::utils::default_scene; pub fn dynamic_demo( device: RenderDevice, app_state: &mut AppState, - _callbacks: &mut Callbacks, + callbacks: &mut Callbacks, ) -> PhysicsContext { let mut rapier_data = RapierData::default(); let device = device.wgpu_device(); @@ -32,14 +31,16 @@ pub fn dynamic_demo( }; default_scene::spawn_ground_and_walls(&mut rapier_data); + let cell_width = 1.0; + let position = vector![0.0, 10f32, 0.0]; callbacks.0.push(Box::new( move |_, physics: &mut PhysicsContext, _, _, queue| { physics.data.particles.push( queue, &Particle { position, - dynamics: ParticleDynamics::with_density(radius, density), + dynamics: ParticleDynamics::with_density(cell_width / 4.0, 2700.0), model: ElasticCoefficients::from_young_modulus(2_000_000_000.0, 0.2), plasticity: Some(DruckerPrager::new(2_000_000_000.0, 0.2)), phase: None, @@ -50,11 +51,11 @@ pub fn dynamic_demo( )); let particles = vec![]; - let cell_width = 1.0; let data = MpmData::new( device, params, &particles, + 1_000, &rapier_data.bodies, &rapier_data.colliders, cell_width, diff --git a/examples3d/src/elastic_cut3.rs b/examples3d/src/elastic_cut3.rs index ccb0876..075a850 100644 --- a/examples3d/src/elastic_cut3.rs +++ b/examples3d/src/elastic_cut3.rs @@ -87,6 +87,7 @@ pub fn elastic_cut_demo( device, params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, cell_width, diff --git a/examples3d/src/glb_to_point_cloud_color.rs b/examples3d/src/glb_to_point_cloud_color.rs index 5e03b8a..d6ff3e1 100644 --- a/examples3d/src/glb_to_point_cloud_color.rs +++ b/examples3d/src/glb_to_point_cloud_color.rs @@ -53,6 +53,7 @@ pub fn elastic_color_model_demo( device.wgpu_device(), params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, 1f32 / default_scene::SAMPLE_PER_UNIT, diff --git a/examples3d/src/heightfield3.rs b/examples3d/src/heightfield3.rs index 152d1ca..b5c8101 100644 --- a/examples3d/src/heightfield3.rs +++ b/examples3d/src/heightfield3.rs @@ -72,6 +72,7 @@ pub fn heightfield_demo( device, params, &particles, + particles.len(), &rapier_data.bodies, &rapier_data.colliders, cell_width, diff --git a/examples3d/src/sand3.rs b/examples3d/src/sand3.rs index 31decbe..f01ac80 100644 --- a/examples3d/src/sand3.rs +++ b/examples3d/src/sand3.rs @@ -1,5 +1,3 @@ -use std::ptr::dangling; - use wgsparkl_testbed3d::{wgsparkl, Callbacks, RapierData}; use bevy::render::renderer::RenderDevice; diff --git a/src/grid/grid.rs b/src/grid/grid.rs index 61769f5..0f5d6f5 100644 --- a/src/grid/grid.rs +++ b/src/grid/grid.rs @@ -69,8 +69,14 @@ impl WgGrid { (grid.debug.buffer(), 8), ], ) - .bind(1, [particles.positions.buffer()]) - .queue((particles.len() as u32).div_ceil(GRID_WORKGROUP_SIZE)); + .bind( + 1, + [ + particles.current_size.buffer(), + particles.positions.buffer(), + ], + ) + .queue(particles.current_size_cached.div_ceil(GRID_WORKGROUP_SIZE)); // Ensure blocks exist wherever we have rigid particles that might affect // other blocks. This is done in two passes: @@ -113,7 +119,6 @@ impl WgGrid { ) .queue((rigid_particles.len() as u32).div_ceil(GRID_WORKGROUP_SIZE)); - // TODO: handle grid buffer resizing sparse_grid_has_the_correct_size = true; } @@ -149,7 +154,13 @@ impl WgGrid { grid.active_blocks.buffer(), ], ) - .bind(1, [particles.positions.buffer()]) + .bind( + 1, + [ + particles.current_size.buffer(), + particles.positions.buffer(), + ], + ) .queue(n_groups); KernelInvocationBuilder::new(queue, &sort_module.copy_particles_len_to_scan_value) @@ -157,7 +168,7 @@ impl WgGrid { 0, [(grid.meta.buffer(), 0), (grid.active_blocks.buffer(), 2)], ) - .bind_at(1, [(grid.scan_values.buffer(), 1)]) + .bind_at(1, [(grid.scan_values.buffer(), 2)]) .queue_indirect(n_block_groups.clone()); prefix_sum_module.queue(queue, prefix_sum, &grid.scan_values); @@ -167,7 +178,7 @@ impl WgGrid { 0, [(grid.meta.buffer(), 0), (grid.active_blocks.buffer(), 2)], ) - .bind_at(1, [(grid.scan_values.buffer(), 1)]) + .bind_at(1, [(grid.scan_values.buffer(), 2)]) .queue_indirect(n_block_groups.clone()); // Reset here so the linked list heads get reset before `finalize_particles_sort` which @@ -196,6 +207,7 @@ impl WgGrid { .bind( 1, [ + particles.current_size.buffer(), particles.positions.buffer(), grid.scan_values.buffer(), particles.sorted_ids.buffer(), @@ -370,7 +382,8 @@ mod test { } } - let particles = GpuParticles::from_particles(gpu.device(), &cpu_particles); + let particles = + GpuParticles::from_particles(gpu.device(), &cpu_particles, cpu_particles.len()); let grid = GpuGrid::with_capacity(gpu.device(), 100_000, cell_width); let mut prefix_sum = PrefixSumWorkspace::with_capacity(gpu.device(), 100_000); let mut queue = KernelInvocationQueue::new(gpu.device()); diff --git a/src/grid/sort.rs b/src/grid/sort.rs index 3d371de..f58b449 100644 --- a/src/grid/sort.rs +++ b/src/grid/sort.rs @@ -45,8 +45,8 @@ impl WgSort { .bind_at( 1, [ - (particles.sample_points.buffer(), 4), - (particles.node_linked_lists.buffer(), 5), + (particles.sample_points.buffer(), 5), + (particles.node_linked_lists.buffer(), 6), ], ) .queue(n_groups); diff --git a/src/grid/sort.wgsl b/src/grid/sort.wgsl index 327c6d4..22c6067 100644 --- a/src/grid/sort.wgsl +++ b/src/grid/sort.wgsl @@ -5,7 +5,7 @@ @group(1) @binding(0) -var num_particles u32; +var num_particles: u32; @group(1) @binding(1) var particles_pos: array; @group(1) @binding(2) @@ -28,7 +28,7 @@ var rigid_particle_needs_block: array>; @compute @workgroup_size(Grid::GRID_WORKGROUP_SIZE, 1, 1) fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; - if id < arrayLength(&particles_pos) { + if id < num_particles { let particle = particles_pos[id]; var blocks = Grid::blocks_associated_to_point(particle.pt); for (var i = 0u; i < Grid::NUM_ASSOC_BLOCKS; i += 1u) { @@ -91,7 +91,7 @@ fn mark_rigid_particles_needing_block(@builtin(global_invocation_id) invocation_ @compute @workgroup_size(Grid::GRID_WORKGROUP_SIZE, 1, 1) fn update_block_particle_count(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; - if id < arrayLength(&particles_pos) { + if id < num_particles { let particle = particles_pos[id]; let block_id = Grid::block_associated_to_point(particle.pt); let active_block_id = Grid::find_block_header_id(block_id); @@ -119,7 +119,7 @@ fn copy_scan_values_to_first_particles(@builtin(global_invocation_id) invocation @compute @workgroup_size(Grid::GRID_WORKGROUP_SIZE, 1, 1) fn finalize_particles_sort(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; - if id < arrayLength(&particles_pos) { + if id < num_particles { let particle = particles_pos[id]; let block_id = Grid::block_associated_to_point(particle.pt); diff --git a/src/grid/touch_particle_blocks2d.wgsl b/src/grid/touch_particle_blocks2d.wgsl index 90c167e..34e1c0d 100644 --- a/src/grid/touch_particle_blocks2d.wgsl +++ b/src/grid/touch_particle_blocks2d.wgsl @@ -231,14 +231,16 @@ struct Position { } @group(1) @binding(0) +var num_particles: u32; +@group(1) @binding(1) var particles_pos: array; -@group(1) @binding(6) +@group(1) @binding(7) var rigid_particle_needs_block: array; @compute @workgroup_size(GRID_WORKGROUP_SIZE, 1, 1) fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; - if id < arrayLength(&particles_pos) { + if id < num_particles { let particle = particles_pos[id]; var blocks = blocks_associated_to_point(particle.pt); for (var i = 0u; i < NUM_ASSOC_BLOCKS; i += 1u) { @@ -250,7 +252,7 @@ fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3 @compute @workgroup_size(GRID_WORKGROUP_SIZE, 1, 1) fn touch_rigid_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; - if id < arrayLength(&particles_pos) { + if id < num_particles { let entry_id = id / 32u; let entry_bit = 1u << (id % 32u); let needs_block = (rigid_particle_needs_block[entry_id] & entry_bit) != 0; diff --git a/src/grid/touch_particle_blocks3d.wgsl b/src/grid/touch_particle_blocks3d.wgsl index af91629..6275258 100644 --- a/src/grid/touch_particle_blocks3d.wgsl +++ b/src/grid/touch_particle_blocks3d.wgsl @@ -240,14 +240,16 @@ struct Position { } @group(1) @binding(0) +var num_particles: u32; +@group(1) @binding(1) var particles_pos: array; -@group(1) @binding(6) +@group(1) @binding(7) var rigid_particle_needs_block: array; @compute @workgroup_size(GRID_WORKGROUP_SIZE, 1, 1) fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; - if id < arrayLength(&particles_pos) { + if id < num_particles { let particle = particles_pos[id]; // PERF: we should look at the local cell coordinates of the point // in the block and only touch adjacent blocks if the @@ -264,7 +266,7 @@ fn touch_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3 @compute @workgroup_size(GRID_WORKGROUP_SIZE, 1, 1) fn touch_rigid_particle_blocks(@builtin(global_invocation_id) invocation_id: vec3) { let id = invocation_id.x; - if id < arrayLength(&particles_pos) { + if id < num_particles { let entry_id = id / 32u; let entry_bit = 1u << (id % 32u); let needs_block = (rigid_particle_needs_block[entry_id] & entry_bit) != 0; diff --git a/src/models/mod.rs b/src/models/mod.rs index 0e8640b..226be4c 100644 --- a/src/models/mod.rs +++ b/src/models/mod.rs @@ -18,17 +18,33 @@ pub struct GpuModels { impl GpuModels { pub fn from_particles(device: &Device, particles: &[Particle], maximum_size: usize) -> Self { - let additional_unused_particles = maximum_size - particles.len(); - let models: Vec<_> = particles.iter().map(|p| p.model).collect(); + let additional_unused_particles = maximum_size - particles.len(); + let models: Vec<_> = particles + .iter() + .map(|p| p.model) + .chain(vec![ + ElasticCoefficients::from_young_modulus( + 2_000_000_000.0, + 0.2 + ); + additional_unused_particles + ]) + .collect(); let plasticity: Vec<_> = particles .iter() .map(|p| p.plasticity.unwrap_or(DruckerPrager::new(-1.0, -1.0))) - .chain(vec![DruckerPrager::new(-1.0, -1.0); additional_unused_particles]) + .chain(vec![ + DruckerPrager::new(-1.0, -1.0); + additional_unused_particles + ]) .collect(); let plastic_states: Vec<_> = particles .iter() .map(|_| DruckerPragerPlasticState::default()) - .chain(vec![DruckerPragerPlasticState::default(); additional_unused_particles]) + .chain(vec![ + DruckerPragerPlasticState::default(); + additional_unused_particles + ]) .collect(); let phases: Vec<_> = particles .iter() @@ -38,25 +54,65 @@ impl GpuModels { max_stretch: -1.0, }) }) - .chain(vec![ParticlePhase { - phase: 0.0, - max_stretch: -1.0, - }; additional_unused_particles]) + .chain(vec![ + ParticlePhase { + phase: 0.0, + max_stretch: -1.0, + }; + additional_unused_particles + ]) .collect(); Self { - linear_elasticity: GpuVector::init(device, &models, BufferUsages::STORAGE), - drucker_prager_plasticity: GpuVector::init(device, &plasticity, BufferUsages::STORAGE), + linear_elasticity: GpuVector::init( + device, + &models, + BufferUsages::STORAGE | BufferUsages::COPY_DST, + ), + drucker_prager_plasticity: GpuVector::init( + device, + &plasticity, + BufferUsages::STORAGE | BufferUsages::COPY_DST, + ), drucker_prager_plastic_state: GpuVector::init( device, &plastic_states, - BufferUsages::STORAGE, + BufferUsages::STORAGE | BufferUsages::COPY_DST, + ), + phases: GpuVector::init( + device, + &phases, + BufferUsages::STORAGE | BufferUsages::COPY_DST, ), - phases: GpuVector::init(device, &phases, BufferUsages::STORAGE), } } - pub fn push() { + pub fn push(&mut self, queue: &wgpu::Queue, particle: &Particle, current_size: u32) { + let model = particle.model; + let plasticity = particle + .plasticity + .unwrap_or(DruckerPrager::new(-1.0, -1.0)); + let plastic_state = DruckerPragerPlasticState::default(); + let phase = particle.phase.unwrap_or(ParticlePhase { + phase: 0.0, + max_stretch: -1.0, + }); + // Push model + let offset = current_size as u64 * size_of::() as u64; + let bytes = bytemuck::bytes_of(&model); + queue.write_buffer(&self.linear_elasticity.buffer(), offset, bytes); + // Push plasticity + let offset = current_size as u64 * size_of::() as u64; + let bytes = bytemuck::bytes_of(&plasticity); + queue.write_buffer(&self.drucker_prager_plasticity.buffer(), offset, bytes); + // Push plastic_state + let offset = current_size as u64 * size_of::() as u64; + let bytes = bytemuck::bytes_of(&plastic_state); + queue.write_buffer(&self.drucker_prager_plastic_state.buffer(), offset, bytes); + // Push phase + let offset = current_size as u64 * size_of::() as u64; + let bytes = bytemuck::bytes_of(&phase); + queue.write_buffer(&self.phases.buffer(), offset, bytes); } } diff --git a/src/pipeline.rs b/src/pipeline.rs index f0bced0..8223618 100644 --- a/src/pipeline.rs +++ b/src/pipeline.rs @@ -173,6 +173,17 @@ impl MpmData { pub fn coupling(&self) -> &[BodyCouplingEntry] { &self.coupling } + + pub fn push_particle(&mut self, queue: &wgpu::Queue, particle: &Particle) -> Result { + if self.particles.maximum_size == self.particles.current_size_cached { + eprintln!("Particle buffers are full."); + return Err(()); + } + self.models + .push(queue, particle, self.particles.current_size_cached); + self.particles.push(queue, particle); + return Ok(1); + } } impl MpmPipeline { diff --git a/src/solver/g2p.rs b/src/solver/g2p.rs index 28ff140..eee6a7f 100644 --- a/src/solver/g2p.rs +++ b/src/solver/g2p.rs @@ -43,7 +43,6 @@ impl WgG2P { 1, [ particles.positions.buffer(), - particles.current_size.buffer(), particles.dynamics.buffer(), particles.sorted_ids.buffer(), sim_params.params.buffer(), diff --git a/src/solver/p2g.wgsl b/src/solver/p2g.wgsl index 87554fe..b56250b 100644 --- a/src/solver/p2g.wgsl +++ b/src/solver/p2g.wgsl @@ -8,13 +8,13 @@ #import wgrapier::body as Body; -@group(1) @binding(1) +@group(1) @binding(0) var particles_pos: array; -@group(1) @binding(2) +@group(1) @binding(1) var particles_dyn: array; -@group(1) @binding(3) +@group(1) @binding(2) var nodes_linked_lists: array; -@group(1) @binding(4) +@group(1) @binding(3) var particle_node_linked_lists: array; @group(2) @binding(0) diff --git a/src/solver/particle2d.rs b/src/solver/particle2d.rs index 5ae49a8..0d13aa4 100644 --- a/src/solver/particle2d.rs +++ b/src/solver/particle2d.rs @@ -227,12 +227,9 @@ impl GpuParticles { } // TODO: push multiple particles at once. + // TODO: merge with particle3d.rs /// This is designed to work in conjunction with [`crate::models::GpuModels::push`] pub fn push(&mut self, queue: &wgpu::Queue, particle: &Particle) { - if self.maximum_size == self.current_size_cached { - eprintln!("Particle buffer is full."); - return; - } println!( "ParticleDynamics min_size: {} bytes", ParticleDynamics::min_size().get() @@ -244,14 +241,12 @@ impl GpuParticles { // Push position let offset = self.current_size_cached as u64 * Vector2::::min_size().get(); let position = particle.position; - // Turn value into &[u8] (raw bytes) let bytes = bytemuck::bytes_of(&position); queue.write_buffer(&self.positions.buffer(), offset, bytes); // Push dynamics. let offset = self.current_size_cached as u64 * ParticleDynamics::min_size().get(); let dynamics = particle.dynamics; - // Turn value into &[u8] (raw bytes) let mut buffer = encase::StorageBuffer::new(Vec::::new()); buffer.write(&dynamics).unwrap(); let bytes = buffer.as_ref(); diff --git a/src/solver/particle3d.rs b/src/solver/particle3d.rs index e737f40..f5faeb1 100644 --- a/src/solver/particle3d.rs +++ b/src/solver/particle3d.rs @@ -6,7 +6,7 @@ use nalgebra::{vector, Matrix3, Point3, Vector3, Vector4}; use rapier::geometry::{Segment, Triangle}; use rapier::prelude::{ColliderSet, TriMesh}; use std::collections::HashSet; -use wgcore::tensor::GpuVector; +use wgcore::tensor::{GpuScalar, GpuVector}; use wgcore::Shader; use wgparry::shape::ShapeBuffers; use wgpu::{BufferUsages, Device}; @@ -179,6 +179,18 @@ pub struct GpuParticles { pub dynamics: GpuVector, pub sorted_ids: GpuVector, pub node_linked_lists: GpuVector, + /// The number of particles in use. + /// + /// As GPU buffers have to be re-created from the ground up to be resized, + /// it's an expensive operation, so we can allocate big buffers and keep track + /// of the used size. + pub current_size: GpuScalar, + /// As current_size is a uniform buffer, we can only change it from CPU, so it's easy to cache. + /// + /// If we don't cache it, we should retrieve it from current_size which is stored on GPU. + pub current_size_cached: u32, + /// Current actual size of the buffer on the GPU. + pub maximum_size: u32, } impl GpuParticles { @@ -190,24 +202,103 @@ impl GpuParticles { self.positions.len() as usize } - pub fn from_particles(device: &Device, particles: &[Particle]) -> Self { - let positions: Vec<_> = particles.iter().map(|p| p.position.push(0.0)).collect(); - let dynamics: Vec<_> = particles.iter().map(|p| p.dynamics).collect(); + /// Creates a new GPU particle buffer. + /// + /// `maximum_size` should be equal or greater than `particles.len()`. + /// Setting a greater value will allow to add more particles later. + pub fn from_particles(device: &Device, particles: &[Particle], maximum_size: usize) -> Self { + let particles_len = particles.len(); + assert!( + maximum_size >= particles_len, + "`maximum_size` should be equal or greater than `particles.len()`." + ); + let additional_unused_particles = maximum_size - particles_len; + let positions: Vec<_> = particles + .iter() + .map(|p| p.position.push(0.0)) + .chain(vec![ + vector![0.0, 0.0, 0.0, 0.0]; + additional_unused_particles + ]) + .collect::>(); + let dynamics: Vec<_> = particles + .iter() + .map(|p| p.dynamics) + .chain(vec![ + ParticleDynamics::with_density(0.0, 0.0); + additional_unused_particles + ]) + .collect(); Self { positions: GpuVector::init( device, &positions, - BufferUsages::STORAGE | BufferUsages::COPY_SRC, + BufferUsages::STORAGE | BufferUsages::COPY_SRC | BufferUsages::COPY_DST, + ), + dynamics: GpuVector::encase( + device, + &dynamics, + BufferUsages::STORAGE | BufferUsages::COPY_DST, + ), + sorted_ids: GpuVector::uninit( + device, + particles.len() as u32, + BufferUsages::STORAGE | BufferUsages::COPY_DST, ), - dynamics: GpuVector::encase(device, &dynamics, BufferUsages::STORAGE), - sorted_ids: GpuVector::uninit(device, particles.len() as u32, BufferUsages::STORAGE), node_linked_lists: GpuVector::uninit( device, particles.len() as u32, - BufferUsages::STORAGE, + BufferUsages::STORAGE | BufferUsages::COPY_DST, + ), + current_size: GpuScalar::init( + device, + particles_len as u32, + BufferUsages::UNIFORM | BufferUsages::COPY_DST, ), + current_size_cached: particles_len as u32, + maximum_size: maximum_size as u32, + } + } + + // TODO: push multiple particles at once. + // TODO: merge with particle2d.rs + /// This is designed to work in conjunction with [`crate::models::GpuModels::push`] + pub fn push(&mut self, queue: &wgpu::Queue, particle: &Particle) { + if self.maximum_size == self.current_size_cached { + eprintln!("Particle buffer is full."); + return; } + println!( + "ParticleDynamics min_size: {} bytes", + ParticleDynamics::min_size().get() + ); + println!("Cdf min_size: {} bytes", Cdf::min_size().get()); + println!("pushed a new particle!"); + // TODO: provide a helper in Tensor to push values? + + // Push position + let offset = self.current_size_cached as u64 * Vector3::::min_size().get(); + let position = particle.position; + // Turn value into &[u8] (raw bytes) + let bytes = bytemuck::bytes_of(&position); + queue.write_buffer(&self.positions.buffer(), offset, bytes); + + // Push dynamics. + let offset = self.current_size_cached as u64 * ParticleDynamics::min_size().get(); + let dynamics = particle.dynamics; + // Turn value into &[u8] (raw bytes) + let mut buffer = encase::StorageBuffer::new(Vec::::new()); + buffer.write(&dynamics).unwrap(); + let bytes = buffer.as_ref(); + queue.write_buffer(&self.dynamics.buffer(), offset, bytes); + + self.current_size_cached += 1; + queue.write_buffer( + &self.current_size.buffer(), + 0, + bytemuck::bytes_of(&self.current_size_cached), + ); } } diff --git a/src/solver/particle_update.wgsl b/src/solver/particle_update.wgsl index 1993e22..918c7ef 100644 --- a/src/solver/particle_update.wgsl +++ b/src/solver/particle_update.wgsl @@ -18,22 +18,20 @@ #import wgparry::cuboid as Cuboid; @group(1) @binding(0) -var num_particles: u32; +var num_particles: u32; @group(1) @binding(1) -var particles_pos: array; -@group(1) @binding(2) var particles_pos: array; -@group(1) @binding(3) +@group(1) @binding(2) var particles_dyn: array; -@group(1) @binding(4) +@group(1) @binding(3) var constitutive_model: array; -@group(1) @binding(5) +@group(1) @binding(4) var plasticity: array; -@group(1) @binding(6) +@group(1) @binding(5) var plastic_state: array; -@group(1) @binding(7) +@group(1) @binding(6) var phases: array; -@group(1) @binding(8) +@group(1) @binding(7) var params: Params::SimulationParams; @group(2) @binding(0) @@ -52,7 +50,7 @@ fn main( ) { let particle_id = gid.x; - if particle_id >= arrayLength(&particles_pos) { + if particle_id >= num_particles { return; } diff --git a/src_testbed/lib.rs b/src_testbed/lib.rs index 5130fed..56596b4 100644 --- a/src_testbed/lib.rs +++ b/src_testbed/lib.rs @@ -54,6 +54,10 @@ pub fn init_testbed(app: &mut App) { .init_resource::() .init_resource::() .add_systems(Startup, startup::setup_app) + .add_systems( + Update, + startup::update_graphics.run_if(resource_exists::), + ) .add_systems( Update, ( @@ -108,6 +112,18 @@ pub struct PhysicsContext { pub data: MpmData, pub rapier_data: RapierData, pub particles: Vec, + /// Flag to reset the graphics pipeline. + // TODO: pass the commands to the callback, so we can schedule a graphics reset, + // or a more fine-grained update of the graphics pipeline. + pub reset_graphics: bool, +} + +impl PhysicsContext { + pub fn push_particle(&mut self, queue: &RenderQueue, particle: &Particle) { + if self.data.push_particle(queue, particle).is_ok() { + self.particles.push(particle.clone()); + } + } } #[derive(Resource, Default)] diff --git a/src_testbed/prep_vertex_buffer2d.wgsl b/src_testbed/prep_vertex_buffer2d.wgsl index 40cace0..4cb0709 100644 --- a/src_testbed/prep_vertex_buffer2d.wgsl +++ b/src_testbed/prep_vertex_buffer2d.wgsl @@ -7,8 +7,10 @@ @group(0) @binding(0) var instances: array; +/// Stores the number of particles in use, because we cannot rely on `arrayLength(particles_pos)`, +/// as the array may be bigger to allow more efficient dynamic addition or removal of particles. @group(0) @binding(1) -var num_particles: u32; +var num_particles: u32; @group(0) @binding(2) var particles_pos: array; @group(0) @binding(3) diff --git a/src_testbed/prep_vertex_buffer3d.wgsl b/src_testbed/prep_vertex_buffer3d.wgsl index 5a1ece8..8148c3b 100644 --- a/src_testbed/prep_vertex_buffer3d.wgsl +++ b/src_testbed/prep_vertex_buffer3d.wgsl @@ -10,7 +10,7 @@ var instances: array; /// Stores the number of particles in use, because we cannot rely on `arrayLength(particles_pos)`, /// as the array may be bigger to allow more efficient dynamic addition or removal of particles. @group(0) @binding(1) -var num_particles: u32; +var num_particles: u32; @group(0) @binding(2) var particles_pos: array; @group(0) @binding(3) diff --git a/src_testbed/startup.rs b/src_testbed/startup.rs index 1f5894c..db88da0 100644 --- a/src_testbed/startup.rs +++ b/src_testbed/startup.rs @@ -284,3 +284,26 @@ pub fn setup_graphics( pub fn setup_app_state(mut callbacks: ResMut) { callbacks.0.clear(); } + +pub fn update_graphics( + mut commands: Commands, + physics_context: Res, + app_state: Res, + device: Res, + physics: Res, + mut rigid_render: ResMut, + mut meshes: ResMut>, + mut materials: ResMut>, + to_clear: Query>, +) { + if physics_context.reset_graphics { + setup_particles_graphics( + &mut commands, + &device, + &app_state, + &physics, + &mut meshes, + &to_clear, + ); + } +} diff --git a/src_testbed/step.rs b/src_testbed/step.rs index c29557a..c3e20ca 100644 --- a/src_testbed/step.rs +++ b/src_testbed/step.rs @@ -20,7 +20,7 @@ pub struct TimestampChannel { } pub fn callbacks( - mut render_queue: ResMut, + render_queue: ResMut, mut render: ResMut, mut physics: ResMut, app_state: ResMut,