From f35711f21c38a3b087f84cb9427814d6b7a1ec44 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?charlotte=20=F0=9F=8C=B8?= Date: Wed, 5 Aug 2026 22:28:27 -0700 Subject: [PATCH 1/2] Particles. --- Cargo.lock | 1 + Cargo.toml | 13 +- crates/processing_ffi/Cargo.toml | 1 + crates/processing_ffi/src/lib.rs | 540 ++++++++++++++++-- .../examples/particles_emit.py | 1 - .../examples/particles_emit_gpu.py | 16 +- .../examples/particles_lifecycle.py | 26 +- .../examples/particles_noise.py | 2 +- .../examples/particles_scatter_volume.py | 56 ++ .../examples/particles_stress.py | 4 +- crates/processing_pyo3/src/compute.rs | 10 +- crates/processing_pyo3/src/lib.rs | 27 +- crates/processing_pyo3/src/particles.rs | 201 ++++++- crates/processing_render/src/compute.rs | 254 +++++++- .../src/geometry/attribute.rs | 100 +++- crates/processing_render/src/graphics.rs | 68 ++- crates/processing_render/src/lib.rs | 390 ++++--------- .../processing_render/src/material/custom.rs | 11 +- .../processing_render/src/particles/emit.rs | 180 ++++++ .../src/particles/kernels/age.wgsl | 20 + .../src/particles/kernels/attr_combine.wgsl | 34 ++ .../src/particles/kernels/attr_linear.wgsl | 15 + .../src/particles/kernels/attr_lookup1d.wgsl | 27 + .../src/particles/kernels/attr_lookup2d.wgsl | 26 + .../src/particles/kernels/attr_mix.wgsl | 23 + .../src/particles/kernels/attract.wgsl | 48 ++ .../src/particles/kernels/bounds_box.wgsl | 67 +++ .../src/particles/kernels/bounds_sphere.wgsl | 60 ++ .../src/particles/kernels/drag.wgsl | 29 + .../src/particles/kernels/field.wgsl | 45 ++ .../src/particles/kernels/flock.wgsl | 127 ++++ .../src/particles/kernels/force.wgsl | 23 + .../src/particles/kernels/impulse.wgsl | 53 ++ .../src/particles/kernels/integrate.wgsl | 19 + .../src/particles/kernels/mod.rs | 311 +++++++++- .../src/particles/kernels/noise.wgsl | 31 +- .../src/particles/kernels/orient.wgsl | 115 ++++ .../particles/kernels/scatter_surface.wgsl | 99 ++++ .../src/particles/kernels/scatter_volume.wgsl | 119 ++++ .../src/particles/kernels/vortex.wgsl | 57 ++ .../src/particles/material.rs | 55 +- crates/processing_render/src/particles/mod.rs | 255 ++++++++- .../processing_render/src/particles/pack.rs | 40 +- .../processing_render/src/particles/pack.wgsl | 10 +- .../src/particles/particles.wgsl | 17 + .../src/particles/scatter.rs | 197 +++++++ crates/processing_render/src/render/mod.rs | 3 +- .../src/render/primitive/shape3d.rs | 1 - crates/processing_render/src/shader_value.rs | 17 +- docs/particles.md | 23 - examples/compute_readback.rs | 17 +- examples/particles_emit.rs | 1 - examples/particles_emit_gpu.rs | 20 +- examples/particles_lifecycle.rs | 27 +- examples/particles_oriented.rs | 1 - examples/particles_scatter.rs | 121 ++++ examples/particles_scatter_volume.rs | 113 ++++ examples/particles_stress.rs | 6 +- examples/particles_text_whirl.rs | 417 ++++++++++++++ 59 files changed, 4034 insertions(+), 556 deletions(-) create mode 100644 crates/processing_pyo3/examples/particles_scatter_volume.py create mode 100644 crates/processing_render/src/particles/emit.rs create mode 100644 crates/processing_render/src/particles/kernels/age.wgsl create mode 100644 crates/processing_render/src/particles/kernels/attr_combine.wgsl create mode 100644 crates/processing_render/src/particles/kernels/attr_linear.wgsl create mode 100644 crates/processing_render/src/particles/kernels/attr_lookup1d.wgsl create mode 100644 crates/processing_render/src/particles/kernels/attr_lookup2d.wgsl create mode 100644 crates/processing_render/src/particles/kernels/attr_mix.wgsl create mode 100644 crates/processing_render/src/particles/kernels/attract.wgsl create mode 100644 crates/processing_render/src/particles/kernels/bounds_box.wgsl create mode 100644 crates/processing_render/src/particles/kernels/bounds_sphere.wgsl create mode 100644 crates/processing_render/src/particles/kernels/drag.wgsl create mode 100644 crates/processing_render/src/particles/kernels/field.wgsl create mode 100644 crates/processing_render/src/particles/kernels/flock.wgsl create mode 100644 crates/processing_render/src/particles/kernels/force.wgsl create mode 100644 crates/processing_render/src/particles/kernels/impulse.wgsl create mode 100644 crates/processing_render/src/particles/kernels/integrate.wgsl create mode 100644 crates/processing_render/src/particles/kernels/orient.wgsl create mode 100644 crates/processing_render/src/particles/kernels/scatter_surface.wgsl create mode 100644 crates/processing_render/src/particles/kernels/scatter_volume.wgsl create mode 100644 crates/processing_render/src/particles/kernels/vortex.wgsl create mode 100644 crates/processing_render/src/particles/scatter.rs delete mode 100644 docs/particles.md create mode 100644 examples/particles_scatter.rs create mode 100644 examples/particles_scatter_volume.rs create mode 100644 examples/particles_text_whirl.rs diff --git a/Cargo.lock b/Cargo.lock index 2704f809..9b293a1b 100644 --- a/Cargo.lock +++ b/Cargo.lock @@ -6124,6 +6124,7 @@ version = "0.0.5" dependencies = [ "bevy", "cbindgen", + "half", "processing", ] diff --git a/Cargo.toml b/Cargo.toml index 4e550b4a..2a2354db 100644 --- a/Cargo.toml +++ b/Cargo.toml @@ -256,6 +256,18 @@ path = "examples/particles_emit_gpu.rs" name = "particles_stress" path = "examples/particles_stress.rs" +[[example]] +name = "particles_scatter" +path = "examples/particles_scatter.rs" + +[[example]] +name = "particles_scatter_volume" +path = "examples/particles_scatter_volume.rs" + +[[example]] +name = "particles_text_whirl" +path = "examples/particles_text_whirl.rs" + [[example]] name = "text" path = "examples/text.rs" @@ -267,7 +279,6 @@ path = "examples/text_3d.rs" [[example]] name = "filter" path = "examples/filter.rs" - [profile.wasm-release] inherits = "release" opt-level = "z" diff --git a/crates/processing_ffi/Cargo.toml b/crates/processing_ffi/Cargo.toml index e034c774..16d41364 100644 --- a/crates/processing_ffi/Cargo.toml +++ b/crates/processing_ffi/Cargo.toml @@ -20,6 +20,7 @@ cuda = ["processing/cuda"] [dependencies] processing = { workspace = true } bevy = { workspace = true } +half = "2.7" [build-dependencies] cbindgen = "0.29" diff --git a/crates/processing_ffi/src/lib.rs b/crates/processing_ffi/src/lib.rs index b81adb1f..f520db75 100644 --- a/crates/processing_ffi/src/lib.rs +++ b/crates/processing_ffi/src/lib.rs @@ -1536,6 +1536,35 @@ pub unsafe extern "C" fn processing_image_create( .unwrap_or(0) } +/// # Safety +/// - `init` has been called. +/// - `floats` is valid for `floats_len` f32 reads. +/// - Called from the same thread as `init`. +#[unsafe(no_mangle)] +pub unsafe extern "C" fn processing_image_create_hdr( + width: u32, + height: u32, + floats: *const f32, + floats_len: usize, +) -> u64 { + error::clear_error(); + let src = unsafe { std::slice::from_raw_parts(floats, floats_len) }; + error::check(|| { + let mut packed = Vec::with_capacity(src.len() * 2); + for &f in src { + packed.extend_from_slice(&half::f16::from_f32(f).to_le_bytes()); + } + let size = Extent3d { + width, + height, + depth_or_array_layers: 1, + }; + image_create(size, packed, TextureFormat::Rgba16Float) + }) + .map(|entity| entity.to_bits()) + .unwrap_or(0) +} + /// Load an image from a file path. /// /// # Safety @@ -2412,6 +2441,70 @@ pub extern "C" fn processing_geometry_attribute_uv() -> u64 { geometry_attribute_uv().to_bits() } +#[unsafe(no_mangle)] +pub extern "C" fn processing_geometry_attribute_rotation() -> u64 { + geometry_attribute_rotation().to_bits() +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_geometry_attribute_scale() -> u64 { + geometry_attribute_scale().to_bits() +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_geometry_attribute_life() -> u64 { + geometry_attribute_life().to_bits() +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_geometry_attribute_velocity() -> u64 { + geometry_attribute_velocity().to_bits() +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_geometry_attribute_age() -> u64 { + geometry_attribute_age().to_bits() +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_geometry_attribute_format(attr_id: u64) -> u8 { + error::clear_error(); + error::check(|| { + let (_name, fmt) = geometry_attribute_info(Entity::from_bits(attr_id))?; + Ok(match fmt { + geometry::AttributeFormat::Float => 1, + geometry::AttributeFormat::Float2 => 2, + geometry::AttributeFormat::Float3 => 3, + geometry::AttributeFormat::Float4 => 4, + }) + }) + .unwrap_or(0) +} + +/// # Safety +/// - `out` is valid for `out_cap` byte writes (may be null when `out_cap == 0` +/// for a length query). +#[unsafe(no_mangle)] +pub unsafe extern "C" fn processing_geometry_attribute_name( + attr_id: u64, + out: *mut u8, + out_cap: u64, +) -> u64 { + error::clear_error(); + let Some((name, _)) = error::check(|| geometry_attribute_info(Entity::from_bits(attr_id))) + else { + return 0; + }; + let name_bytes = name.as_bytes(); + let name_len = name_bytes.len(); + if out_cap > 0 && !out.is_null() { + let copy_len = name_len.min((out_cap - 1) as usize); + unsafe { std::ptr::copy_nonoverlapping(name_bytes.as_ptr(), out, copy_len) }; + unsafe { *out.add(copy_len) = 0 }; + } + name_len as u64 +} + #[unsafe(no_mangle)] pub extern "C" fn processing_geometry_attribute_float(geo_id: u64, attr_id: u64, v: f32) { error::clear_error(); @@ -2748,10 +2841,18 @@ pub extern "C" fn processing_material_create_pbr() -> u64 { .unwrap_or(0) } -/// Set float value for `name` field on Material. +#[unsafe(no_mangle)] +pub extern "C" fn processing_material_create_custom(shader_id: u64) -> u64 { + error::clear_error(); + error::check(|| material_create_custom(Entity::from_bits(shader_id))) + .map(|e| e.to_bits()) + .unwrap_or(0) +} + +/// Set a float field on a material. /// /// # Safety -/// - `name` must be non-null +/// - `name` is a valid null-terminated C string. #[unsafe(no_mangle)] pub unsafe extern "C" fn processing_material_set_float( mat_id: u64, @@ -2769,10 +2870,10 @@ pub unsafe extern "C" fn processing_material_set_float( }); } -/// Set float4 value for `name` field on Material. +/// Set a float4 field on a material. /// /// # Safety -/// - `name` must be non-null +/// - `name` is a valid null-terminated C string. #[unsafe(no_mangle)] pub unsafe extern "C" fn processing_material_set_float4( mat_id: u64, @@ -2807,8 +2908,10 @@ pub extern "C" fn processing_material(window_id: u64, mat_id: u64) { error::check(|| graphics_record_command(window_entity, DrawCommand::Material(mat_entity))); } +/// Create a shader from WGSL source. +/// /// # Safety -/// - `source` must be non-null +/// - `source` is a valid null-terminated C string. #[unsafe(no_mangle)] pub unsafe extern "C" fn processing_shader_create(source: *const std::ffi::c_char) -> u64 { error::clear_error(); @@ -2821,7 +2924,7 @@ pub unsafe extern "C" fn processing_shader_create(source: *const std::ffi::c_cha } /// # Safety -/// - `path` must be non-null +/// - `path` is a valid null-terminated C string. #[unsafe(no_mangle)] pub unsafe extern "C" fn processing_shader_load(path: *const std::ffi::c_char) -> u64 { error::clear_error(); @@ -2848,7 +2951,7 @@ pub extern "C" fn processing_buffer_create(size: u64) -> u64 { } /// # Safety -/// - `data` must point to `len` valid bytes +/// - `data` is valid for `len` byte reads. #[unsafe(no_mangle)] pub unsafe extern "C" fn processing_buffer_create_with_data(data: *const u8, len: u64) -> u64 { error::clear_error(); @@ -2859,7 +2962,7 @@ pub unsafe extern "C" fn processing_buffer_create_with_data(data: *const u8, len } /// # Safety -/// - `data` must point to `len` valid bytes +/// - `data` is valid for `len` byte reads. #[unsafe(no_mangle)] pub unsafe extern "C" fn processing_buffer_write(buf_id: u64, data: *const u8, len: u64) { error::clear_error(); @@ -2867,16 +2970,19 @@ pub unsafe extern "C" fn processing_buffer_write(buf_id: u64, data: *const u8, l error::check(|| buffer_write(Entity::from_bits(buf_id), bytes)); } -/// returns the byte length of a buffer, or 0 if not found. +/// Returns the byte length of a buffer, or 0 if not found (error is set). #[unsafe(no_mangle)] pub extern "C" fn processing_buffer_size(buf_id: u64) -> u64 { error::clear_error(); error::check(|| buffer_size(Entity::from_bits(buf_id))).unwrap_or(0) } +/// Read buffer contents into `out`. Returns the buffer's byte length; +/// `out` is only written when it fits in `out_len`. Pass `out_len == 0` for a +/// size query. +/// /// # Safety -/// - `out` must be valid for writes of `out_len` bytes (may be null if -/// `out_len == 0`, in which case this acts as a size query). +/// - `out` is valid for `out_len` byte writes (may be null when `out_len == 0`). #[unsafe(no_mangle)] pub unsafe extern "C" fn processing_buffer_read(buf_id: u64, out: *mut u8, out_len: u64) -> u64 { error::clear_error(); @@ -3078,94 +3184,446 @@ pub extern "C" fn processing_compute_destroy(compute_id: u64) { error::check(|| compute_destroy(Entity::from_bits(compute_id))); } -/// Create a filter from a shader entity. +/// # Safety +/// - `attr_ids` is valid for `attr_count` u64 reads. #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_create(shader_id: u64) -> u64 { +pub unsafe extern "C" fn processing_particles_create( + capacity: u32, + attr_ids: *const u64, + attr_count: u32, +) -> u64 { error::clear_error(); - error::check(|| filter_create(Entity::from_bits(shader_id))) + let attrs = if attr_count > 0 && !attr_ids.is_null() { + unsafe { std::slice::from_raw_parts(attr_ids, attr_count as usize) } + .iter() + .map(|&id| Entity::from_bits(id)) + .collect() + } else { + vec![] + }; + error::check(|| particles_create(capacity, attrs)) .map(|e| e.to_bits()) .unwrap_or(0) } +/// # Safety +/// - `attr_ids` is valid for `attr_count` u64 reads. #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_invert() -> u64 { +pub unsafe extern "C" fn processing_particles_create_from_geometry( + geo_id: u64, + attr_ids: *const u64, + attr_count: u32, +) -> u64 { error::clear_error(); - error::check(filter_invert) + let attrs = if attr_count > 0 && !attr_ids.is_null() { + unsafe { std::slice::from_raw_parts(attr_ids, attr_count as usize) } + .iter() + .map(|&id| Entity::from_bits(id)) + .collect() + } else { + vec![] + }; + error::check(|| particles_create_from_geometry(Entity::from_bits(geo_id), attrs)) .map(|e| e.to_bits()) .unwrap_or(0) } #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_gray() -> u64 { +pub extern "C" fn processing_particles_destroy(particles_id: u64) { error::clear_error(); - error::check(filter_gray).map(|e| e.to_bits()).unwrap_or(0) + error::check(|| particles_destroy(Entity::from_bits(particles_id))); } #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_threshold() -> u64 { +pub extern "C" fn processing_particles_capacity(particles_id: u64) -> u32 { error::clear_error(); - error::check(filter_threshold) - .map(|e| e.to_bits()) - .unwrap_or(0) + error::check(|| particles_capacity(Entity::from_bits(particles_id))).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_buffer(particles_id: u64, attr_id: u64) -> u64 { + error::clear_error(); + error::check(|| { + particles_buffer(Entity::from_bits(particles_id), Entity::from_bits(attr_id)) + }) + .flatten() + .map(|e| e.to_bits()) + .unwrap_or(0) +} + +/// # Safety +/// - `attr_ids` is valid for `attr_count` u64 reads. +/// - `attr_byte_lengths` is valid for `attr_count` u64 reads. +/// - `data` is valid for `sum(attr_byte_lengths)` byte reads. +#[unsafe(no_mangle)] +pub unsafe extern "C" fn processing_particles_emit( + particles_id: u64, + n: u32, + attr_ids: *const u64, + data: *const u8, + attr_byte_lengths: *const u64, + attr_count: u32, +) { + error::clear_error(); + error::check(|| { + if attr_count == 0 { + return particles_emit(Entity::from_bits(particles_id), n, vec![]); + } + let ids = unsafe { std::slice::from_raw_parts(attr_ids, attr_count as usize) }; + let lens = unsafe { std::slice::from_raw_parts(attr_byte_lengths, attr_count as usize) }; + let mut offset: usize = 0; + let mut attribute_data = Vec::with_capacity(attr_count as usize); + for i in 0..attr_count as usize { + let len = lens[i] as usize; + let bytes = unsafe { std::slice::from_raw_parts(data.add(offset), len) }.to_vec(); + attribute_data.push((Entity::from_bits(ids[i]), bytes)); + offset += len; + } + particles_emit(Entity::from_bits(particles_id), n, attribute_data) + }); +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_emit_gpu( + particles_id: u64, + n: u32, + compute_id: u64, +) { + error::clear_error(); + error::check(|| { + particles_emit_gpu( + Entity::from_bits(particles_id), + n, + Entity::from_bits(compute_id), + ) + }); +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_noise() -> u64 { + error::clear_error(); + error::check(particles_kernel_noise).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_transform() -> u64 { + error::clear_error(); + error::check(particles_kernel_transform).map(|e| e.to_bits()).unwrap_or(0) } #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_posterize() -> u64 { +pub extern "C" fn processing_particles_kernel_attract() -> u64 { error::clear_error(); - error::check(filter_posterize) + error::check(particles_kernel_attract).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_drag() -> u64 { + error::clear_error(); + error::check(particles_kernel_drag).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_vortex() -> u64 { + error::clear_error(); + error::check(particles_kernel_vortex).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_attribute_add( + particles_id: u64, + attribute_id: u64, +) -> i32 { + error::clear_error(); + error::check(|| { + particles_attribute_add( + Entity::from_bits(particles_id), + Entity::from_bits(attribute_id), + None, + ) + }) + .map(|_| 0) + .unwrap_or(-1) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_force() -> u64 { + error::clear_error(); + error::check(particles_kernel_force).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_integrate() -> u64 { + error::clear_error(); + error::check(particles_kernel_integrate).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_age() -> u64 { + error::clear_error(); + error::check(particles_kernel_age).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_bounds_sphere() -> u64 { + error::clear_error(); + error::check(particles_kernel_bounds_sphere).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_bounds_box() -> u64 { + error::clear_error(); + error::check(particles_kernel_bounds_box).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_bounds_geometry(geometry_entity: u64) -> u64 { + error::clear_error(); + error::check(|| particles_kernel_bounds_geometry(Entity::from_bits(geometry_entity))) .map(|e| e.to_bits()) .unwrap_or(0) } #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_blur() -> u64 { +pub extern "C" fn processing_particles_kernel_impulse() -> u64 { + error::clear_error(); + error::check(particles_kernel_impulse).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_flock() -> u64 { + error::clear_error(); + error::check(particles_kernel_flock).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_orient() -> u64 { + error::clear_error(); + error::check(particles_kernel_orient).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_field() -> u64 { + error::clear_error(); + error::check(particles_kernel_field).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_attr_linear() -> u64 { + error::clear_error(); + error::check(particles_kernel_attr_linear).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_attr_combine() -> u64 { + error::clear_error(); + error::check(particles_kernel_attr_combine).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_attr_mix() -> u64 { + error::clear_error(); + error::check(particles_kernel_attr_mix).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_attr_lookup1d() -> u64 { + error::clear_error(); + error::check(particles_kernel_attr_lookup1d).map(|e| e.to_bits()).unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_kernel_attr_lookup2d() -> u64 { error::clear_error(); - error::check(filter_blur).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_attr_lookup2d).map(|e| e.to_bits()).unwrap_or(0) } #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_opaque() -> u64 { +pub extern "C" fn processing_particles_scatter_create(geometry_id: u64) -> u64 { error::clear_error(); - error::check(filter_opaque) + error::check(|| particles_scatter_create(Entity::from_bits(geometry_id))) .map(|e| e.to_bits()) .unwrap_or(0) } #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_erode() -> u64 { +pub extern "C" fn processing_particles_scatter_volume_create(geometry_id: u64) -> u64 { + error::clear_error(); + error::check(|| particles_scatter_volume_create(Entity::from_bits(geometry_id))) + .map(|e| e.to_bits()) + .unwrap_or(0) +} + +/// # Safety +/// - `path` is a valid null-terminated C string. +#[unsafe(no_mangle)] +pub unsafe extern "C" fn processing_gltf_load( + graphics_id: u64, + path: *const std::ffi::c_char, +) -> u64 { + error::clear_error(); + error::check(|| { + let path = unsafe { cstr_to_str(path) }?; + gltf_load(Entity::from_bits(graphics_id), path) + }) + .map(|e| e.to_bits()) + .unwrap_or(0) +} + +/// # Safety +/// - `name` is a valid null-terminated C string. +#[unsafe(no_mangle)] +pub unsafe extern "C" fn processing_gltf_geometry( + gltf_id: u64, + name: *const std::ffi::c_char, +) -> u64 { + error::clear_error(); + error::check(|| { + let name = unsafe { cstr_to_str(name) }?; + gltf_geometry(Entity::from_bits(gltf_id), name) + }) + .map(|e| e.to_bits()) + .unwrap_or(0) +} + +/// # Safety +/// - `name` is a valid null-terminated C string. +#[unsafe(no_mangle)] +pub unsafe extern "C" fn processing_gltf_material( + gltf_id: u64, + name: *const std::ffi::c_char, +) -> u64 { + error::clear_error(); + error::check(|| { + let name = unsafe { cstr_to_str(name) }?; + gltf_material(Entity::from_bits(gltf_id), name) + }) + .map(|e| e.to_bits()) + .unwrap_or(0) +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_gltf_camera(gltf_id: u64, index: u32) { error::clear_error(); - error::check(filter_erode).map(|e| e.to_bits()).unwrap_or(0) + error::check(|| gltf_camera(Entity::from_bits(gltf_id), index as usize)); } #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_dilate() -> u64 { +pub extern "C" fn processing_gltf_light(gltf_id: u64, index: u32) -> u64 { error::clear_error(); - error::check(filter_dilate) + error::check(|| gltf_light(Entity::from_bits(gltf_id), index as usize)) .map(|e| e.to_bits()) .unwrap_or(0) } -/// Set the number of fullscreen passes a filter runs. #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_set_passes(filter_id: u64, passes: u32) { +pub extern "C" fn processing_particles_apply(particles_id: u64, compute_id: u64) { + error::clear_error(); + error::check(|| { + particles_apply( + Entity::from_bits(particles_id), + Entity::from_bits(compute_id), + ) + }); +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_particles_draw( + graphics_id: u64, + particles_id: u64, + geometry_id: u64, +) { + error::clear_error(); + let graphics_entity = Entity::from_bits(graphics_id); + error::check(|| { + graphics_record_command( + graphics_entity, + DrawCommand::Particles { + particles: Entity::from_bits(particles_id), + geometry: Entity::from_bits(geometry_id), + }, + ) + }); +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_fill_buffer(graphics_id: u64, buffer_id: u64) { + error::clear_error(); + let graphics_entity = Entity::from_bits(graphics_id); + error::check(|| { + graphics_record_command( + graphics_entity, + DrawCommand::FillBuffer(Entity::from_bits(buffer_id)), + ) + }); +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_material_set_albedo_color( + mat_id: u64, + r: f32, + g: f32, + b: f32, + a: f32, +) { + error::clear_error(); + error::check(|| material_set_albedo_color(Entity::from_bits(mat_id), [r, g, b, a])); +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_material_set_albedo_buffer(mat_id: u64, buffer_id: u64) { error::clear_error(); - error::check(|| filter_set_passes(Entity::from_bits(filter_id), passes)); + error::check(|| { + material_set_albedo_buffer(Entity::from_bits(mat_id), Entity::from_bits(buffer_id)) + }); } -/// Apply a filter to a graphics canvas. #[unsafe(no_mangle)] -pub extern "C" fn processing_graphics_apply_filter(graphics_id: u64, filter_id: u64) { +pub extern "C" fn processing_material_set_emissive_buffer(mat_id: u64, buffer_id: u64) { error::clear_error(); error::check(|| { - graphics_apply_filter(Entity::from_bits(graphics_id), Entity::from_bits(filter_id)) + material_set_emissive_buffer(Entity::from_bits(mat_id), Entity::from_bits(buffer_id)) }); } -/// Destroy a filter entity. +/// # Safety +/// - `out_x`, `out_y`, `out_z` are each valid for one f32 write. +#[unsafe(no_mangle)] +pub unsafe extern "C" fn processing_graphics_world_from_screen( + graphics_id: u64, + sx: f32, + sy: f32, + depth: f32, + out_x: *mut f32, + out_y: *mut f32, + out_z: *mut f32, +) { + error::clear_error(); + if let Some(world) = error::check(|| { + graphics_world_from_screen(Entity::from_bits(graphics_id), sx, sy, depth) + }) { + unsafe { + *out_x = world.x; + *out_y = world.y; + *out_z = world.z; + } + } +} + +#[unsafe(no_mangle)] +pub extern "C" fn processing_graphics_set_bloom( + graphics_id: u64, + intensity: f32, + threshold: f32, +) { + error::clear_error(); + error::check(|| graphics_set_bloom(Entity::from_bits(graphics_id), intensity, threshold)); +} + #[unsafe(no_mangle)] -pub extern "C" fn processing_filter_destroy(filter_id: u64) { +pub extern "C" fn processing_graphics_remove_bloom(graphics_id: u64) { error::clear_error(); - error::check(|| filter_destroy(Entity::from_bits(filter_id))); + error::check(|| graphics_remove_bloom(Entity::from_bits(graphics_id))); } // Mouse buttons diff --git a/crates/processing_pyo3/examples/particles_emit.py b/crates/processing_pyo3/examples/particles_emit.py index cea86f02..23dad914 100644 --- a/crates/processing_pyo3/examples/particles_emit.py +++ b/crates/processing_pyo3/examples/particles_emit.py @@ -21,7 +21,6 @@ def setup(): attributes=[Attribute.position(), Attribute.color()], ) - # park unemitted slots far off-screen until the ring buffer fills. pos_buf = p.buffer(Attribute.position()) pos_buf.write([1.0e6] * (capacity * 3)) diff --git a/crates/processing_pyo3/examples/particles_emit_gpu.py b/crates/processing_pyo3/examples/particles_emit_gpu.py index 5224fcf1..943a991b 100644 --- a/crates/processing_pyo3/examples/particles_emit_gpu.py +++ b/crates/processing_pyo3/examples/particles_emit_gpu.py @@ -25,7 +25,7 @@ @group(0) @binding(2) var color: array; @group(0) @binding(3) var scale: array; @group(0) @binding(4) var age: array; -@group(0) @binding(5) var dead: array; +@group(0) @binding(5) var life: array; @group(0) @binding(6) var spawn: Spawn; @group(0) @binding(7) var emit_range: vec4; @@ -78,7 +78,7 @@ scale[slot * 3u + 2u] = 1.0; age[slot] = 0.0; - dead[slot] = 0.0; + life[slot] = 1.0; } """ @@ -94,7 +94,7 @@ @group(0) @binding(1) var velocity: array; @group(0) @binding(2) var scale: array; @group(0) @binding(3) var age: array; -@group(0) @binding(4) var dead: array; +@group(0) @binding(4) var life: array; @group(0) @binding(5) var params: Params; @compute @workgroup_size(64) @@ -102,7 +102,7 @@ let i = gid.x; let count = arrayLength(&age); if i >= count { return; } - if dead[i] != 0.0 { return; } + if life[i] <= 0.0 { return; } age[i] = age[i] + params.dt; @@ -118,7 +118,7 @@ scale[i * 3u + 1u] = s; scale[i * 3u + 2u] = s; - if age[i] > params.ttl { dead[i] = 1.0; } + if age[i] > params.ttl { life[i] = 0.0; } } """ @@ -142,16 +142,12 @@ def setup(): Attribute.position(), Attribute.color(), Attribute.scale(), - Attribute.dead(), + Attribute.life(), velocity_attr, age_attr, ], ) - # park unemitted slots until the spawn kernel fills them. - dead_buf = p.buffer(Attribute.dead()) - dead_buf.write([1.0] * CAPACITY) - color_buf = p.buffer(Attribute.color()) mat = Material.pbr(albedo=color_buf) diff --git a/crates/processing_pyo3/examples/particles_lifecycle.py b/crates/processing_pyo3/examples/particles_lifecycle.py index 2ee30ba5..eae7df01 100644 --- a/crates/processing_pyo3/examples/particles_lifecycle.py +++ b/crates/processing_pyo3/examples/particles_lifecycle.py @@ -8,7 +8,7 @@ position_attr = None color_attr = None scale_attr = None -dead_attr = None +life_attr = None age_attr = None frame = 0 @@ -18,7 +18,7 @@ AGING_SHADER = """ @group(0) @binding(0) var age: array; -@group(0) @binding(1) var dead: array; +@group(0) @binding(1) var life: array; @group(0) @binding(2) var position: array; @group(0) @binding(3) var scale: array; @group(0) @binding(4) var params: vec4; // x = dt, y = ttl @@ -33,21 +33,21 @@ let dt = params.x; let ttl = params.y; - if dead[i] != 0.0 { + if life[i] <= 0.0 { return; } age[i] = age[i] + dt; position[i * 3u + 1u] = position[i * 3u + 1u] - dt * 1.5; - let life = clamp(1.0 - age[i] / ttl, 0.0, 1.0); - let s = life * life; + let remaining = clamp(1.0 - age[i] / ttl, 0.0, 1.0); + let s = remaining * remaining; scale[i * 3u + 0u] = s; scale[i * 3u + 1u] = s; scale[i * 3u + 2u] = s; if age[i] > ttl { - dead[i] = 1.0; + life[i] = 0.0; } } """ @@ -55,7 +55,7 @@ def setup(): global p, sphere, mat, aging - global position_attr, color_attr, scale_attr, dead_attr, age_attr + global position_attr, color_attr, scale_attr, life_attr, age_attr size(900, 700) mode_3d() @@ -66,18 +66,13 @@ def setup(): position_attr = Attribute.position() color_attr = Attribute.color() scale_attr = Attribute.scale() - dead_attr = Attribute.dead() + life_attr = Attribute.life() age_attr = Attribute("age", AttributeFormat.Float) p = Particles( capacity=capacity, - attributes=[position_attr, color_attr, scale_attr, dead_attr, age_attr], + attributes=[position_attr, color_attr, scale_attr, life_attr, age_attr], ) - - # park unemitted slots until the spawn loop fills them. - dead_buf = p.buffer(dead_attr) - dead_buf.write([1.0] * capacity) - color_buf = p.buffer(color_attr) mat = Material.unlit(albedo=color_buf) aging = Compute(Shader(AGING_SHADER)) @@ -105,6 +100,7 @@ def draw(): colors.extend([c.r, c.g, c.b, 1.0]) zeros = [0.0] * BURST + ones = [1.0] * BURST ones_scale = [1.0] * (BURST * 3) p.emit( BURST, @@ -112,7 +108,7 @@ def draw(): color=colors, scale=ones_scale, age=zeros, - dead=zeros, + life=ones, ) aging.set(params=[DT, TTL, 0.0, 0.0]) diff --git a/crates/processing_pyo3/examples/particles_noise.py b/crates/processing_pyo3/examples/particles_noise.py index c32bbb13..c1e7dd8b 100644 --- a/crates/processing_pyo3/examples/particles_noise.py +++ b/crates/processing_pyo3/examples/particles_noise.py @@ -31,7 +31,7 @@ def setup(): particle = Geometry.sphere(0.18, 10, 8) mat = Material.pbr(albedo=color_buf) - noise = kernel_noise() + noise = Particles.noise() def draw(): diff --git a/crates/processing_pyo3/examples/particles_scatter_volume.py b/crates/processing_pyo3/examples/particles_scatter_volume.py new file mode 100644 index 00000000..30ca4490 --- /dev/null +++ b/crates/processing_pyo3/examples/particles_scatter_volume.py @@ -0,0 +1,56 @@ +from mewnala import * + +CAPACITY = 30_000 +BURST = 250 + +p = None +particle = None +mat = None +scatter = None +decay = None + + +def setup(): + global p, particle, mat, scatter, decay + + size(900, 700) + mode_3d() + + camera_position(0.0, 100.0, 400.0) + camera_look_at(0.0, 80.0, 0.0) + orbit_camera() + + gltf = load_gltf("gltf/Duck.glb") + duck = gltf.geometry("LOD3spShape") + scatter = Particles.scatter_volume(duck) + + particle = Geometry.sphere(0.15, 4, 3) + + age_attr = Attribute("age", AttributeFormat.Float) + p = Particles( + capacity=CAPACITY, + attributes=[ + Attribute.position(), + Attribute.scale(), + Attribute.life(), + age_attr, + ], + ) + mat = Material.unlit(albedo=[1.0, 1.0, 1.0, 1.0]) + + decay = Particles.attr_linear() + decay.set(op=p.buffer(Attribute.scale()), scale=0.985, offset=0.0) + + +def draw(): + background(8, 8, 13) + use_material(mat) + particles(p, particle) + + seed = (int(elapsed_time * 1000.0) ^ 0xC0FFEE) & 0xFFFFFFFF + scatter.set(seed=seed) + p.emit_gpu(BURST, scatter) + p.apply(decay) + + +run() diff --git a/crates/processing_pyo3/examples/particles_stress.py b/crates/processing_pyo3/examples/particles_stress.py index 1f8c5176..a0de597f 100644 --- a/crates/processing_pyo3/examples/particles_stress.py +++ b/crates/processing_pyo3/examples/particles_stress.py @@ -29,7 +29,7 @@ def setup(): attributes=[Attribute.position(), Attribute.uv(), Attribute.color()], ) - p.apply(kernel_noise(), scale=1.0 / SPACING, strength=SPACING * 0.6) + p.apply(Particles.noise(), scale=1.0 / SPACING, strength=SPACING * 0.6) color_buf = p.buffer(Attribute.color()) color_buf.write([ @@ -40,7 +40,7 @@ def setup(): fill(color_buf) cube = Geometry.box(0.35, 0.35, 0.35) - spin = kernel_transform() + spin = Particles.transform() def draw(): diff --git a/crates/processing_pyo3/src/compute.rs b/crates/processing_pyo3/src/compute.rs index 97124b24..efad0a69 100644 --- a/crates/processing_pyo3/src/compute.rs +++ b/crates/processing_pyo3/src/compute.rs @@ -16,13 +16,10 @@ pub struct Buffer { pub(crate) entity: Entity, element_type: Option, size: u64, - /// `true` for borrowed wrappers (e.g. `Particles.buffer()`) where the - /// underlying entity belongs elsewhere; `Drop` skips destroy in that case. borrowed: bool, } impl Buffer { - /// borrowed wrapper: `Drop` will not destroy the underlying entity. pub(crate) fn from_entity(entity: Entity, element_type: Option) -> Self { let size = buffer_size(entity).unwrap_or(0); Self { @@ -258,8 +255,11 @@ fn shader_value_to_py<'py>(py: Python<'py>, sv: &ShaderValue) -> PyResult list(py, v), ShaderValue::Int4(v) => list(py, v), ShaderValue::Mat4(v) => list(py, v), - ShaderValue::Texture(_) | ShaderValue::Buffer(_) => Err(PyRuntimeError::new_err( - "cannot convert Texture/Buffer to Python value", + ShaderValue::Texture(_) + | ShaderValue::Buffer(_) + | ShaderValue::MeshAttribute(..) + | ShaderValue::MeshIndex(_) => Err(PyRuntimeError::new_err( + "cannot convert Texture/Buffer/Mesh* to Python value", )), } } diff --git a/crates/processing_pyo3/src/lib.rs b/crates/processing_pyo3/src/lib.rs index 3ce155f5..ba08e6f2 100644 --- a/crates/processing_pyo3/src/lib.rs +++ b/crates/processing_pyo3/src/lib.rs @@ -334,6 +334,12 @@ mod mewnala { #[pymodule_export] use super::Compute; #[pymodule_export] + use super::particles::Attribute; + #[pymodule_export] + use super::particles::AttributeFormat; + #[pymodule_export] + use super::particles::Particles; + #[pymodule_export] use super::Font; #[pymodule_export] use super::Geometry; @@ -369,12 +375,6 @@ mod mewnala { #[pymodule_export] use super::monitor::Monitor; #[pymodule_export] - use super::particles::Attribute; - #[pymodule_export] - use super::particles::AttributeFormat; - #[pymodule_export] - use super::particles::Particles; - #[pymodule_export] use super::surface::Surface; #[pymodule_init] @@ -426,7 +426,8 @@ mod mewnala { } } - // color constructors live at module level: a `color` submodule conflicted with `color()` + // top-level so `from mewnala import *` exposes hsva/srgb/etc. directly; + // a `color` submodule would clash with the `color()` function #[pyfunction] fn color_hex(s: &str) -> PyResult { @@ -888,7 +889,7 @@ mod mewnala { Ok(()) }); - // tear the app down here while the TLS is still alive; the eager + // tear down the app while the thread-local is still alive; the eager // TLS destructor aborts inside a Bevy resource drop let _ = ::processing::exit(0); @@ -1097,16 +1098,6 @@ mod mewnala { ) } - #[pyfunction] - fn kernel_noise() -> PyResult { - super::particles::kernel_noise() - } - - #[pyfunction] - fn kernel_transform() -> PyResult { - super::particles::kernel_transform() - } - #[pyfunction(name = "color")] #[pyo3(pass_module, signature = (*args))] fn create_color( diff --git a/crates/processing_pyo3/src/particles.rs b/crates/processing_pyo3/src/particles.rs index 936992ac..afb7e699 100644 --- a/crates/processing_pyo3/src/particles.rs +++ b/crates/processing_pyo3/src/particles.rs @@ -46,8 +46,6 @@ impl AttributeFormat { } } -/// named typed attribute. use the `position()`/`color()`/etc. classmethods for -/// builtins or `Attribute(name, format)` for custom ones. #[pyclass(unsendable, frozen, hash, eq, from_py_object)] #[derive(Clone, PartialEq, Eq, Hash)] pub struct Attribute { @@ -100,9 +98,9 @@ impl Attribute { } } #[staticmethod] - pub fn dead() -> Self { + pub fn life() -> Self { Self { - entity: geometry_attribute_dead(), + entity: geometry_attribute_life(), } } @@ -124,7 +122,6 @@ impl Attribute { #[pyclass(unsendable)] pub struct Particles { pub(crate) entity: Entity, - // name → (entity, format); used by `emit(**kwargs)` to route kwargs and pack bytes name_to_attr: HashMap, } @@ -187,7 +184,24 @@ impl Particles { particles_capacity(self.entity).map_err(|e| PyRuntimeError::new_err(format!("{e}"))) } - /// backing `Buffer` for a registered attribute, or `None` if not registered. + #[pyo3(signature = (attribute, default=None))] + pub fn add_attribute( + &mut self, + attribute: PyRef, + default: Option<&Bound<'_, PyAny>>, + ) -> PyResult<()> { + let default_value = default + .map(crate::material::py_to_shader_value) + .transpose()?; + particles_attribute_add(self.entity, attribute.entity, default_value) + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let (name, fmt) = geometry_attribute_info(attribute.entity) + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + self.name_to_attr + .insert(name, (attribute.entity, AttributeFormat::from_inner(fmt))); + Ok(()) + } + pub fn buffer(&self, attribute: &Attribute) -> PyResult> { let buf = particles_buffer(self.entity, attribute.entity) .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; @@ -202,8 +216,6 @@ impl Particles { Ok(buf.map(|e| Buffer::from_entity(e, Some(element_type)))) } - /// dispatch a compute kernel against these particles' buffers. buffers are - /// auto-bound by attribute name; kwargs are forwarded to `compute.set(...)`. #[pyo3(signature = (compute, **kwargs))] pub fn apply(&self, compute: &Compute, kwargs: Option<&Bound<'_, PyDict>>) -> PyResult<()> { if let Some(kwargs) = kwargs { @@ -213,9 +225,6 @@ impl Particles { .map_err(|e| PyRuntimeError::new_err(format!("{e}"))) } - /// emit `n` particles into the next ring-buffer slots. per-attribute data - /// is a kwarg keyed by attribute name; each value is a flat list of - /// `n * format.float_count()` floats. #[pyo3(signature = (n, **kwargs))] pub fn emit(&self, n: u32, kwargs: Option<&Bound<'_, PyDict>>) -> PyResult<()> { let Some(kwargs) = kwargs else { @@ -246,12 +255,164 @@ impl Particles { particles_emit(self.entity, n, data).map_err(|e| PyRuntimeError::new_err(format!("{e}"))) } - /// emit `n` particles via a GPU kernel. auto-binds buffers and an - /// `emit_range: vec4 = (base_slot, n, capacity, 0)` uniform. pub fn emit_gpu(&self, n: u32, compute: &Compute) -> PyResult<()> { particles_emit_gpu(self.entity, n, compute.entity) .map_err(|e| PyRuntimeError::new_err(format!("{e}"))) } + + #[staticmethod] + pub fn noise() -> PyResult { + let entity = particles_kernel_noise() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn transform() -> PyResult { + let entity = particles_kernel_transform() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn attract() -> PyResult { + let entity = particles_kernel_attract() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn drag() -> PyResult { + let entity = particles_kernel_drag() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn vortex() -> PyResult { + let entity = particles_kernel_vortex() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn force() -> PyResult { + let entity = particles_kernel_force() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn integrate() -> PyResult { + let entity = particles_kernel_integrate() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn age() -> PyResult { + let entity = particles_kernel_age() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn bounds_sphere() -> PyResult { + let entity = particles_kernel_bounds_sphere() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn bounds_box() -> PyResult { + let entity = particles_kernel_bounds_box() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn bounds_geometry(geometry: &Geometry) -> PyResult { + let entity = particles_kernel_bounds_geometry(geometry.entity) + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn impulse() -> PyResult { + let entity = particles_kernel_impulse() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn flock() -> PyResult { + let entity = particles_kernel_flock() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn orient() -> PyResult { + let entity = particles_kernel_orient() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn field() -> PyResult { + let entity = particles_kernel_field() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn attr_linear() -> PyResult { + let entity = particles_kernel_attr_linear() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn attr_combine() -> PyResult { + let entity = particles_kernel_attr_combine() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn attr_mix() -> PyResult { + let entity = particles_kernel_attr_mix() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn attr_lookup1d() -> PyResult { + let entity = particles_kernel_attr_lookup1d() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn attr_lookup2d() -> PyResult { + let entity = particles_kernel_attr_lookup2d() + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn scatter_surface(geometry: &Geometry) -> PyResult { + let entity = particles_scatter_create(geometry.entity) + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } + + #[staticmethod] + pub fn scatter_volume(geometry: &Geometry) -> PyResult { + let entity = particles_scatter_volume_create(geometry.entity) + .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + Ok(Compute::from_entity(entity)) + } } impl Drop for Particles { @@ -259,17 +420,3 @@ impl Drop for Particles { let _ = particles_destroy(self.entity); } } - -/// built-in noise kernel. uniforms: `scale`, `strength`, `time`. -pub fn kernel_noise() -> PyResult { - let entity = particles_kernel_noise().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; - Ok(Compute::from_entity(entity)) -} - -/// built-in transform kernel: scale → axis-angle rotate → translate. uniforms: -/// `translate: vec3`, `rotation_axis: vec3`, `rotation_angle: f32`, `scale: vec3`. -pub fn kernel_transform() -> PyResult { - let entity = - particles_kernel_transform().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; - Ok(Compute::from_entity(entity)) -} diff --git a/crates/processing_render/src/compute.rs b/crates/processing_render/src/compute.rs index 22518d31..cfa0cd9a 100644 --- a/crates/processing_render/src/compute.rs +++ b/crates/processing_render/src/compute.rs @@ -1,15 +1,19 @@ -use std::collections::BTreeSet; +use std::collections::{BTreeSet, HashMap}; -use bevy::asset::RenderAssetUsages; +use bevy::asset::{AssetId, RenderAssetUsages}; +use bevy::mesh::MeshVertexAttribute; +use bevy::reflect::PartialReflect; use bevy::{ prelude::*, render::{ RenderApp, + mesh::{RenderMesh, allocator::MeshAllocator}, render_asset::RenderAssets, render_resource::{ BindGroupLayoutDescriptor, Buffer as WgpuBuffer, BufferDescriptor, BufferUsages, CachedComputePipelineId, CachedPipelineState, CommandEncoderDescriptor, - ComputePassDescriptor, ComputePipelineDescriptor, MapMode, PipelineCache, PollType, + ComputePassDescriptor, ComputePipelineDescriptor, MapMode, OwnedBindingResource, + PipelineCache, PollType, }, renderer::{RenderDevice, RenderQueue}, storage::{GpuShaderBuffer, ShaderBuffer}, @@ -19,7 +23,10 @@ use bevy::{ use bevy_naga_reflect::dynamic_shader::DynamicShader; -use crate::material::custom::Shader; +use crate::geometry::{Attribute, Geometry}; +use crate::image::Image as PImage; +use crate::material::custom::{Shader, apply_reflect_field, shader_value_to_reflect}; +use crate::shader_value::ShaderValue; use processing_core::error::{ProcessingError, Result}; pub struct ComputePlugin; @@ -149,12 +156,30 @@ pub fn destroy_buffer(In(entity): In, mut commands: Commands) -> Result< Ok(()) } +#[derive(Clone, Copy)] +pub enum MeshBindingRef { + Attribute { geom: Entity, attribute: Entity }, + Index { geom: Entity }, +} + +#[derive(Clone)] +pub enum ResolvedMeshBinding { + Attribute { + mesh_id: AssetId, + attribute: MeshVertexAttribute, + }, + Index { + mesh_id: AssetId, + }, +} + #[derive(Component)] pub struct Compute { pub shader: DynamicShader, pub entry_point: String, pub pipeline_id: CachedComputePipelineId, pub bind_group_layout_descriptors: Vec<(u32, BindGroupLayoutDescriptor)>, + pub mesh_bindings: HashMap, } fn queue_pipeline( @@ -196,8 +221,15 @@ pub fn create_compute(app: &mut App, shader_entity: Entity) -> Result { })?; let entry_point = compute_ep.name.clone(); - let mut shader = DynamicShader::new(module) - .map_err(|e| ProcessingError::ShaderCompilationError(e.to_string()))?; + let mut shader = DynamicShader::new(module).map_err(|e| { + let mut msg = e.to_string(); + let mut src: &dyn std::error::Error = &e; + while let Some(s) = src.source() { + msg.push_str(&format!("\n caused by: {s}")); + src = s; + } + ProcessingError::ShaderCompilationError(msg) + })?; shader.init(); let reflection = shader.reflection(); @@ -255,6 +287,7 @@ pub fn create_compute(app: &mut App, shader_entity: Entity) -> Result { entry_point, pipeline_id, bind_group_layout_descriptors, + mesh_bindings: HashMap::new(), }) .id()); } @@ -262,11 +295,124 @@ pub fn create_compute(app: &mut App, shader_entity: Entity) -> Result { Err(ProcessingError::PipelineNotReady(MAX_WAIT)) } +pub fn set_compute_property( + In((entity, name, value)): In<(Entity, String, ShaderValue)>, + mut computes: Query<&mut Compute>, + mut p_buffers: Query<&mut Buffer>, + p_images: Query<&PImage>, +) -> Result<()> { + use bevy_naga_reflect::reflect::ParameterCategory; + + let mut compute = computes + .get_mut(entity) + .map_err(|_| ProcessingError::ComputeNotFound)?; + + match value { + ShaderValue::Buffer(buf_entity) => { + let category = compute + .shader + .reflection() + .parameter(&name) + .map(|p| p.category()) + .ok_or_else(|| ProcessingError::UnknownShaderProperty(name.clone()))?; + let ParameterCategory::Storage { read_only } = category else { + return Err(ProcessingError::InvalidArgument(format!( + "property `{name}` expects {category:?}, got Buffer", + ))); + }; + let mut buffer = p_buffers + .get_mut(buf_entity) + .map_err(|_| ProcessingError::BufferNotFound)?; + compute.shader.insert(&name, buffer.handle.clone()); + if !read_only { + buffer.bound_rw = true; + } + Ok(()) + } + ShaderValue::MeshAttribute(geom_entity, attribute_entity) => { + let category = compute + .shader + .reflection() + .parameter(&name) + .map(|p| p.category()) + .ok_or_else(|| ProcessingError::UnknownShaderProperty(name.clone()))?; + let ParameterCategory::Storage { read_only } = category else { + return Err(ProcessingError::InvalidArgument(format!( + "property `{name}` expects {category:?}, got MeshAttribute", + ))); + }; + if !read_only { + return Err(ProcessingError::InvalidArgument(format!( + "property `{name}` is read-write; mesh attribute buffers can only bind as read-only", + ))); + } + compute.mesh_bindings.insert( + name, + MeshBindingRef::Attribute { + geom: geom_entity, + attribute: attribute_entity, + }, + ); + Ok(()) + } + ShaderValue::MeshIndex(geom_entity) => { + let category = compute + .shader + .reflection() + .parameter(&name) + .map(|p| p.category()) + .ok_or_else(|| ProcessingError::UnknownShaderProperty(name.clone()))?; + let ParameterCategory::Storage { read_only } = category else { + return Err(ProcessingError::InvalidArgument(format!( + "property `{name}` expects {category:?}, got MeshIndex", + ))); + }; + if !read_only { + return Err(ProcessingError::InvalidArgument(format!( + "property `{name}` is read-write; mesh index buffer can only bind as read-only", + ))); + } + compute + .mesh_bindings + .insert(name, MeshBindingRef::Index { geom: geom_entity }); + Ok(()) + } + ShaderValue::Texture(img_entity) => { + let category = compute + .shader + .reflection() + .parameter(&name) + .map(|p| p.category()) + .ok_or_else(|| ProcessingError::UnknownShaderProperty(name.clone()))?; + if !matches!( + category, + ParameterCategory::Texture + | ParameterCategory::StorageTexture + | ParameterCategory::Sampler + ) { + return Err(ProcessingError::InvalidArgument(format!( + "property `{name}` expects {category:?}, got Texture", + ))); + } + let image = p_images + .get(img_entity) + .map_err(|_| ProcessingError::ImageNotFound)?; + compute.shader.insert(&name, image.handle.clone()); + Ok(()) + } + v => { + let reflect_value: Box = shader_value_to_reflect(&v)?; + apply_reflect_field(&mut compute.shader, &name, &*reflect_value) + } + } +} + pub fn dispatch( - In((pipeline_id, layout_descriptors, shader, x, y, z)): In<( + In((pipeline_id, layout_descriptors, shader, mesh_bindings, x, y, z)): In<( CachedComputePipelineId, Vec<(u32, BindGroupLayoutDescriptor)>, DynamicShader, + Vec<(String, ResolvedMeshBinding)>, u32, u32, u32, @@ -276,6 +422,8 @@ pub fn dispatch( render_queue: Res, gpu_images: Res>, gpu_buffers: Res>, + render_meshes: Res>, + mesh_allocator: Res, ) -> Result<()> { let pipeline = pipeline_cache .get_compute_pipeline(pipeline_id) @@ -287,9 +435,25 @@ pub fn dispatch( let mut bind_groups = Vec::new(); for (group, desc) in &layout_descriptors { let layout = pipeline_cache.get_bind_group_layout(desc); - let bindings = + let mut bindings = reflection.create_bindings(*group, &shader, &render_device, &gpu_images, &gpu_buffers); + for (name, resolved) in &mesh_bindings { + let Some(param) = reflection.parameter(name) else { + return Err(ProcessingError::UnknownShaderProperty(name.clone())); + }; + if param.group() != *group { + continue; + } + let buffer = resolve_mesh_binding(resolved, &render_meshes, &mesh_allocator)?; + let binding_idx = param.binding(); + if let Some(slot) = bindings.iter_mut().find(|(b, _)| *b == binding_idx) { + slot.1 = OwnedBindingResource::Buffer(buffer); + } else { + bindings.push((binding_idx, OwnedBindingResource::Buffer(buffer))); + } + } + let bind_group_entries: Vec<_> = bindings .iter() .map( @@ -329,3 +493,77 @@ pub fn destroy_compute(In(entity): In, mut commands: Commands) -> Result commands.entity(entity).despawn(); Ok(()) } + +fn resolve_mesh_binding( + resolved: &ResolvedMeshBinding, + render_meshes: &RenderAssets, + mesh_allocator: &MeshAllocator, +) -> Result { + match resolved { + ResolvedMeshBinding::Attribute { mesh_id, attribute } => { + let render_mesh = render_meshes + .get(*mesh_id) + .ok_or(ProcessingError::GeometryNotFound)?; + let binding_idx = render_mesh + .layout + .0 + .binding_index_for_attribute(attribute.id) + .ok_or_else(|| { + ProcessingError::InvalidArgument(format!( + "mesh has no `{}` attribute (deinterleave required?)", + attribute.name + )) + })?; + let slice = mesh_allocator + .mesh_vertex_slice(mesh_id, binding_idx as u8) + .ok_or_else(|| { + ProcessingError::InvalidArgument(format!( + "mesh attribute `{}` not yet allocated", + attribute.name + )) + })?; + Ok(slice.buffer.clone()) + } + ResolvedMeshBinding::Index { mesh_id } => { + let slice = mesh_allocator.mesh_index_slice(mesh_id).ok_or_else(|| { + ProcessingError::InvalidArgument( + "mesh has no index buffer or it is not yet allocated".to_string(), + ) + })?; + Ok(slice.buffer.clone()) + } + } +} + +pub fn resolve_mesh_bindings( + world: &World, + compute: &Compute, +) -> Result> { + let mut out = Vec::with_capacity(compute.mesh_bindings.len()); + for (name, mesh_ref) in &compute.mesh_bindings { + let resolved = match mesh_ref { + MeshBindingRef::Attribute { geom, attribute } => { + let g = world + .get::(*geom) + .ok_or(ProcessingError::GeometryNotFound)?; + let a = world + .get::(*attribute) + .ok_or(ProcessingError::InvalidEntity)?; + ResolvedMeshBinding::Attribute { + mesh_id: g.handle.id(), + attribute: a.inner, + } + } + MeshBindingRef::Index { geom } => { + let g = world + .get::(*geom) + .ok_or(ProcessingError::GeometryNotFound)?; + ResolvedMeshBinding::Index { + mesh_id: g.handle.id(), + } + } + }; + out.push((name.clone(), resolved)); + } + Ok(out) +} diff --git a/crates/processing_render/src/geometry/attribute.rs b/crates/processing_render/src/geometry/attribute.rs index 6ca8d466..d7f153b6 100644 --- a/crates/processing_render/src/geometry/attribute.rs +++ b/crates/processing_render/src/geometry/attribute.rs @@ -157,6 +157,10 @@ impl AttributeFormat { } } + pub fn components(self) -> usize { + self.byte_size() / 4 + } + pub fn from_u8(value: u8) -> Option { match value { 1 => Some(Self::Float), @@ -177,9 +181,7 @@ pub struct Attribute { impl Attribute { pub fn new(name: impl Into, format: AttributeFormat) -> Self { - // we leak here to get a 'static str for the attribute name, but this is okay because - // we never expect to unload attributes during the lifetime of the application - // and attribute names are generally small in number + // leaked for a 'static name; attributes are never unloaded and are few. let name: &'static str = Box::leak(name.into().into_boxed_str()); let id = hash_attr_name(name); let inner = MeshVertexAttribute::new(name, id, format.to_vertex_format()); @@ -198,8 +200,6 @@ impl Attribute { } } - /// like [`Self::from_builtin`], but with a user-facing `name` distinct - /// from `inner.name`. shaders bind by `name`. pub fn from_builtin_with_name( name: &'static str, inner: MeshVertexAttribute, @@ -223,12 +223,11 @@ pub struct BuiltinAttributes { pub normal: Entity, pub color: Entity, pub uv: Entity, - /// per-instance rotation as a quaternion `(x, y, z, w)`. pub rotation: Entity, - /// per-instance scale `(x, y, z)`. pub scale: Entity, - /// per-particle lifecycle flag: `0.0` = alive, non-zero = dead. - pub dead: Entity, + pub life: Entity, + pub velocity: Entity, + pub age: Entity, } impl FromWorld for BuiltinAttributes { @@ -267,8 +266,14 @@ impl FromWorld for BuiltinAttributes { let scale = world .spawn(Attribute::new("scale", AttributeFormat::Float3)) .id(); - let dead = world - .spawn(Attribute::new("dead", AttributeFormat::Float)) + let life = world + .spawn(Attribute::new("life", AttributeFormat::Float)) + .id(); + let velocity = world + .spawn(Attribute::new("velocity", AttributeFormat::Float3)) + .id(); + let age = world + .spawn(Attribute::new("age", AttributeFormat::Float)) .id(); Self { @@ -278,11 +283,40 @@ impl FromWorld for BuiltinAttributes { uv, rotation, scale, - dead, + life, + velocity, + age, } } } +impl BuiltinAttributes { + pub fn by_name(&self, name: &str) -> Option { + Some(match name { + "position" => self.position, + "normal" => self.normal, + "color" => self.color, + "uv" => self.uv, + "rotation" => self.rotation, + "scale" => self.scale, + "life" => self.life, + "velocity" => self.velocity, + "age" => self.age, + _ => return None, + }) + } +} + +pub fn default_attribute_init(name: &str, format: AttributeFormat) -> Vec { + match name { + "life" => vec![1.0], + "scale" => vec![1.0, 1.0, 1.0], + "color" => vec![1.0, 1.0, 1.0, 1.0], + "rotation" => vec![0.0, 0.0, 0.0, 1.0], + _ => vec![0.0; format.components()], + } +} + pub fn create( In((name, format)): In<(String, AttributeFormat)>, mut commands: Commands, @@ -423,3 +457,45 @@ pub fn set_attribute( )), } } + +#[cfg(test)] +mod tests { + use super::*; + + #[test] + fn builtin_defaults_match_format() { + use AttributeFormat::*; + let builtins = [ + ("position", Float3), + ("normal", Float3), + ("color", Float4), + ("uv", Float2), + ("rotation", Float4), + ("scale", Float3), + ("life", Float), + ("velocity", Float3), + ("age", Float), + ]; + for (name, format) in builtins { + assert_eq!( + default_attribute_init(name, format).len(), + format.components(), + "default_attribute_init({name:?}) does not match {format:?} component count", + ); + } + } + + #[test] + fn unknown_name_seeds_zero_for_its_format() { + for format in [ + AttributeFormat::Float, + AttributeFormat::Float2, + AttributeFormat::Float3, + AttributeFormat::Float4, + ] { + let seed = default_attribute_init("custom_thing", format); + assert_eq!(seed.len(), format.components()); + assert!(seed.iter().all(|&f| f == 0.0)); + } + } +} diff --git a/crates/processing_render/src/graphics.rs b/crates/processing_render/src/graphics.rs index 5ebae48d..9cb39ee0 100644 --- a/crates/processing_render/src/graphics.rs +++ b/crates/processing_render/src/graphics.rs @@ -9,6 +9,7 @@ use bevy::{ ImageRenderTarget, MsaaWriteback, Projection, RenderTarget, visibility::RenderLayers, }, core_pipeline::tonemapping::Tonemapping, + post_process::bloom::Bloom, ecs::query::QueryEntityError, math::{Mat4, Vec3A}, prelude::*, @@ -209,7 +210,7 @@ pub fn create( ..default() }, target, - // tonemapping prevents color accurate readback, so we disable it + // overridden below for hdr targets Tonemapping::None, // we need to be able to write to the texture CameraMainTextureUsages::default().with(TextureUsages::COPY_DST), @@ -227,9 +228,8 @@ pub fn create( }, )); - // only enable Hdr for floating-point texture formats if is_hdr { - entity_commands.insert(Hdr); + entity_commands.insert((Hdr, Bloom::NATURAL, Tonemapping::TonyMcMapface)); } let entity = entity_commands.id(); @@ -426,6 +426,68 @@ pub fn ortho( Ok(()) } +pub fn world_from_screen( + In((entity, sx, sy, depth)): In<(Entity, f32, f32, f32)>, + cameras: Query<(&bevy::camera::Camera, &GlobalTransform)>, +) -> Result { + let (camera, transform) = cameras + .get(entity) + .map_err(|_| ProcessingError::GraphicsNotFound)?; + + let ndc_xy = camera + .viewport_to_ndc(Vec2::new(sx, sy)) + .map_err(|_| ProcessingError::GraphicsNotFound)?; + let ndc_z = (1.0 - depth).max(f32::EPSILON); + let world: Vec3 = camera + .ndc_to_world(transform, ndc_xy.extend(ndc_z)) + .ok_or(ProcessingError::GraphicsNotFound)?; + Ok(world) +} + +pub fn set_bloom( + In((entity, intensity, threshold)): In<(Entity, f32, f32)>, + mut commands: Commands, + mut tonemapping_query: Query<&mut Tonemapping>, +) -> Result<()> { + use bevy::post_process::bloom::{Bloom, BloomCompositeMode, BloomPrefilter}; + + let mut bloom = Bloom::NATURAL; + bloom.intensity = intensity; + if threshold > 0.0 { + bloom.composite_mode = BloomCompositeMode::Additive; + bloom.prefilter = BloomPrefilter { + threshold, + threshold_softness: 0.5, + }; + } + + commands.entity(entity).insert((bloom, Hdr)); + + if let Ok(mut tm) = tonemapping_query.get_mut(entity) { + if *tm == Tonemapping::None { + *tm = Tonemapping::TonyMcMapface; + } + } + + Ok(()) +} + +pub fn remove_bloom( + In(entity): In, + mut commands: Commands, + mut tonemapping_query: Query<&mut Tonemapping>, +) -> Result<()> { + use bevy::post_process::bloom::Bloom; + + commands.entity(entity).remove::(); + + if let Ok(mut tm) = tonemapping_query.get_mut(entity) { + *tm = Tonemapping::None; + } + + Ok(()) +} + pub fn destroy( In(entity): In, mut commands: Commands, diff --git a/crates/processing_render/src/lib.rs b/crates/processing_render/src/lib.rs index 094e9d49..b7ef15ce 100644 --- a/crates/processing_render/src/lib.rs +++ b/crates/processing_render/src/lib.rs @@ -20,6 +20,22 @@ pub mod text; pub mod time; pub mod transform; +pub use particles::{ + BOUNDS_CLAMP, BOUNDS_REFLECT, BOUNDS_SOFT, BOUNDS_WRAP, COMBINE_ADD, COMBINE_DIV, COMBINE_MAX, + COMBINE_MIN, COMBINE_MUL, COMBINE_POW, COMBINE_SUB, FALLOFF_CONST, FALLOFF_CUBIC, + FALLOFF_INVERSE, FALLOFF_LINEAR, FALLOFF_QUADRATIC, FALLOFF_SMOOTHSTEP, particles_apply, + particles_attribute_add, particles_buffer, particles_capacity, particles_create, + particles_create_from_geometry, particles_destroy, particles_emit, particles_emit_gpu, + particles_kernel_age, + particles_kernel_attr_combine, particles_kernel_attr_linear, particles_kernel_attr_lookup1d, + particles_kernel_attr_lookup2d, particles_kernel_attr_mix, particles_kernel_attract, + particles_kernel_bounds_box, particles_kernel_bounds_geometry, particles_kernel_bounds_sphere, + particles_kernel_drag, particles_kernel_field, particles_kernel_flock, particles_kernel_force, + particles_kernel_impulse, particles_kernel_integrate, particles_kernel_noise, + particles_kernel_orient, particles_kernel_transform, particles_kernel_vortex, + particles_scatter_create, particles_scatter_volume_create, +}; + use std::path::PathBuf; use bevy::{ @@ -851,6 +867,42 @@ pub fn graphics_ortho( }) } +pub fn graphics_world_from_screen( + graphics_entity: Entity, + sx: f32, + sy: f32, + depth: f32, +) -> error::Result { + app_mut(|app| { + app.world_mut() + .run_system_cached_with(graphics::world_from_screen, (graphics_entity, sx, sy, depth)) + .unwrap() + }) +} + +pub fn graphics_set_bloom( + graphics_entity: Entity, + intensity: f32, + threshold: f32, +) -> error::Result<()> { + app_mut(|app| { + app.world_mut() + .run_system_cached_with( + graphics::set_bloom, + (graphics_entity, intensity, threshold), + ) + .unwrap() + }) +} + +pub fn graphics_remove_bloom(graphics_entity: Entity) -> error::Result<()> { + app_mut(|app| { + app.world_mut() + .run_system_cached_with(graphics::remove_bloom, graphics_entity) + .unwrap() + }) +} + pub fn transform_set_position(entity: Entity, position: Vec3) -> error::Result<()> { app_mut(|app| { app.world_mut() @@ -1277,8 +1329,22 @@ pub fn geometry_attribute_scale() -> Entity { app_mut(|app| Ok(app.world().resource::().scale)).unwrap() } -pub fn geometry_attribute_dead() -> Entity { - app_mut(|app| Ok(app.world().resource::().dead)).unwrap() +pub fn geometry_attribute_life() -> Entity { + app_mut(|app| Ok(app.world().resource::().life)).unwrap() +} + +pub fn geometry_attribute_velocity() -> Entity { + app_mut(|app| { + Ok(app + .world() + .resource::() + .velocity) + }) + .unwrap() +} + +pub fn geometry_attribute_age() -> Entity { + app_mut(|app| Ok(app.world().resource::().age)).unwrap() } pub fn geometry_attribute_destroy(entity: Entity) -> error::Result<()> { @@ -1715,12 +1781,16 @@ pub fn material_set_albedo_color(entity: Entity, color: [f32; 4]) -> error::Resu }) } -/// set the albedo source to a per-particle color buffer (`Float4` per slot, -/// indexed by `mesh.tag`). Preserves all other `StandardMaterial` fields; -/// `base_color` modulates the buffer color. -pub fn material_set_albedo_buffer( +#[derive(Copy, Clone)] +enum ParticlesBufferSlot { + Albedo, + Emissive, +} + +fn material_set_particles_buffer( entity: Entity, - color_buffer_entity: Entity, + buffer_entity: Entity, + slot: ParticlesBufferSlot, ) -> error::Result<()> { use crate::material::ProcessingMaterial; use crate::particles::material::{ParticlesExtension, ParticlesMaterial}; @@ -1732,7 +1802,7 @@ pub fn material_set_albedo_buffer( app_mut(|app| { let buffer_handle = app .world() - .get::(color_buffer_entity) + .get::(buffer_entity) .ok_or(error::ProcessingError::BufferNotFound)? .handle .clone(); @@ -1748,7 +1818,11 @@ pub fn material_set_albedo_buffer( let mat = mats .get_mut(&handle) .ok_or(error::ProcessingError::MaterialNotFound)?; - mat.into_inner().extension.colors = buffer_handle; + let ext = &mut mat.into_inner().extension; + match slot { + ParticlesBufferSlot::Albedo => ext.colors = Some(buffer_handle), + ParticlesBufferSlot::Emissive => ext.emissive_colors = Some(buffer_handle), + } return Ok(()); } @@ -1766,13 +1840,21 @@ pub fn material_set_albedo_buffer( mats.remove(&handle); base }; + let extension = match slot { + ParticlesBufferSlot::Albedo => ParticlesExtension { + colors: Some(buffer_handle), + emissive_colors: None, + }, + ParticlesBufferSlot::Emissive => ParticlesExtension { + colors: None, + emissive_colors: Some(buffer_handle), + }, + }; let new_handle = world .resource_mut::>() .add(ExtendedMaterial { base: preserved, - extension: ParticlesExtension { - colors: buffer_handle, - }, + extension, }); world .entity_mut(entity) @@ -1781,6 +1863,20 @@ pub fn material_set_albedo_buffer( }) } +pub fn material_set_albedo_buffer( + entity: Entity, + color_buffer_entity: Entity, +) -> error::Result<()> { + material_set_particles_buffer(entity, color_buffer_entity, ParticlesBufferSlot::Albedo) +} + +pub fn material_set_emissive_buffer( + entity: Entity, + emissive_buffer_entity: Entity, +) -> error::Result<()> { + material_set_particles_buffer(entity, emissive_buffer_entity, ParticlesBufferSlot::Emissive) +} + pub fn material_set( entity: Entity, name: impl Into, @@ -2228,14 +2324,16 @@ pub fn compute_dispatch(entity: Entity, x: u32, y: u32, z: u32) -> error::Result app.update(); let args = { - let c = app - .world() + let world = app.world(); + let c = world .get::(entity) .ok_or(error::ProcessingError::ComputeNotFound)?; + let mesh_bindings = compute::resolve_mesh_bindings(world, c)?; ( c.pipeline_id, c.bind_group_layout_descriptors.clone(), c.shader.clone(), + mesh_bindings, x, y, z, @@ -2256,270 +2354,6 @@ pub fn compute_destroy(entity: Entity) -> error::Result<()> { }) } -pub fn particles_create(capacity: u32, attribute_entities: Vec) -> error::Result { - app_mut(|app| { - app.world_mut() - .run_system_cached_with(particles::create, (capacity, attribute_entities)) - .unwrap() - }) -} - -/// capacity = `geometry`'s vertex count. Builtin attributes (`position`, -/// `normal`, `color`, `uv`) are seeded from the matching mesh attribute when -/// formats line up; everything else is zero-initialized. -pub fn particles_create_from_geometry( - geometry_entity: Entity, - attribute_entities: Vec, -) -> error::Result { - app_mut(|app| { - app.world_mut() - .run_system_cached_with( - particles::create_from_geometry, - (geometry_entity, attribute_entities), - ) - .unwrap() - }) -} - -pub fn particles_destroy(entity: Entity) -> error::Result<()> { - app_mut(|app| { - app.world_mut() - .run_system_cached_with(particles::destroy, entity) - .unwrap() - }) -} - -pub fn particles_capacity(entity: Entity) -> error::Result { - app_mut(|app| { - Ok(app - .world() - .get::(entity) - .ok_or(error::ProcessingError::ParticlesNotFound)? - .capacity) - }) -} - -pub fn particles_buffer(entity: Entity, attribute_entity: Entity) -> error::Result> { - app_mut(|app| { - Ok(app - .world() - .get::(entity) - .ok_or(error::ProcessingError::ParticlesNotFound)? - .buffer(attribute_entity)) - }) -} - -/// GPU-driven emission into the next `count` ring-buffer slots. Auto-binds -/// attribute buffers (same convention as [`particles_apply`]) and an -/// `emit_range: vec4 = (base_slot, count, capacity, 0)` uniform. -pub fn particles_emit_gpu( - particles_entity: Entity, - count: u32, - compute_entity: Entity, -) -> error::Result<()> { - if count == 0 { - return Ok(()); - } - const WORKGROUP_SIZE: u32 = 64; - - let (capacity, head, buffers) = app_mut(|app| { - let world = app.world(); - let field = world - .get::(particles_entity) - .ok_or(error::ProcessingError::ParticlesNotFound)?; - if count > field.capacity { - return Err(error::ProcessingError::InvalidArgument(format!( - "particles_emit_gpu count={} exceeds field capacity {}", - count, field.capacity - ))); - } - let mut buffers: Vec<(String, Entity)> = Vec::with_capacity(field.buffers.len()); - for (&attr_entity, &buf_entity) in &field.buffers { - let attr = world - .get::(attr_entity) - .ok_or(error::ProcessingError::InvalidEntity)?; - buffers.push((attr.name.to_string(), buf_entity)); - } - Ok((field.capacity, field.emit_head, buffers)) - })?; - - for (name, buf_entity) in buffers { - match compute_set( - compute_entity, - name, - shader_value::ShaderValue::Buffer(buf_entity), - ) { - Ok(()) => {} - Err(error::ProcessingError::UnknownShaderProperty(_)) => {} - Err(e) => return Err(e), - } - } - - match compute_set( - compute_entity, - "emit_range", - shader_value::ShaderValue::Float4([head as f32, count as f32, capacity as f32, 0.0]), - ) { - Ok(()) => {} - Err(error::ProcessingError::UnknownShaderProperty(_)) => {} - Err(e) => return Err(e), - } - - let workgroup_count = count.div_ceil(WORKGROUP_SIZE); - compute_dispatch(compute_entity, workgroup_count, 1, 1)?; - - app_mut(|app| { - let mut field = app - .world_mut() - .get_mut::(particles_entity) - .ok_or(error::ProcessingError::ParticlesNotFound)?; - field.emit_head = (field.emit_head + count) % field.capacity; - Ok(()) - }) -} - -/// CPU-driven emission. Writes per-attribute byte payloads into the next `n` -/// ring-buffer slots. Each entry in `attribute_data` must be exactly -/// `attr.byte_size * n` bytes. On wrap, oldest slots are overwritten. -pub fn particles_emit( - particles_entity: Entity, - n: u32, - attribute_data: Vec<(Entity, Vec)>, -) -> error::Result<()> { - if n == 0 { - return Ok(()); - } - - let (capacity, head, attr_specs) = app_mut(|app| { - let world = app.world(); - let field = world - .get::(particles_entity) - .ok_or(error::ProcessingError::ParticlesNotFound)?; - if n > field.capacity { - return Err(error::ProcessingError::InvalidArgument(format!( - "particles_emit n={} exceeds field capacity {}", - n, field.capacity - ))); - } - let mut specs: Vec<(Entity, u32, Entity)> = Vec::with_capacity(attribute_data.len()); - for (attr_entity, _) in &attribute_data { - let attr = world - .get::(*attr_entity) - .ok_or(error::ProcessingError::InvalidEntity)?; - let buf = field.buffer(*attr_entity).ok_or_else(|| { - error::ProcessingError::InvalidArgument(format!( - "particles have no buffer for attribute {:?}", - attr_entity - )) - })?; - specs.push((*attr_entity, attr.format.byte_size() as u32, buf)); - } - Ok((field.capacity, field.emit_head, specs)) - })?; - - for ((_, bytes), &(_, byte_size, buf)) in attribute_data.iter().zip(attr_specs.iter()) { - let expected = (n as usize) * (byte_size as usize); - if bytes.len() != expected { - return Err(error::ProcessingError::InvalidArgument(format!( - "expected {} bytes ({} particles * {} bytes), got {}", - expected, - n, - byte_size, - bytes.len() - ))); - } - let first_chunk_n = (capacity - head).min(n); - let split = (first_chunk_n as usize) * (byte_size as usize); - let first_offset = (head as u64) * (byte_size as u64); - buffer_write_element(buf, first_offset, bytes[..split].to_vec())?; - if first_chunk_n < n { - buffer_write_element(buf, 0, bytes[split..].to_vec())?; - } - } - - app_mut(|app| { - let mut field = app - .world_mut() - .get_mut::(particles_entity) - .ok_or(error::ProcessingError::ParticlesNotFound)?; - field.emit_head = (field.emit_head + n) % field.capacity; - Ok(()) - }) -} - -/// built-in noise kernel: displaces `position` by 3d value noise. Uniforms: -/// `scale: f32`, `strength: f32`, `time: f32`. -pub fn particles_kernel_noise() -> error::Result { - let shader = shader_load(particles::kernels::NOISE_PATH)?; - compute_create(shader) -} - -/// built-in transform kernel: scale → axis-angle rotate → translate on -/// `position`. Uniforms: `translate: vec3`, `rotation_axis: vec3`, -/// `rotation_angle: f32`, `scale: vec3`. Identity defaults are seeded. -pub fn particles_kernel_transform() -> error::Result { - let shader = shader_load(particles::kernels::TRANSFORM_PATH)?; - let entity = compute_create(shader)?; - compute_set( - entity, - "translate", - shader_value::ShaderValue::Float3([0.0; 3]), - )?; - compute_set( - entity, - "rotation_axis", - shader_value::ShaderValue::Float3([0.0, 1.0, 0.0]), - )?; - compute_set( - entity, - "rotation_angle", - shader_value::ShaderValue::Float(0.0), - )?; - compute_set( - entity, - "scale", - shader_value::ShaderValue::Float3([1.0, 1.0, 1.0]), - )?; - Ok(entity) -} - -/// dispatch `compute_entity` against the [`Particles`]'s buffers. Each buffer -/// is auto-bound by attribute name; undeclared bindings are skipped. Kernels -/// must declare `@workgroup_size(64)`. Set uniforms via `compute_set` first. -pub fn particles_apply(particles_entity: Entity, compute_entity: Entity) -> error::Result<()> { - const WORKGROUP_SIZE: u32 = 64; - - let (capacity, buffers) = app_mut(|app| { - let world = app.world(); - let field = world - .get::(particles_entity) - .ok_or(error::ProcessingError::ParticlesNotFound)?; - let mut buffers: Vec<(String, Entity)> = Vec::with_capacity(field.buffers.len()); - for (&attr_entity, &buf_entity) in &field.buffers { - let attr = world - .get::(attr_entity) - .ok_or(error::ProcessingError::InvalidEntity)?; - buffers.push((attr.name.to_string(), buf_entity)); - } - Ok((field.capacity, buffers)) - })?; - - for (name, buf_entity) in buffers { - match compute_set( - compute_entity, - name, - shader_value::ShaderValue::Buffer(buf_entity), - ) { - Ok(()) => {} - Err(error::ProcessingError::UnknownShaderProperty(_)) => {} - Err(e) => return Err(e), - } - } - - let workgroup_count = capacity.div_ceil(WORKGROUP_SIZE); - compute_dispatch(compute_entity, workgroup_count, 1, 1) -} - // --- Font API --- /// Load a font file and return a font entity handle. diff --git a/crates/processing_render/src/material/custom.rs b/crates/processing_render/src/material/custom.rs index 7f895b14..443e9c83 100644 --- a/crates/processing_render/src/material/custom.rs +++ b/crates/processing_render/src/material/custom.rs @@ -181,7 +181,8 @@ pub fn load_shader(In(path): In, world: &mut World) -> Result { }; use bevy::ecs::system::RunSystemOnce; - // url-scheme paths parse as-is; others go through the configured asset dir + // url-scheme paths (e.g. `embedded://crate/foo.wgsl`) carry their own + // source; relative paths route through the configured asset directory let asset_path: AssetPath = if path.contains("://") { AssetPath::parse(&path).into_owned() } else { @@ -311,9 +312,13 @@ pub(crate) fn shader_value_to_reflect(value: &ShaderValue) -> Result Box::new(IVec4::from_array(*v)), ShaderValue::UInt(v) => Box::new(*v), ShaderValue::Mat4(v) => Box::new(Mat4::from_cols_array(v)), - ShaderValue::Texture(_) | ShaderValue::Buffer(_) => { + ShaderValue::Texture(_) + | ShaderValue::Buffer(_) + | ShaderValue::MeshAttribute(..) + | ShaderValue::MeshIndex(_) => { return Err(ProcessingError::InvalidArgument( - "Texture/Buffer must be bound via set_property, not as a uniform value".to_string(), + "Texture/Buffer/Mesh* must be bound via set_property, not as a uniform value" + .to_string(), )); } }) diff --git a/crates/processing_render/src/particles/emit.rs b/crates/processing_render/src/particles/emit.rs new file mode 100644 index 00000000..12df47d8 --- /dev/null +++ b/crates/processing_render/src/particles/emit.rs @@ -0,0 +1,180 @@ +use bevy::prelude::*; + +use processing_core::app_mut; +use processing_core::error; + +use crate::geometry; +use crate::particles::kernels::KernelRequires; +use crate::particles::{Particles, particles_ensure_attribute}; +use crate::shader_value::ShaderValue; +use crate::{buffer_write_element, compute_dispatch, compute_set}; + +const WORKGROUP_SIZE: u32 = 64; + +pub fn particles_emit_gpu( + particles_entity: Entity, + count: u32, + compute_entity: Entity, +) -> error::Result<()> { + if count == 0 { + return Ok(()); + } + + let (capacity, head, buffers) = app_mut(|app| { + let world = app.world(); + let field = world + .get::(particles_entity) + .ok_or(error::ProcessingError::ParticlesNotFound)?; + if count > field.capacity { + return Err(error::ProcessingError::InvalidArgument(format!( + "particles_emit_gpu count={} exceeds field capacity {}", + count, field.capacity + ))); + } + let mut buffers: Vec<(String, Entity)> = Vec::with_capacity(field.buffers.len()); + for (&attr_entity, &buf_entity) in &field.buffers { + let attr = world + .get::(attr_entity) + .ok_or(error::ProcessingError::InvalidEntity)?; + buffers.push((attr.name.to_string(), buf_entity)); + } + Ok((field.capacity, field.emit_head, buffers)) + })?; + + for (name, buf_entity) in buffers { + match compute_set(compute_entity, name, ShaderValue::Buffer(buf_entity)) { + Ok(()) => {} + Err(error::ProcessingError::UnknownShaderProperty(_)) => {} + Err(e) => return Err(e), + } + } + + for (name, value) in [ + ("emit_base", ShaderValue::UInt(head)), + ("emit_count", ShaderValue::UInt(count)), + ("emit_capacity", ShaderValue::UInt(capacity)), + ] { + match compute_set(compute_entity, name, value) { + Ok(()) => {} + Err(error::ProcessingError::UnknownShaderProperty(_)) => {} + Err(e) => return Err(e), + } + } + + let workgroup_count = count.div_ceil(WORKGROUP_SIZE); + compute_dispatch(compute_entity, workgroup_count, 1, 1)?; + + app_mut(|app| { + let mut field = app + .world_mut() + .get_mut::(particles_entity) + .ok_or(error::ProcessingError::ParticlesNotFound)?; + field.emit_head = (field.emit_head + count) % field.capacity; + Ok(()) + }) +} + +pub fn particles_emit( + particles_entity: Entity, + n: u32, + attribute_data: Vec<(Entity, Vec)>, +) -> error::Result<()> { + if n == 0 { + return Ok(()); + } + + let (capacity, head, attr_specs) = app_mut(|app| { + let world = app.world(); + let field = world + .get::(particles_entity) + .ok_or(error::ProcessingError::ParticlesNotFound)?; + if n > field.capacity { + return Err(error::ProcessingError::InvalidArgument(format!( + "particles_emit n={} exceeds field capacity {}", + n, field.capacity + ))); + } + let mut specs: Vec<(Entity, u32, Entity)> = Vec::with_capacity(attribute_data.len()); + for (attr_entity, _) in &attribute_data { + let attr = world + .get::(*attr_entity) + .ok_or(error::ProcessingError::InvalidEntity)?; + let buf = field.buffer(*attr_entity).ok_or_else(|| { + error::ProcessingError::InvalidArgument(format!( + "particles have no buffer for attribute {:?}", + attr_entity + )) + })?; + specs.push((*attr_entity, attr.format.byte_size() as u32, buf)); + } + Ok((field.capacity, field.emit_head, specs)) + })?; + + for ((_, bytes), &(_, byte_size, buf)) in attribute_data.iter().zip(attr_specs.iter()) { + let expected = (n as usize) * (byte_size as usize); + if bytes.len() != expected { + return Err(error::ProcessingError::InvalidArgument(format!( + "expected {} bytes ({} particles * {} bytes), got {}", + expected, + n, + byte_size, + bytes.len() + ))); + } + let first_chunk_n = (capacity - head).min(n); + let split = (first_chunk_n as usize) * (byte_size as usize); + let first_offset = (head as u64) * (byte_size as u64); + buffer_write_element(buf, first_offset, bytes[..split].to_vec())?; + if first_chunk_n < n { + buffer_write_element(buf, 0, bytes[split..].to_vec())?; + } + } + + app_mut(|app| { + let mut field = app + .world_mut() + .get_mut::(particles_entity) + .ok_or(error::ProcessingError::ParticlesNotFound)?; + field.emit_head = (field.emit_head + n) % field.capacity; + Ok(()) + }) +} + +pub fn particles_apply(particles_entity: Entity, compute_entity: Entity) -> error::Result<()> { + let required: Vec = app_mut(|app| { + Ok(app + .world() + .get::(compute_entity) + .map(|r| r.0.clone()) + .unwrap_or_default()) + })?; + for attr_entity in required { + particles_ensure_attribute(particles_entity, attr_entity)?; + } + + let (capacity, buffers) = app_mut(|app| { + let world = app.world(); + let field = world + .get::(particles_entity) + .ok_or(error::ProcessingError::ParticlesNotFound)?; + let mut buffers: Vec<(String, Entity)> = Vec::with_capacity(field.buffers.len()); + for (&attr_entity, &buf_entity) in &field.buffers { + let attr = world + .get::(attr_entity) + .ok_or(error::ProcessingError::InvalidEntity)?; + buffers.push((attr.name.to_string(), buf_entity)); + } + Ok((field.capacity, buffers)) + })?; + + for (name, buf_entity) in buffers { + match compute_set(compute_entity, name, ShaderValue::Buffer(buf_entity)) { + Ok(()) => {} + Err(error::ProcessingError::UnknownShaderProperty(_)) => {} + Err(e) => return Err(e), + } + } + + let workgroup_count = capacity.div_ceil(WORKGROUP_SIZE); + compute_dispatch(compute_entity, workgroup_count, 1, 1) +} diff --git a/crates/processing_render/src/particles/kernels/age.wgsl b/crates/processing_render/src/particles/kernels/age.wgsl new file mode 100644 index 00000000..65cea10f --- /dev/null +++ b/crates/processing_render/src/particles/kernels/age.wgsl @@ -0,0 +1,20 @@ +struct Params { + dt: f32, +} + +@group(0) @binding(0) var age: array; +@group(0) @binding(1) var life: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&age); + if i >= count { return; } + + if life[i] <= 0.0 { return; } + age[i] = age[i] + params.dt; + if age[i] >= life[i] { + life[i] = 0.0; + } +} diff --git a/crates/processing_render/src/particles/kernels/attr_combine.wgsl b/crates/processing_render/src/particles/kernels/attr_combine.wgsl new file mode 100644 index 00000000..826d0e7c --- /dev/null +++ b/crates/processing_render/src/particles/kernels/attr_combine.wgsl @@ -0,0 +1,34 @@ +struct Params { + op: u32, + b_scale: f32, + b_offset: f32, +} + +@group(0) @binding(0) var op_a: array; +@group(0) @binding(1) var op_b: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&op_a); + if i >= count { return; } + + let a = op_a[i]; + let b = op_b[i] * params.b_scale + params.b_offset; + + var r: f32; + switch params.op { + case 0u: { r = a + b; } + case 1u: { r = a - b; } + case 2u: { r = a * b; } + case 3u: { + if b == 0.0 { r = a; } else { r = a / b; } + } + case 4u: { r = min(a, b); } + case 5u: { r = max(a, b); } + case 6u: { r = pow(max(a, 0.0), b); } + default: { r = a; } + } + op_a[i] = r; +} diff --git a/crates/processing_render/src/particles/kernels/attr_linear.wgsl b/crates/processing_render/src/particles/kernels/attr_linear.wgsl new file mode 100644 index 00000000..ed931511 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/attr_linear.wgsl @@ -0,0 +1,15 @@ +struct Params { + scale: f32, + offset: f32, +} + +@group(0) @binding(0) var op: array; +@group(0) @binding(1) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&op); + if i >= count { return; } + op[i] = op[i] * params.scale + params.offset; +} diff --git a/crates/processing_render/src/particles/kernels/attr_lookup1d.wgsl b/crates/processing_render/src/particles/kernels/attr_lookup1d.wgsl new file mode 100644 index 00000000..b1551c46 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/attr_lookup1d.wgsl @@ -0,0 +1,27 @@ +struct Params { + scale: f32, + offset: f32, +} + +@group(0) @binding(0) var op_in: array; +@group(0) @binding(1) var op_out_r: array; +@group(0) @binding(2) var op_out_g: array; +@group(0) @binding(3) var op_out_b: array; +@group(0) @binding(4) var op_out_a: array; +@group(0) @binding(5) var params: Params; +@group(0) @binding(6) var ramp: texture_2d; +@group(0) @binding(7) var ramp_sampler: sampler; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&op_in); + if i >= count { return; } + + let t = clamp(op_in[i] * params.scale + params.offset, 0.0, 1.0); + let c = textureSampleLevel(ramp, ramp_sampler, vec2(t, 0.5), 0.0); + op_out_r[i] = c.r; + op_out_g[i] = c.g; + op_out_b[i] = c.b; + op_out_a[i] = c.a; +} diff --git a/crates/processing_render/src/particles/kernels/attr_lookup2d.wgsl b/crates/processing_render/src/particles/kernels/attr_lookup2d.wgsl new file mode 100644 index 00000000..f5a61d0d --- /dev/null +++ b/crates/processing_render/src/particles/kernels/attr_lookup2d.wgsl @@ -0,0 +1,26 @@ +struct Params { + u_scale: f32, + u_offset: f32, + v_scale: f32, + v_offset: f32, + color_scale: f32, +} + +@group(0) @binding(0) var op_in_u: array; +@group(0) @binding(1) var op_in_v: array; +@group(0) @binding(2) var op_out: array>; +@group(0) @binding(3) var params: Params; +@group(0) @binding(4) var lookup: texture_2d; +@group(0) @binding(5) var lookup_sampler: sampler; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&op_out); + if i >= count { return; } + + let u = clamp(op_in_u[i] * params.u_scale + params.u_offset, 0.0, 1.0); + let v = clamp(op_in_v[i] * params.v_scale + params.v_offset, 0.0, 1.0); + let c = textureSampleLevel(lookup, lookup_sampler, vec2(u, v), 0.0); + op_out[i] = vec4(c.rgb * params.color_scale, c.a); +} diff --git a/crates/processing_render/src/particles/kernels/attr_mix.wgsl b/crates/processing_render/src/particles/kernels/attr_mix.wgsl new file mode 100644 index 00000000..e9c61376 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/attr_mix.wgsl @@ -0,0 +1,23 @@ +struct Params { + t_scale: f32, + t_offset: f32, + t_clamp: u32, +} + +@group(0) @binding(0) var op_a: array; +@group(0) @binding(1) var op_b: array; +@group(0) @binding(2) var op_t: array; +@group(0) @binding(3) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&op_a); + if i >= count { return; } + + var t = op_t[i] * params.t_scale + params.t_offset; + if params.t_clamp != 0u { + t = clamp(t, 0.0, 1.0); + } + op_a[i] = mix(op_a[i], op_b[i], t); +} diff --git a/crates/processing_render/src/particles/kernels/attract.wgsl b/crates/processing_render/src/particles/kernels/attract.wgsl new file mode 100644 index 00000000..bcf45693 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/attract.wgsl @@ -0,0 +1,48 @@ +struct Params { + center: vec3, + _pad0: f32, + strength: f32, + radius: f32, + falloff_mode: u32, + _pad1: u32, +} + +@group(0) @binding(0) var position: array; +@group(0) @binding(1) var velocity: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + if i >= count { return; } + + let pi = i * 3u; + let pos = vec3(position[pi], position[pi + 1u], position[pi + 2u]); + let diff = params.center - pos; + let d2 = dot(diff, diff); + let r2 = params.radius * params.radius; + if d2 > r2 || d2 < 0.000001 { return; } + + let d = sqrt(d2); + let dir = diff / d; + + var fall: f32 = 1.0; + let n = 1.0 - d / params.radius; + if params.falloff_mode == 1u { + fall = n; + } else if params.falloff_mode == 2u { + fall = n * n * (3.0 - 2.0 * n); + } else if params.falloff_mode == 3u { + fall = n * n; + } else if params.falloff_mode == 4u { + fall = n * n * n; + } else if params.falloff_mode == 5u { + fall = params.radius / (d + params.radius); + } + + let kick = dir * (params.strength * fall); + velocity[pi] = velocity[pi] + kick.x; + velocity[pi + 1u] = velocity[pi + 1u] + kick.y; + velocity[pi + 2u] = velocity[pi + 2u] + kick.z; +} diff --git a/crates/processing_render/src/particles/kernels/bounds_box.wgsl b/crates/processing_render/src/particles/kernels/bounds_box.wgsl new file mode 100644 index 00000000..65c75bba --- /dev/null +++ b/crates/processing_render/src/particles/kernels/bounds_box.wgsl @@ -0,0 +1,67 @@ +struct Params { + aabb_min: vec3, + soft_strength: f32, + aabb_max: vec3, + max_speed: f32, + mode: u32, + _pad0: u32, + _pad1: u32, + _pad2: u32, +} + +@group(0) @binding(0) var position: array; +@group(0) @binding(1) var velocity: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + if i >= count { return; } + + let pi = i * 3u; + var pos = vec3(position[pi], position[pi + 1u], position[pi + 2u]); + var vel = vec3(velocity[pi], velocity[pi + 1u], velocity[pi + 2u]); + + let lo = params.aabb_min; + let hi = params.aabb_max; + let size = max(hi - lo, vec3(0.000001)); + + if params.mode == 0u { + if pos.x > hi.x { pos.x = hi.x; vel.x = min(vel.x, 0.0); } + if pos.x < lo.x { pos.x = lo.x; vel.x = max(vel.x, 0.0); } + if pos.y > hi.y { pos.y = hi.y; vel.y = min(vel.y, 0.0); } + if pos.y < lo.y { pos.y = lo.y; vel.y = max(vel.y, 0.0); } + if pos.z > hi.z { pos.z = hi.z; vel.z = min(vel.z, 0.0); } + if pos.z < lo.z { pos.z = lo.z; vel.z = max(vel.z, 0.0); } + } else if params.mode == 1u { + if pos.x > hi.x { pos.x = 2.0 * hi.x - pos.x; if vel.x > 0.0 { vel.x = -vel.x; } } + if pos.x < lo.x { pos.x = 2.0 * lo.x - pos.x; if vel.x < 0.0 { vel.x = -vel.x; } } + if pos.y > hi.y { pos.y = 2.0 * hi.y - pos.y; if vel.y > 0.0 { vel.y = -vel.y; } } + if pos.y < lo.y { pos.y = 2.0 * lo.y - pos.y; if vel.y < 0.0 { vel.y = -vel.y; } } + if pos.z > hi.z { pos.z = 2.0 * hi.z - pos.z; if vel.z > 0.0 { vel.z = -vel.z; } } + if pos.z < lo.z { pos.z = 2.0 * lo.z - pos.z; if vel.z < 0.0 { vel.z = -vel.z; } } + } else if params.mode == 2u { + let rel = pos - lo; + pos = lo + (rel - size * floor(rel / size)); + } else { + let over_hi = max(pos - hi, vec3(0.0)); + let over_lo = max(lo - pos, vec3(0.0)); + vel = vel - params.soft_strength * (over_hi - over_lo); + } + + if params.max_speed > 0.0 { + let speed2 = dot(vel, vel); + let cap2 = params.max_speed * params.max_speed; + if speed2 > cap2 { + vel = vel * (params.max_speed * inverseSqrt(speed2)); + } + } + + position[pi] = pos.x; + position[pi + 1u] = pos.y; + position[pi + 2u] = pos.z; + velocity[pi] = vel.x; + velocity[pi + 1u] = vel.y; + velocity[pi + 2u] = vel.z; +} diff --git a/crates/processing_render/src/particles/kernels/bounds_sphere.wgsl b/crates/processing_render/src/particles/kernels/bounds_sphere.wgsl new file mode 100644 index 00000000..6c342eb2 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/bounds_sphere.wgsl @@ -0,0 +1,60 @@ +struct Params { + center: vec3, + radius: f32, + max_speed: f32, + soft_strength: f32, + mode: u32, + _pad: u32, +} + +@group(0) @binding(0) var position: array; +@group(0) @binding(1) var velocity: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + if i >= count { return; } + + let pi = i * 3u; + var pos = vec3(position[pi], position[pi + 1u], position[pi + 2u]); + var vel = vec3(velocity[pi], velocity[pi + 1u], velocity[pi + 2u]); + + let offset = pos - params.center; + let d2 = dot(offset, offset); + let r2 = params.radius * params.radius; + + if d2 > r2 && d2 > 0.000001 { + let d = sqrt(d2); + let normal = offset / d; + + if params.mode == 0u { + pos = params.center + normal * params.radius; + let outward = max(0.0, dot(vel, normal)); + vel = vel - normal * outward; + } else if params.mode == 1u { + let overshoot = d - params.radius; + pos = pos - normal * (2.0 * overshoot); + let outward = max(0.0, dot(vel, normal)); + vel = vel - normal * (2.0 * outward); + } else if params.mode == 3u { + vel = vel - normal * (params.soft_strength * (d - params.radius)); + } + } + + if params.max_speed > 0.0 { + let speed2 = dot(vel, vel); + let cap2 = params.max_speed * params.max_speed; + if speed2 > cap2 { + vel = vel * (params.max_speed * inverseSqrt(speed2)); + } + } + + position[pi] = pos.x; + position[pi + 1u] = pos.y; + position[pi + 2u] = pos.z; + velocity[pi] = vel.x; + velocity[pi + 1u] = vel.y; + velocity[pi + 2u] = vel.z; +} diff --git a/crates/processing_render/src/particles/kernels/drag.wgsl b/crates/processing_render/src/particles/kernels/drag.wgsl new file mode 100644 index 00000000..93737d15 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/drag.wgsl @@ -0,0 +1,29 @@ +struct Params { + damping: f32, + max_speed: f32, +} + +@group(0) @binding(0) var velocity: array; +@group(0) @binding(1) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&velocity) / 3u; + if i >= count { return; } + + let pi = i * 3u; + let v = vec3(velocity[pi], velocity[pi + 1u], velocity[pi + 2u]); + let k = clamp(params.damping, 0.0, 1.0); + var new_v = v * (1.0 - k); + if params.max_speed > 0.0 { + let speed2 = dot(new_v, new_v); + let cap2 = params.max_speed * params.max_speed; + if speed2 > cap2 { + new_v = new_v * (params.max_speed * inverseSqrt(speed2)); + } + } + velocity[pi] = new_v.x; + velocity[pi + 1u] = new_v.y; + velocity[pi + 2u] = new_v.z; +} diff --git a/crates/processing_render/src/particles/kernels/field.wgsl b/crates/processing_render/src/particles/kernels/field.wgsl new file mode 100644 index 00000000..ef079d4d --- /dev/null +++ b/crates/processing_render/src/particles/kernels/field.wgsl @@ -0,0 +1,45 @@ +struct Params { + center: vec3, + radius: f32, + falloff_mode: u32, + _pad0: u32, + _pad1: u32, + _pad2: u32, +} + +@group(0) @binding(0) var position: array; +@group(0) @binding(1) var weight: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + if i >= count { return; } + + let pi = i * 3u; + let pos = vec3(position[pi], position[pi + 1u], position[pi + 2u]); + let diff = pos - params.center; + let d2 = dot(diff, diff); + let r2 = params.radius * params.radius; + + var w: f32 = 0.0; + if d2 < r2 { + let d = sqrt(d2); + let n = 1.0 - d / params.radius; + w = 1.0; + if params.falloff_mode == 1u { + w = n; + } else if params.falloff_mode == 2u { + w = n * n * (3.0 - 2.0 * n); + } else if params.falloff_mode == 3u { + w = n * n; + } else if params.falloff_mode == 4u { + w = n * n * n; + } else if params.falloff_mode == 5u { + w = params.radius / (d + params.radius); + } + } + + weight[i] = w; +} diff --git a/crates/processing_render/src/particles/kernels/flock.wgsl b/crates/processing_render/src/particles/kernels/flock.wgsl new file mode 100644 index 00000000..c2d6311a --- /dev/null +++ b/crates/processing_render/src/particles/kernels/flock.wgsl @@ -0,0 +1,127 @@ +struct Params { + sep_distance: f32, + neighbor_distance: f32, + weight_separation: f32, + weight_alignment: f32, + weight_cohesion: f32, + max_speed: f32, + max_force: f32, + min_speed: f32, +} + +@group(0) @binding(0) var position: array; +@group(0) @binding(1) var velocity: array; +@group(0) @binding(2) var params: Params; + +const TILE: u32 = 64u; +var s_pos: array, 64>; +var s_vel: array, 64>; + +fn limit_mag(v: vec3, m: f32) -> vec3 { + let len2 = dot(v, v); + if len2 > m * m { return v * (m * inverseSqrt(len2)); } + return v; +} + +fn steer_toward(desired: vec3, vel: vec3, max_speed: f32, max_force: f32) -> vec3 { + let m2 = dot(desired, desired); + if m2 < 0.00000001 { return vec3(0.0); } + return limit_mag(desired * (max_speed * inverseSqrt(m2)) - vel, max_force); +} + +@compute @workgroup_size(64) +fn main( + @builtin(global_invocation_id) gid: vec3, + @builtin(local_invocation_id) lid: vec3, +) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + let alive = i < count; + + let sep_d2 = params.sep_distance * params.sep_distance; + let neighbor_d2 = params.neighbor_distance * params.neighbor_distance; + + var pos = vec3(0.0); + var vel = vec3(0.0); + if alive { + let pi = i * 3u; + pos = vec3(position[pi], position[pi + 1u], position[pi + 2u]); + vel = vec3(velocity[pi], velocity[pi + 1u], velocity[pi + 2u]); + } + + var sep_steer = vec3(0.0); + var sep_count = 0u; + var ali_sum = vec3(0.0); + var coh_sum = vec3(0.0); + var flock_count = 0u; + + let num_tiles = (count + TILE - 1u) / TILE; + + for (var t = 0u; t < num_tiles; t++) { + let j = t * TILE + lid.x; + if j < count { + let pj = j * 3u; + s_pos[lid.x] = vec3(position[pj], position[pj + 1u], position[pj + 2u]); + s_vel[lid.x] = vec3(velocity[pj], velocity[pj + 1u], velocity[pj + 2u]); + } + workgroupBarrier(); + + if alive { + let tile_end = min(TILE, count - t * TILE); + for (var k = 0u; k < tile_end; k++) { + let global_j = t * TILE + k; + if global_j == i { continue; } + + let diff = pos - s_pos[k]; + let d2 = dot(diff, diff); + + if d2 > 0.000001 && d2 < neighbor_d2 { + if d2 < sep_d2 { + sep_steer += diff / d2; + sep_count += 1u; + } + ali_sum += s_vel[k]; + coh_sum += diff; + flock_count += 1u; + } + } + } + workgroupBarrier(); + } + + if !alive { return; } + + var force = vec3(0.0); + if sep_count > 0u { + force += steer_toward(sep_steer / f32(sep_count), vel, + params.max_speed, params.max_force) * params.weight_separation; + } + if flock_count > 0u { + force += steer_toward(ali_sum / f32(flock_count), vel, + params.max_speed, params.max_force) * params.weight_alignment; + force += steer_toward(-coh_sum / f32(flock_count), vel, + params.max_speed, params.max_force) * params.weight_cohesion; + } + + var new_vel = vel + force; + let speed2 = dot(new_vel, new_vel); + let max_speed_sq = params.max_speed * params.max_speed; + if speed2 > max_speed_sq { + new_vel = new_vel * (params.max_speed * inverseSqrt(speed2)); + } else if params.min_speed > 0.0 { + let min_speed_sq = params.min_speed * params.min_speed; + if speed2 < min_speed_sq && speed2 > 0.0 { + new_vel = new_vel * sqrt(min_speed_sq / speed2); + } + } + + let new_pos = pos + new_vel; + + let pi = i * 3u; + position[pi] = new_pos.x; + position[pi + 1u] = new_pos.y; + position[pi + 2u] = new_pos.z; + velocity[pi] = new_vel.x; + velocity[pi + 1u] = new_vel.y; + velocity[pi + 2u] = new_vel.z; +} diff --git a/crates/processing_render/src/particles/kernels/force.wgsl b/crates/processing_render/src/particles/kernels/force.wgsl new file mode 100644 index 00000000..5a9f55f0 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/force.wgsl @@ -0,0 +1,23 @@ +struct Params { + direction: vec3, + strength: f32, +} + +@group(0) @binding(0) var velocity: array; +@group(0) @binding(1) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&velocity) / 3u; + if i >= count { return; } + + let len2 = dot(params.direction, params.direction); + if len2 < 0.000001 || params.strength == 0.0 { return; } + let kick = params.direction * (params.strength * inverseSqrt(len2)); + + let pi = i * 3u; + velocity[pi] = velocity[pi] + kick.x; + velocity[pi + 1u] = velocity[pi + 1u] + kick.y; + velocity[pi + 2u] = velocity[pi + 2u] + kick.z; +} diff --git a/crates/processing_render/src/particles/kernels/impulse.wgsl b/crates/processing_render/src/particles/kernels/impulse.wgsl new file mode 100644 index 00000000..3368d748 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/impulse.wgsl @@ -0,0 +1,53 @@ +struct Params { + center: vec3, + radius: f32, + position_kick: f32, + velocity_kick: f32, + falloff_mode: u32, + _pad: u32, +} + +@group(0) @binding(0) var position: array; +@group(0) @binding(1) var velocity: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + if i >= count { return; } + + let pi = i * 3u; + let pos = vec3(position[pi], position[pi + 1u], position[pi + 2u]); + let diff = pos - params.center; + let d2 = dot(diff, diff); + let r2 = params.radius * params.radius; + if d2 > r2 || d2 < 0.000001 { return; } + + let d = sqrt(d2); + let dir = diff / d; + + var fall: f32 = 1.0; + let n = 1.0 - d / params.radius; + if params.falloff_mode == 1u { + fall = n; + } else if params.falloff_mode == 2u { + fall = n * n * (3.0 - 2.0 * n); + } else if params.falloff_mode == 3u { + fall = n * n; + } else if params.falloff_mode == 4u { + fall = n * n * n; + } else if params.falloff_mode == 5u { + fall = params.radius / (d + params.radius); + } + + let pos_push = dir * (params.position_kick * fall); + let vel_push = dir * (params.velocity_kick * fall); + + position[pi] = position[pi] + pos_push.x; + position[pi + 1u] = position[pi + 1u] + pos_push.y; + position[pi + 2u] = position[pi + 2u] + pos_push.z; + velocity[pi] = velocity[pi] + vel_push.x; + velocity[pi + 1u] = velocity[pi + 1u] + vel_push.y; + velocity[pi + 2u] = velocity[pi + 2u] + vel_push.z; +} diff --git a/crates/processing_render/src/particles/kernels/integrate.wgsl b/crates/processing_render/src/particles/kernels/integrate.wgsl new file mode 100644 index 00000000..22a82afd --- /dev/null +++ b/crates/processing_render/src/particles/kernels/integrate.wgsl @@ -0,0 +1,19 @@ +struct Params { + dt: f32, +} + +@group(0) @binding(0) var velocity: array; +@group(0) @binding(1) var position: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + if i >= count { return; } + + let pi = i * 3u; + position[pi] = position[pi] + velocity[pi] * params.dt; + position[pi + 1u] = position[pi + 1u] + velocity[pi + 1u] * params.dt; + position[pi + 2u] = position[pi + 2u] + velocity[pi + 2u] * params.dt; +} diff --git a/crates/processing_render/src/particles/kernels/mod.rs b/crates/processing_render/src/particles/kernels/mod.rs index 8f64f699..c0068373 100644 --- a/crates/processing_render/src/particles/kernels/mod.rs +++ b/crates/processing_render/src/particles/kernels/mod.rs @@ -2,16 +2,323 @@ //! assets and dispatched via `particles_apply`. use bevy::asset::embedded_asset; +use bevy::mesh::VertexAttributeValues; use bevy::prelude::*; +use processing_core::app_mut; +use processing_core::error::{self, ProcessingError}; + +use crate::geometry::{BuiltinAttributes, Geometry}; +use crate::shader_value::ShaderValue; +use crate::{compute_create, compute_set, shader_load}; + +#[derive(Component, Default, Clone)] +pub struct KernelRequires(pub Vec); + +fn set_requires(compute: Entity, names: &[&str]) -> error::Result<()> { + app_mut(|app| { + let world = app.world_mut(); + let attrs: Vec = { + let builtins = world.resource::(); + names.iter().filter_map(|n| builtins.by_name(n)).collect() + }; + world.entity_mut(compute).insert(KernelRequires(attrs)); + Ok(()) + }) +} + pub struct ParticlesKernelsPlugin; impl Plugin for ParticlesKernelsPlugin { fn build(&self, app: &mut App) { embedded_asset!(app, "noise.wgsl"); embedded_asset!(app, "transform.wgsl"); + embedded_asset!(app, "attract.wgsl"); + embedded_asset!(app, "drag.wgsl"); + embedded_asset!(app, "force.wgsl"); + embedded_asset!(app, "integrate.wgsl"); + embedded_asset!(app, "age.wgsl"); + embedded_asset!(app, "vortex.wgsl"); + embedded_asset!(app, "bounds_sphere.wgsl"); + embedded_asset!(app, "bounds_box.wgsl"); + embedded_asset!(app, "impulse.wgsl"); + embedded_asset!(app, "flock.wgsl"); + embedded_asset!(app, "orient.wgsl"); + embedded_asset!(app, "field.wgsl"); + embedded_asset!(app, "attr_linear.wgsl"); + embedded_asset!(app, "attr_combine.wgsl"); + embedded_asset!(app, "attr_mix.wgsl"); + embedded_asset!(app, "attr_lookup1d.wgsl"); + embedded_asset!(app, "attr_lookup2d.wgsl"); + embedded_asset!(app, "scatter_surface.wgsl"); + embedded_asset!(app, "scatter_volume.wgsl"); } } -pub const NOISE_PATH: &str = "embedded://processing_render/particles/kernels/noise.wgsl"; -pub const TRANSFORM_PATH: &str = "embedded://processing_render/particles/kernels/transform.wgsl"; +pub const FALLOFF_CONST: u32 = 0; +pub const FALLOFF_LINEAR: u32 = 1; +pub const FALLOFF_SMOOTHSTEP: u32 = 2; +pub const FALLOFF_QUADRATIC: u32 = 3; +pub const FALLOFF_CUBIC: u32 = 4; +pub const FALLOFF_INVERSE: u32 = 5; + +pub const BOUNDS_CLAMP: u32 = 0; +pub const BOUNDS_REFLECT: u32 = 1; +pub const BOUNDS_WRAP: u32 = 2; +pub const BOUNDS_SOFT: u32 = 3; + +pub const COMBINE_ADD: u32 = 0; +pub const COMBINE_SUB: u32 = 1; +pub const COMBINE_MUL: u32 = 2; +pub const COMBINE_DIV: u32 = 3; +pub const COMBINE_MIN: u32 = 4; +pub const COMBINE_MAX: u32 = 5; +pub const COMBINE_POW: u32 = 6; + +pub fn particles_kernel_noise() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/noise.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position"])?; + compute_set(entity, "divergence_free", ShaderValue::UInt(0))?; + Ok(entity) +} + +pub fn particles_kernel_transform() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/transform.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position"])?; + compute_set(entity, "translate", ShaderValue::Float3([0.0; 3]))?; + compute_set( + entity, + "rotation_axis", + ShaderValue::Float3([0.0, 1.0, 0.0]), + )?; + compute_set(entity, "rotation_angle", ShaderValue::Float(0.0))?; + compute_set(entity, "scale", ShaderValue::Float3([1.0, 1.0, 1.0]))?; + Ok(entity) +} + +pub fn particles_kernel_attract() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/attract.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position", "velocity"])?; + compute_set(entity, "center", ShaderValue::Float3([0.0; 3]))?; + compute_set(entity, "strength", ShaderValue::Float(0.0))?; + compute_set(entity, "radius", ShaderValue::Float(1.0))?; + compute_set(entity, "falloff_mode", ShaderValue::UInt(FALLOFF_LINEAR))?; + Ok(entity) +} + +pub fn particles_kernel_drag() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/drag.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["velocity"])?; + compute_set(entity, "damping", ShaderValue::Float(0.0))?; + compute_set(entity, "max_speed", ShaderValue::Float(0.0))?; + Ok(entity) +} + +pub fn particles_kernel_force() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/force.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["velocity"])?; + compute_set(entity, "direction", ShaderValue::Float3([0.0, -1.0, 0.0]))?; + compute_set(entity, "strength", ShaderValue::Float(0.0))?; + Ok(entity) +} + +pub fn particles_kernel_integrate() -> error::Result { + let shader = + shader_load("embedded://processing_render/particles/kernels/integrate.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position", "velocity"])?; + compute_set(entity, "dt", ShaderValue::Float(1.0))?; + Ok(entity) +} + +pub fn particles_kernel_age() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/age.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["age", "life"])?; + compute_set(entity, "dt", ShaderValue::Float(1.0))?; + Ok(entity) +} + +pub fn particles_kernel_vortex() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/vortex.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position", "velocity"])?; + compute_set(entity, "center", ShaderValue::Float3([0.0; 3]))?; + compute_set(entity, "axis", ShaderValue::Float3([0.0, 1.0, 0.0]))?; + compute_set(entity, "strength", ShaderValue::Float(0.0))?; + compute_set(entity, "radius", ShaderValue::Float(1.0))?; + compute_set(entity, "falloff_mode", ShaderValue::UInt(FALLOFF_LINEAR))?; + Ok(entity) +} + +pub fn particles_kernel_bounds_sphere() -> error::Result { + let shader = + shader_load("embedded://processing_render/particles/kernels/bounds_sphere.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position", "velocity"])?; + compute_set(entity, "center", ShaderValue::Float3([0.0; 3]))?; + compute_set(entity, "radius", ShaderValue::Float(1.0))?; + compute_set(entity, "mode", ShaderValue::UInt(BOUNDS_SOFT))?; + compute_set(entity, "soft_strength", ShaderValue::Float(0.01))?; + compute_set(entity, "max_speed", ShaderValue::Float(0.0))?; + Ok(entity) +} + +pub fn particles_kernel_bounds_box() -> error::Result { + let shader = + shader_load("embedded://processing_render/particles/kernels/bounds_box.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position", "velocity"])?; + compute_set(entity, "aabb_min", ShaderValue::Float3([-1.0, -1.0, -1.0]))?; + compute_set(entity, "aabb_max", ShaderValue::Float3([1.0, 1.0, 1.0]))?; + compute_set(entity, "mode", ShaderValue::UInt(BOUNDS_SOFT))?; + compute_set(entity, "soft_strength", ShaderValue::Float(0.01))?; + compute_set(entity, "max_speed", ShaderValue::Float(0.0))?; + Ok(entity) +} + +pub fn particles_kernel_bounds_geometry(geometry_entity: Entity) -> error::Result { + let (aabb_min, aabb_max) = app_mut(|app| { + app.world_mut() + .run_system_cached_with(extract_geometry_aabb, geometry_entity) + .unwrap() + })?; + let entity = particles_kernel_bounds_box()?; + compute_set(entity, "aabb_min", ShaderValue::Float3(aabb_min))?; + compute_set(entity, "aabb_max", ShaderValue::Float3(aabb_max))?; + Ok(entity) +} + +fn extract_geometry_aabb( + In(geom_entity): In, + geometries: Query<&Geometry>, + meshes: Res>, +) -> error::Result<([f32; 3], [f32; 3])> { + let geom = geometries + .get(geom_entity) + .map_err(|_| ProcessingError::GeometryNotFound)?; + let mesh = meshes + .get(&geom.handle) + .ok_or(ProcessingError::GeometryNotFound)?; + let positions = match mesh.attribute(Mesh::ATTRIBUTE_POSITION) { + Some(VertexAttributeValues::Float32x3(p)) => p, + _ => { + return Err(ProcessingError::InvalidArgument( + "bounds geometry has no Float32x3 position attribute".to_string(), + )); + } + }; + if positions.is_empty() { + return Err(ProcessingError::InvalidArgument( + "bounds geometry has no vertices".to_string(), + )); + } + let mut min = Vec3::splat(f32::INFINITY); + let mut max = Vec3::splat(f32::NEG_INFINITY); + for p in positions { + let v = Vec3::from_array(*p); + min = min.min(v); + max = max.max(v); + } + Ok((min.to_array(), max.to_array())) +} + +pub fn particles_kernel_impulse() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/impulse.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position", "velocity"])?; + compute_set(entity, "center", ShaderValue::Float3([0.0; 3]))?; + compute_set(entity, "radius", ShaderValue::Float(1.0))?; + compute_set(entity, "position_kick", ShaderValue::Float(0.0))?; + compute_set(entity, "velocity_kick", ShaderValue::Float(0.0))?; + compute_set(entity, "falloff_mode", ShaderValue::UInt(FALLOFF_SMOOTHSTEP))?; + Ok(entity) +} + +pub fn particles_kernel_flock() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/flock.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position", "velocity"])?; + compute_set(entity, "sep_distance", ShaderValue::Float(1.2))?; + compute_set(entity, "neighbor_distance", ShaderValue::Float(2.5))?; + compute_set(entity, "weight_separation", ShaderValue::Float(1.5))?; + compute_set(entity, "weight_alignment", ShaderValue::Float(1.0))?; + compute_set(entity, "weight_cohesion", ShaderValue::Float(1.0))?; + compute_set(entity, "max_speed", ShaderValue::Float(0.1))?; + compute_set(entity, "max_force", ShaderValue::Float(0.003))?; + compute_set(entity, "min_speed", ShaderValue::Float(0.02))?; + Ok(entity) +} + +pub fn particles_kernel_orient() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/orient.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["velocity", "rotation"])?; + compute_set(entity, "forward", ShaderValue::Float3([0.0, 0.0, 1.0]))?; + compute_set(entity, "up", ShaderValue::Float3([0.0, 1.0, 0.0]))?; + Ok(entity) +} + +pub fn particles_kernel_field() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/field.wgsl")?; + let entity = compute_create(shader)?; + set_requires(entity, &["position"])?; + compute_set(entity, "center", ShaderValue::Float3([0.0; 3]))?; + compute_set(entity, "radius", ShaderValue::Float(1.0))?; + compute_set(entity, "falloff_mode", ShaderValue::UInt(FALLOFF_SMOOTHSTEP))?; + Ok(entity) +} + +pub fn particles_kernel_attr_linear() -> error::Result { + let shader = + shader_load("embedded://processing_render/particles/kernels/attr_linear.wgsl")?; + let entity = compute_create(shader)?; + compute_set(entity, "scale", ShaderValue::Float(1.0))?; + compute_set(entity, "offset", ShaderValue::Float(0.0))?; + Ok(entity) +} + +pub fn particles_kernel_attr_combine() -> error::Result { + let shader = + shader_load("embedded://processing_render/particles/kernels/attr_combine.wgsl")?; + let entity = compute_create(shader)?; + compute_set(entity, "op", ShaderValue::UInt(COMBINE_ADD))?; + compute_set(entity, "b_scale", ShaderValue::Float(1.0))?; + compute_set(entity, "b_offset", ShaderValue::Float(0.0))?; + Ok(entity) +} + +pub fn particles_kernel_attr_mix() -> error::Result { + let shader = shader_load("embedded://processing_render/particles/kernels/attr_mix.wgsl")?; + let entity = compute_create(shader)?; + compute_set(entity, "t_scale", ShaderValue::Float(1.0))?; + compute_set(entity, "t_offset", ShaderValue::Float(0.0))?; + compute_set(entity, "t_clamp", ShaderValue::UInt(1))?; + Ok(entity) +} + +pub fn particles_kernel_attr_lookup1d() -> error::Result { + let shader = + shader_load("embedded://processing_render/particles/kernels/attr_lookup1d.wgsl")?; + let entity = compute_create(shader)?; + compute_set(entity, "scale", ShaderValue::Float(1.0))?; + compute_set(entity, "offset", ShaderValue::Float(0.0))?; + Ok(entity) +} + +pub fn particles_kernel_attr_lookup2d() -> error::Result { + let shader = + shader_load("embedded://processing_render/particles/kernels/attr_lookup2d.wgsl")?; + let entity = compute_create(shader)?; + compute_set(entity, "u_scale", ShaderValue::Float(1.0))?; + compute_set(entity, "u_offset", ShaderValue::Float(0.0))?; + compute_set(entity, "v_scale", ShaderValue::Float(1.0))?; + compute_set(entity, "v_offset", ShaderValue::Float(0.0))?; + compute_set(entity, "color_scale", ShaderValue::Float(1.0))?; + Ok(entity) +} diff --git a/crates/processing_render/src/particles/kernels/noise.wgsl b/crates/processing_render/src/particles/kernels/noise.wgsl index 8aec001b..65053475 100644 --- a/crates/processing_render/src/particles/kernels/noise.wgsl +++ b/crates/processing_render/src/particles/kernels/noise.wgsl @@ -2,7 +2,7 @@ struct Params { scale: f32, strength: f32, time: f32, - _pad: f32, + divergence_free: u32, } @group(0) @binding(0) var position: array; @@ -42,6 +42,27 @@ fn noise3(p: vec3) -> vec3 { ) * 2.0 - 1.0; } +fn curl_noise(p: vec3, eps: f32) -> vec3 { + let dx = vec3(eps, 0.0, 0.0); + let dy = vec3(0.0, eps, 0.0); + let dz = vec3(0.0, 0.0, eps); + + let n_xp = noise3(p + dx); let n_xm = noise3(p - dx); + let n_yp = noise3(p + dy); let n_ym = noise3(p - dy); + let n_zp = noise3(p + dz); let n_zm = noise3(p - dz); + + let inv = 1.0 / (2.0 * eps); + let dn_dx = (n_xp - n_xm) * inv; + let dn_dy = (n_yp - n_ym) * inv; + let dn_dz = (n_zp - n_zm) * inv; + + return vec3( + dn_dy.z - dn_dz.y, + dn_dz.x - dn_dx.z, + dn_dx.y - dn_dy.x, + ); +} + @compute @workgroup_size(64) fn main(@builtin(global_invocation_id) gid: vec3) { let i = gid.x; @@ -55,7 +76,13 @@ fn main(@builtin(global_invocation_id) gid: vec3) { position[i * 3u + 2u], ); let sample = p * params.scale + vec3(params.time, params.time * 0.7, params.time * 1.3); - let n = noise3(sample); + var n: vec3; + if params.divergence_free == 1u { + let eps = 0.5 / max(params.scale, 1.0); + n = curl_noise(sample, eps); + } else { + n = noise3(sample); + } let new_p = p + n * params.strength; position[i * 3u + 0u] = new_p.x; position[i * 3u + 1u] = new_p.y; diff --git a/crates/processing_render/src/particles/kernels/orient.wgsl b/crates/processing_render/src/particles/kernels/orient.wgsl new file mode 100644 index 00000000..6260cc64 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/orient.wgsl @@ -0,0 +1,115 @@ +struct Params { + forward: vec3, + _pad0: f32, + up: vec3, + _pad1: f32, +} + +@group(0) @binding(0) var velocity: array; +@group(0) @binding(1) var rotation: array; +@group(0) @binding(2) var params: Params; + +fn quat_from_to(src: vec3, dst: vec3) -> vec4 { + let d = dot(src, dst); + if d > 0.999999 { + return vec4(0.0, 0.0, 0.0, 1.0); + } + if d < -0.999999 { + var axis = cross(src, vec3(1.0, 0.0, 0.0)); + if dot(axis, axis) < 0.001 { + axis = cross(src, vec3(0.0, 1.0, 0.0)); + } + return vec4(normalize(axis), 0.0); + } + let axis = cross(src, dst); + return normalize(vec4(axis, 1.0 + d)); +} + +fn mat3_to_quat(m: mat3x3) -> vec4 { + let trace = m[0][0] + m[1][1] + m[2][2]; + if trace > 0.0 { + let s = sqrt(trace + 1.0) * 2.0; + return vec4( + (m[1][2] - m[2][1]) / s, + (m[2][0] - m[0][2]) / s, + (m[0][1] - m[1][0]) / s, + 0.25 * s, + ); + } + if m[0][0] > m[1][1] && m[0][0] > m[2][2] { + let s = sqrt(1.0 + m[0][0] - m[1][1] - m[2][2]) * 2.0; + return vec4( + 0.25 * s, + (m[1][0] + m[0][1]) / s, + (m[2][0] + m[0][2]) / s, + (m[1][2] - m[2][1]) / s, + ); + } + if m[1][1] > m[2][2] { + let s = sqrt(1.0 + m[1][1] - m[0][0] - m[2][2]) * 2.0; + return vec4( + (m[1][0] + m[0][1]) / s, + 0.25 * s, + (m[2][1] + m[1][2]) / s, + (m[2][0] - m[0][2]) / s, + ); + } + let s = sqrt(1.0 + m[2][2] - m[0][0] - m[1][1]) * 2.0; + return vec4( + (m[2][0] + m[0][2]) / s, + (m[2][1] + m[1][2]) / s, + 0.25 * s, + (m[0][1] - m[1][0]) / s, + ); +} + +fn write_quat(i: u32, q: vec4) { + let ri = i * 4u; + rotation[ri] = q.x; + rotation[ri + 1u] = q.y; + rotation[ri + 2u] = q.z; + rotation[ri + 3u] = q.w; +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&velocity) / 3u; + if i >= count { return; } + + let pi = i * 3u; + let v = vec3(velocity[pi], velocity[pi + 1u], velocity[pi + 2u]); + let v_len2 = dot(v, v); + if v_len2 < 0.000001 { return; } + let vel_dir = v * inverseSqrt(v_len2); + + let fwd_len2 = dot(params.forward, params.forward); + if fwd_len2 < 0.000001 { return; } + let fwd = params.forward * inverseSqrt(fwd_len2); + + let up_param_len2 = dot(params.up, params.up); + if up_param_len2 < 0.000001 { + write_quat(i, quat_from_to(fwd, vel_dir)); + return; + } + + var up_local = params.up - dot(params.up, fwd) * fwd; + let up_local_len2 = dot(up_local, up_local); + var up_target = params.up - dot(params.up, vel_dir) * vel_dir; + let up_target_len2 = dot(up_target, up_target); + if up_local_len2 < 0.000001 || up_target_len2 < 0.000001 { + write_quat(i, quat_from_to(fwd, vel_dir)); + return; + } + up_local = up_local * inverseSqrt(up_local_len2); + up_target = up_target * inverseSqrt(up_target_len2); + + let right_local = cross(fwd, up_local); + let right_target = cross(vel_dir, up_target); + + let local_col = mat3x3(fwd, up_local, right_local); + let target_col = mat3x3(vel_dir, up_target, right_target); + let r = target_col * transpose(local_col); + + write_quat(i, mat3_to_quat(r)); +} diff --git a/crates/processing_render/src/particles/kernels/scatter_surface.wgsl b/crates/processing_render/src/particles/kernels/scatter_surface.wgsl new file mode 100644 index 00000000..a9fa3b60 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/scatter_surface.wgsl @@ -0,0 +1,99 @@ +struct Params { + face_count: u32, + seed: u32, + _pad0: u32, + _pad1: u32, +} + +struct EmitRange { + emit_base: u32, + emit_count: u32, + emit_capacity: u32, + _pad: u32, +} + +@group(0) @binding(0) var source_position: array; +@group(0) @binding(1) var source_indices: array; +@group(0) @binding(2) var cdf: array; +@group(0) @binding(3) var position: array; +@group(0) @binding(4) var scale: array; +@group(0) @binding(5) var age: array; +@group(0) @binding(6) var life: array; +@group(0) @binding(7) var params: Params; +@group(0) @binding(8) var emit_range: EmitRange; + +fn hash(n: u32) -> u32 { + var x = n; + x = (x ^ 61u) ^ (x >> 16u); + x = x + (x << 3u); + x = x ^ (x >> 4u); + x = x * 0x27d4eb2du; + x = x ^ (x >> 15u); + return x; +} + +fn hash_unit(n: u32) -> f32 { + return f32(hash(n)) / f32(0xffffffffu); +} + +fn cdf_search(u: f32) -> u32 { + var lo: u32 = 0u; + var hi: u32 = params.face_count; + loop { + if lo >= hi { break; } + let mid = (lo + hi) >> 1u; + if cdf[mid] < u { + lo = mid + 1u; + } else { + hi = mid; + } + } + return min(lo, params.face_count - 1u); +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let local_i = gid.x; + if local_i >= emit_range.emit_count { return; } + let base = emit_range.emit_base; + let cap = emit_range.emit_capacity; + let slot = (base + local_i) % cap; + let seed = params.seed ^ (base + local_i); + + let u01 = hash_unit(seed * 7u + 13u); + let face = cdf_search(u01); + + let i0 = source_indices[face * 3u + 0u]; + let i1 = source_indices[face * 3u + 1u]; + let i2 = source_indices[face * 3u + 2u]; + + let p0 = vec3( + source_position[i0 * 3u + 0u], + source_position[i0 * 3u + 1u], + source_position[i0 * 3u + 2u], + ); + let p1 = vec3( + source_position[i1 * 3u + 0u], + source_position[i1 * 3u + 1u], + source_position[i1 * 3u + 2u], + ); + let p2 = vec3( + source_position[i2 * 3u + 0u], + source_position[i2 * 3u + 1u], + source_position[i2 * 3u + 2u], + ); + + var u = hash_unit(seed * 31u + 23u); + var v = hash_unit(seed * 47u + 29u); + if u + v > 1.0 { u = 1.0 - u; v = 1.0 - v; } + let p = (1.0 - u - v) * p0 + u * p1 + v * p2; + + position[slot * 3u + 0u] = p.x; + position[slot * 3u + 1u] = p.y; + position[slot * 3u + 2u] = p.z; + scale[slot * 3u + 0u] = 1.0; + scale[slot * 3u + 1u] = 1.0; + scale[slot * 3u + 2u] = 1.0; + age[slot] = 0.0; + life[slot] = 1.0; +} diff --git a/crates/processing_render/src/particles/kernels/scatter_volume.wgsl b/crates/processing_render/src/particles/kernels/scatter_volume.wgsl new file mode 100644 index 00000000..47f0b2de --- /dev/null +++ b/crates/processing_render/src/particles/kernels/scatter_volume.wgsl @@ -0,0 +1,119 @@ +struct Params { + aabb_min: vec4, + aabb_max: vec4, + face_count: u32, + max_attempts: u32, + seed: u32, + _pad: u32, +} + +struct EmitRange { + emit_base: u32, + emit_count: u32, + emit_capacity: u32, + _pad: u32, +} + +@group(0) @binding(0) var source_position: array; +@group(0) @binding(1) var source_indices: array; +@group(0) @binding(2) var position: array; +@group(0) @binding(3) var scale: array; +@group(0) @binding(4) var age: array; +@group(0) @binding(5) var life: array; +@group(0) @binding(6) var params: Params; +@group(0) @binding(7) var emit_range: EmitRange; + +fn hash(n: u32) -> u32 { + var x = n; + x = (x ^ 61u) ^ (x >> 16u); + x = x + (x << 3u); + x = x ^ (x >> 4u); + x = x * 0x27d4eb2du; + x = x ^ (x >> 15u); + return x; +} + +fn hash_unit(n: u32) -> f32 { + return f32(hash(n)) / f32(0xffffffffu); +} + +fn fetch_vertex(i: u32) -> vec3 { + return vec3( + source_position[i * 3u + 0u], + source_position[i * 3u + 1u], + source_position[i * 3u + 2u], + ); +} + +fn ray_triangle( + ro: vec3, rd: vec3, + p0: vec3, p1: vec3, p2: vec3, +) -> f32 { + let e1 = p1 - p0; + let e2 = p2 - p0; + let h = cross(rd, e2); + let a = dot(e1, h); + if abs(a) < 1e-7 { return -1.0; } + let f = 1.0 / a; + let s = ro - p0; + let u = f * dot(s, h); + if u < 0.0 || u > 1.0 { return -1.0; } + let q = cross(s, e1); + let v = f * dot(rd, q); + if v < 0.0 || (u + v) > 1.0 { return -1.0; } + return f * dot(e2, q); +} + +fn point_inside(p: vec3) -> bool { + let rd = normalize(vec3(0.5773, 0.5774, 0.5775)); + var hits = 0u; + for (var f = 0u; f < params.face_count; f = f + 1u) { + let i0 = source_indices[f * 3u + 0u]; + let i1 = source_indices[f * 3u + 1u]; + let i2 = source_indices[f * 3u + 2u]; + let t = ray_triangle(p, rd, fetch_vertex(i0), fetch_vertex(i1), fetch_vertex(i2)); + if t >= 0.0 { + hits = hits + 1u; + } + } + return (hits & 1u) == 1u; +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let local_i = gid.x; + if local_i >= emit_range.emit_count { return; } + let base = emit_range.emit_base; + let cap = emit_range.emit_capacity; + let slot = (base + local_i) % cap; + let seed_base = params.seed ^ (base + local_i); + + let lo = params.aabb_min.xyz; + let hi = params.aabb_max.xyz; + + var p = vec3(0.0); + var found = false; + for (var attempt = 0u; attempt < params.max_attempts; attempt = attempt + 1u) { + let s = seed_base ^ (attempt * 0x9e3779b9u); + p = vec3( + mix(lo.x, hi.x, hash_unit(s * 7u + 13u)), + mix(lo.y, hi.y, hash_unit(s * 31u + 23u)), + mix(lo.z, hi.z, hash_unit(s * 47u + 29u)), + ); + if point_inside(p) { + found = true; + break; + } + } + + if !found { return; } + + position[slot * 3u + 0u] = p.x; + position[slot * 3u + 1u] = p.y; + position[slot * 3u + 2u] = p.z; + scale[slot * 3u + 0u] = 1.0; + scale[slot * 3u + 1u] = 1.0; + scale[slot * 3u + 2u] = 1.0; + age[slot] = 0.0; + life[slot] = 1.0; +} diff --git a/crates/processing_render/src/particles/kernels/vortex.wgsl b/crates/processing_render/src/particles/kernels/vortex.wgsl new file mode 100644 index 00000000..59ca61e4 --- /dev/null +++ b/crates/processing_render/src/particles/kernels/vortex.wgsl @@ -0,0 +1,57 @@ +struct Params { + center: vec3, + _pad0: f32, + axis: vec3, + _pad1: f32, + strength: f32, + radius: f32, + falloff_mode: u32, + _pad2: u32, +} + +@group(0) @binding(0) var position: array; +@group(0) @binding(1) var velocity: array; +@group(0) @binding(2) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&position) / 3u; + if i >= count { return; } + + let axis_len2 = dot(params.axis, params.axis); + if axis_len2 < 0.000001 { return; } + let axis = params.axis * inverseSqrt(axis_len2); + + let pi = i * 3u; + let pos = vec3(position[pi], position[pi + 1u], position[pi + 2u]); + let offset = pos - params.center; + + let parallel = dot(offset, axis) * axis; + let radial = offset - parallel; + let r2 = dot(radial, radial); + let max_r2 = params.radius * params.radius; + if r2 > max_r2 || r2 < 0.000001 { return; } + + let r = sqrt(r2); + let tangent = cross(axis, radial / r); + + var fall: f32 = 1.0; + let n = 1.0 - r / params.radius; + if params.falloff_mode == 1u { + fall = n; + } else if params.falloff_mode == 2u { + fall = n * n * (3.0 - 2.0 * n); + } else if params.falloff_mode == 3u { + fall = n * n; + } else if params.falloff_mode == 4u { + fall = n * n * n; + } else if params.falloff_mode == 5u { + fall = params.radius / (r + params.radius); + } + + let kick = tangent * (params.strength * fall); + velocity[pi] = velocity[pi] + kick.x; + velocity[pi + 1u] = velocity[pi + 1u] + kick.y; + velocity[pi + 2u] = velocity[pi + 2u] + kick.z; +} diff --git a/crates/processing_render/src/particles/material.rs b/crates/processing_render/src/particles/material.rs index 2518fcf3..90b55a98 100644 --- a/crates/processing_render/src/particles/material.rs +++ b/crates/processing_render/src/particles/material.rs @@ -1,13 +1,17 @@ -//! per-particle albedo on top of `StandardMaterial`. the `unlit` flag on the -//! base material toggles between lit and unlit; `apply_pbr_lighting` -//! short-circuits when set. - use std::ops::Deref; use bevy::asset::embedded_asset; -use bevy::pbr::{ExtendedMaterial, MaterialExtension, MaterialPlugin}; +use bevy::material::specialize::SpecializedMeshPipelineError; +use bevy::pbr::{ + ExtendedMaterial, MaterialExtension, MaterialExtensionKey, MaterialExtensionPipeline, + MaterialPlugin, +}; use bevy::prelude::*; -use bevy::render::{render_resource::AsBindGroup, storage::ShaderBuffer}; +use bevy::render::{ + mesh::MeshVertexBufferLayoutRef, + render_resource::{AsBindGroup, RenderPipelineDescriptor}, + storage::ShaderBuffer, +}; use bevy::shader::ShaderRef; use crate::render::material::UntypedMaterial; @@ -23,10 +27,28 @@ impl Plugin for ParticlesMaterialPlugin { pub type ParticlesMaterial = ExtendedMaterial; +#[derive(Copy, Clone, PartialEq, Eq, Hash)] +pub struct ParticlesExtensionKey { + pub has_albedo: bool, + pub has_emissive: bool, +} + +impl From<&ParticlesExtension> for ParticlesExtensionKey { + fn from(ext: &ParticlesExtension) -> Self { + Self { + has_albedo: ext.colors.is_some(), + has_emissive: ext.emissive_colors.is_some(), + } + } +} + #[derive(Asset, AsBindGroup, Reflect, Debug, Clone)] +#[bind_group_data(ParticlesExtensionKey)] pub struct ParticlesExtension { #[storage(100, read_only)] - pub colors: Handle, + pub colors: Option>, + #[storage(101, read_only)] + pub emissive_colors: Option>, } impl MaterialExtension for ParticlesExtension { @@ -37,10 +59,25 @@ impl MaterialExtension for ParticlesExtension { fn deferred_fragment_shader() -> ShaderRef { "embedded://processing_render/particles/particles.wgsl".into() } + + fn specialize( + _pipeline: &MaterialExtensionPipeline, + descriptor: &mut RenderPipelineDescriptor, + _layout: &MeshVertexBufferLayoutRef, + key: MaterialExtensionKey, + ) -> Result<(), SpecializedMeshPipelineError> { + if let Some(ref mut fragment) = descriptor.fragment { + if key.bind_group_data.has_albedo { + fragment.shader_defs.push("HAS_COLORS".into()); + } + if key.bind_group_data.has_emissive { + fragment.shader_defs.push("HAS_EMISSIVE_COLORS".into()); + } + } + Ok(()) + } } -/// promote `UntypedMaterial(handle)` to `MeshMaterial3d` -/// where the handle's type matches. pub fn add_particles_materials(mut commands: Commands, meshes: Query<(Entity, &UntypedMaterial)>) { for (entity, handle) in meshes.iter() { let handle = handle.deref().clone(); diff --git a/crates/processing_render/src/particles/mod.rs b/crates/processing_render/src/particles/mod.rs index 02cbd9a9..28efe687 100644 --- a/crates/processing_render/src/particles/mod.rs +++ b/crates/processing_render/src/particles/mod.rs @@ -1,22 +1,45 @@ -//! gpu-resident particle / instancing container. See `docs/particles.md`. +//! See `docs/particles.md`. +mod emit; pub mod kernels; pub mod material; pub mod pack; +mod scatter; + +pub use emit::{particles_apply, particles_emit, particles_emit_gpu}; +pub use kernels::{ + BOUNDS_CLAMP, BOUNDS_REFLECT, BOUNDS_SOFT, BOUNDS_WRAP, COMBINE_ADD, COMBINE_DIV, COMBINE_MAX, + COMBINE_MIN, COMBINE_MUL, COMBINE_POW, COMBINE_SUB, FALLOFF_CONST, FALLOFF_CUBIC, + FALLOFF_INVERSE, FALLOFF_LINEAR, FALLOFF_QUADRATIC, FALLOFF_SMOOTHSTEP, + particles_kernel_age, particles_kernel_attr_combine, particles_kernel_attr_linear, + particles_kernel_attr_lookup1d, particles_kernel_attr_lookup2d, particles_kernel_attr_mix, + particles_kernel_attract, particles_kernel_bounds_box, particles_kernel_bounds_geometry, + particles_kernel_bounds_sphere, particles_kernel_drag, particles_kernel_field, + particles_kernel_flock, particles_kernel_force, particles_kernel_impulse, + particles_kernel_integrate, particles_kernel_noise, particles_kernel_orient, + particles_kernel_transform, particles_kernel_vortex, +}; +pub use scatter::{ + particles_scatter_create, particles_scatter_volume_create, prepare_scatter_source, + prepare_scatter_volume_source, +}; use bevy::asset::RenderAssetUsages; use bevy::mesh::VertexAttributeValues; use bevy::pbr::gpu_instance_batch::GpuInstanceBatchPlugin; use bevy::platform::collections::HashMap; use bevy::prelude::*; +use bevy::render::RenderApp; +use bevy::render::mesh::allocator::MeshAllocatorSettings; use bevy::render::render_resource::{BufferDescriptor, BufferUsages}; use bevy::render::renderer::RenderDevice; use bevy::render::storage::ShaderBuffer; -use processing_core::error::{ProcessingError, Result}; +use processing_core::app_mut; +use processing_core::error::{self, ProcessingError, Result}; use crate::compute; -use crate::geometry::{Attribute, AttributeFormat, Geometry}; +use crate::geometry::{Attribute, AttributeFormat, Geometry, default_attribute_init}; pub struct ParticlesPlugin; @@ -27,18 +50,29 @@ impl Plugin for ParticlesPlugin { app.add_plugins(material::ParticlesMaterialPlugin); app.add_plugins(kernels::ParticlesKernelsPlugin); } + + fn finish(&self, app: &mut App) { + // The mesh allocator emits its GPU buffers before the render device + // exists, so the STORAGE flag must be set in finish(), not build(). + let Some(render_app) = app.get_sub_app_mut(RenderApp) else { + return; + }; + render_app + .world_mut() + .resource_mut::() + .extra_buffer_usages |= BufferUsages::STORAGE; + } } #[derive(Component)] pub struct Particles { pub capacity: u32, - /// `Attribute` entity → backing `compute::Buffer` entity. + /// `Attribute` entity to backing `compute::Buffer` entity. pub buffers: HashMap, - /// lazy persistent rasterization entity. Must outlive the per-frame draw - /// because `GpuInstanceBatchReservations` queue mesh batches one frame - /// behind, so respawning per-frame loses the reservation. + /// Must outlive the per-frame draw: `GpuInstanceBatchReservations` queues + /// mesh batches one frame behind, so respawning per-frame loses the reservation. pub draw_entity: Option, - /// ring-buffer write cursor for `particles_emit`. Wraps at `capacity`. + /// Ring-buffer write cursor; wraps at `capacity`. pub emit_head: u32, } @@ -48,7 +82,6 @@ impl Particles { } } -/// render-side marker pointing at the [`Particles`] entity to pack from. #[derive(Component, Clone, Copy)] pub struct ParticlesDraw { pub particles: Entity, @@ -87,9 +120,6 @@ pub fn create( Ok(entity) } -/// capacity = source mesh's vertex count. Registered attributes are seeded -/// from the matching mesh attribute (by name + format); unmatched ones are -/// zero-initialized. pub fn create_from_geometry( In((geom_entity, attribute_entities)): In<(Entity, Vec)>, mut commands: Commands, @@ -205,3 +235,204 @@ pub fn destroy( commands.entity(entity).despawn(); Ok(()) } + +pub enum AttributeSeed { + Ensure, + Declare(Option>), +} + +fn tile_seed(per_element: &[u8], capacity: usize) -> Vec { + let mut bytes = Vec::with_capacity(capacity * per_element.len()); + for _ in 0..capacity { + bytes.extend_from_slice(per_element); + } + bytes +} + +fn convention_seed_bytes(name: &str, format: AttributeFormat) -> Vec { + default_attribute_init(name, format) + .iter() + .flat_map(|f| f.to_le_bytes()) + .collect() +} + +pub fn materialize_attribute( + In((particles_entity, attribute_entity, seed)): In<(Entity, Entity, AttributeSeed)>, + mut commands: Commands, + mut particles_q: Query<&mut Particles>, + attributes: Query<&Attribute>, + mut shader_buffers: ResMut>, + render_device: Res, +) -> Result { + let attr = attributes + .get(attribute_entity) + .map_err(|_| ProcessingError::InvalidEntity)? + .clone(); + + let capacity = { + let particles = particles_q + .get(particles_entity) + .map_err(|_| ProcessingError::ParticlesNotFound)?; + let existing = match particles.buffers.get(&attribute_entity).copied() { + Some(buf) => Some(buf), + None => { + let mut hit = None; + for (&e, &buf) in &particles.buffers { + let Ok(other) = attributes.get(e) else { continue }; + if other.name == attr.name { + if other.format != attr.format { + return Err(ProcessingError::InvalidArgument(format!( + "attribute '{}' already present with a different format", + attr.name + ))); + } + hit = Some(buf); + break; + } + } + hit + } + }; + if let Some(buf) = existing { + return match seed { + AttributeSeed::Ensure => Ok(buf), + AttributeSeed::Declare(_) => Err(ProcessingError::InvalidArgument(format!( + "particles already have attribute '{}'", + attr.name + ))), + }; + } + particles.capacity as usize + }; + + let elem_size = attr.format.byte_size(); + let per_element = match &seed { + AttributeSeed::Declare(Some(default_bytes)) => { + if default_bytes.len() != elem_size { + return Err(ProcessingError::InvalidArgument(format!( + "default value byte size {} does not match attribute '{}' format byte size {}", + default_bytes.len(), + attr.name, + elem_size, + ))); + } + default_bytes.clone() + } + AttributeSeed::Ensure | AttributeSeed::Declare(None) => { + let seed = convention_seed_bytes(attr.name, attr.format); + if seed.len() != elem_size { + return Err(ProcessingError::InvalidArgument(format!( + "attribute '{}' reuses a builtin name with an incompatible format; \ + declare it with an explicit default", + attr.name + ))); + } + seed + } + }; + + let initial = tile_seed(&per_element, capacity); + let buffer_entity = make_buffer(&mut commands, &mut shader_buffers, &render_device, &initial); + particles_q + .get_mut(particles_entity) + .map_err(|_| ProcessingError::ParticlesNotFound)? + .buffers + .insert(attribute_entity, buffer_entity); + Ok(buffer_entity) +} + +pub fn particles_create( + capacity: u32, + attribute_entities: Vec, +) -> error::Result { + app_mut(|app| { + app.world_mut() + .run_system_cached_with(create, (capacity, attribute_entities)) + .unwrap() + }) +} + +pub fn particles_create_from_geometry( + geometry_entity: Entity, + attribute_entities: Vec, +) -> error::Result { + app_mut(|app| { + app.world_mut() + .run_system_cached_with(create_from_geometry, (geometry_entity, attribute_entities)) + .unwrap() + }) +} + +pub fn particles_destroy(entity: Entity) -> error::Result<()> { + app_mut(|app| { + app.world_mut() + .run_system_cached_with(destroy, entity) + .unwrap() + }) +} + +pub fn particles_capacity(entity: Entity) -> error::Result { + app_mut(|app| { + Ok(app + .world() + .get::(entity) + .ok_or(error::ProcessingError::ParticlesNotFound)? + .capacity) + }) +} + +pub fn particles_buffer( + entity: Entity, + attribute_entity: Entity, +) -> error::Result> { + app_mut(|app| { + Ok(app + .world() + .get::(entity) + .ok_or(error::ProcessingError::ParticlesNotFound)? + .buffer(attribute_entity)) + }) +} + +pub fn particles_ensure_attribute( + particles_entity: Entity, + attribute_entity: Entity, +) -> error::Result { + app_mut(|app| { + app.world_mut() + .run_system_cached_with( + materialize_attribute, + (particles_entity, attribute_entity, AttributeSeed::Ensure), + ) + .unwrap() + }) +} + +pub fn particles_attribute_add( + particles_entity: Entity, + attribute_entity: Entity, + default: Option, +) -> error::Result<()> { + let default_bytes = match default { + None => None, + Some(v) => Some(v.to_bytes().ok_or_else(|| { + error::ProcessingError::InvalidArgument( + "default must be a scalar/vector ShaderValue, not a Buffer/Texture/Mesh*" + .to_string(), + ) + })?), + }; + app_mut(|app| { + app.world_mut() + .run_system_cached_with( + materialize_attribute, + ( + particles_entity, + attribute_entity, + AttributeSeed::Declare(default_bytes), + ), + ) + .unwrap() + .map(|_| ()) + }) +} diff --git a/crates/processing_render/src/particles/pack.rs b/crates/processing_render/src/particles/pack.rs index 253aee00..0d844113 100644 --- a/crates/processing_render/src/particles/pack.rs +++ b/crates/processing_render/src/particles/pack.rs @@ -1,7 +1,3 @@ -//! compute pass that writes [`Particles`] position/rotation/scale/dead into -//! the per-instance slots reserved by [`GpuBatchedMesh3d`]. pipelines are -//! cached per `(HAS_ROTATION, HAS_SCALE, HAS_DEAD)` shader_def combination. - use std::num::NonZeroU64; use bevy::core_pipeline::Core3d; @@ -66,13 +62,11 @@ impl Plugin for ParticlesPackPlugin { #[derive(Resource, Clone)] pub struct ParticlesPackShader(pub Handle); -/// specialization key — controls which `#ifdef`s are set when compiling the pack shader, -/// and which bindings are present in the bind-group layout. #[derive(Hash, Eq, PartialEq, Clone, Copy, Debug)] pub struct PackPipelineKey { pub has_rotation: bool, pub has_scale: bool, - pub has_dead: bool, + pub has_life: bool, } pub struct CachedPackPipeline { @@ -98,7 +92,7 @@ pub struct ExtractedParticlesData { pub position: Handle, pub rotation: Option>, pub scale: Option>, - pub dead: Option>, + pub life: Option>, } #[derive(Resource, Default)] @@ -145,7 +139,7 @@ fn pack_layout_entries(key: PackPipelineKey) -> Vec { if key.has_scale { entries.push(layout_entry(4, storage_r)); } - if key.has_dead { + if key.has_life { entries.push(layout_entry(5, storage_r)); } entries.push(layout_entry(6, uniform)); @@ -169,8 +163,8 @@ fn shader_defs_for(key: PackPipelineKey) -> Vec { if key.has_scale { defs.push("HAS_SCALE".into()); } - if key.has_dead { - defs.push("HAS_DEAD".into()); + if key.has_life { + defs.push("HAS_LIFE".into()); } defs } @@ -186,16 +180,16 @@ fn get_or_create_pipeline( } let bind_group_layout = BindGroupLayoutDescriptor::new( format!( - "ParticlesPackBindGroupLayout(rot={},scale={},dead={})", - key.has_rotation, key.has_scale, key.has_dead + "ParticlesPackBindGroupLayout(rot={},scale={},life={})", + key.has_rotation, key.has_scale, key.has_life ), &pack_layout_entries(key), ); let pipeline = pipeline_cache.queue_compute_pipeline(ComputePipelineDescriptor { label: Some( format!( - "particles_pack_pipeline(rot={},scale={},dead={})", - key.has_rotation, key.has_scale, key.has_dead + "particles_pack_pipeline(rot={},scale={},life={})", + key.has_rotation, key.has_scale, key.has_life ) .into(), ), @@ -241,15 +235,15 @@ fn extract_particles_draws( .buffer(builtins.scale) .and_then(|e| buffers.get(e).ok()) .map(|b| b.handle.clone()); - let dead = p - .buffer(builtins.dead) + let life = p + .buffer(builtins.life) .and_then(|e| buffers.get(e).ok()) .map(|b| b.handle.clone()); let key = PackPipelineKey { has_rotation: rotation.is_some(), has_scale: scale.is_some(), - has_dead: dead.is_some(), + has_life: life.is_some(), }; extracted.by_main.insert( MainEntity::from(entity), @@ -258,7 +252,7 @@ fn extract_particles_draws( position: pos_buf.handle.clone(), rotation, scale, - dead, + life, }, ); } @@ -305,8 +299,8 @@ fn prepare_pack_bind_groups( if data.key.has_scale && gpu_scale.is_none() { continue; } - let gpu_dead = data.dead.as_ref().and_then(|h| gpu_buffers.get(h)); - if data.key.has_dead && gpu_dead.is_none() { + let gpu_life = data.life.as_ref().and_then(|h| gpu_buffers.get(h)); + if data.key.has_life && gpu_life.is_none() { continue; } @@ -354,10 +348,10 @@ fn prepare_pack_bind_groups( resource: gpu_scale.buffer.as_entire_binding(), }); } - if let Some(gpu_dead) = gpu_dead { + if let Some(gpu_life) = gpu_life { entries.push(BindGroupEntry { binding: 5, - resource: gpu_dead.buffer.as_entire_binding(), + resource: gpu_life.buffer.as_entire_binding(), }); } entries.push(BindGroupEntry { diff --git a/crates/processing_render/src/particles/pack.wgsl b/crates/processing_render/src/particles/pack.wgsl index 020cfb96..c1b0e460 100644 --- a/crates/processing_render/src/particles/pack.wgsl +++ b/crates/processing_render/src/particles/pack.wgsl @@ -17,7 +17,6 @@ struct MeshCullingData { aabb_center: vec3, _pad: f32, aabb_half_extents: vec3, - // > 0.0 = render, <= 0.0 = skip in preprocessing. life: f32, } @@ -37,8 +36,8 @@ struct PackParams { #ifdef HAS_SCALE @group(0) @binding(4) var scale: array; #endif -#ifdef HAS_DEAD -@group(0) @binding(5) var dead: array; +#ifdef HAS_LIFE +@group(0) @binding(5) var life: array; #endif @group(0) @binding(6) var params: PackParams; @@ -106,9 +105,8 @@ fn pack(@builtin(global_invocation_id) gid: vec3) { mesh_culling_buffer[slot].aabb_center = vec3(0.0, 0.0, 0.0); mesh_culling_buffer[slot].aabb_half_extents = vec3(1.0, 1.0, 1.0); -#ifdef HAS_DEAD - // dead[i]: 0.0 = alive, nonzero = dead -> life > 0.0 renders. - mesh_culling_buffer[slot].life = select(1.0, 0.0, dead[i] != 0.0); +#ifdef HAS_LIFE + mesh_culling_buffer[slot].life = life[i]; #else mesh_culling_buffer[slot].life = 1.0; #endif diff --git a/crates/processing_render/src/particles/particles.wgsl b/crates/processing_render/src/particles/particles.wgsl index a3b2ecaf..05a7419d 100644 --- a/crates/processing_render/src/particles/particles.wgsl +++ b/crates/processing_render/src/particles/particles.wgsl @@ -16,8 +16,15 @@ } #endif +#ifdef HAS_COLORS @group(#{MATERIAL_BIND_GROUP}) @binding(100) var particle_colors: array>; +#endif + +#ifdef HAS_EMISSIVE_COLORS +@group(#{MATERIAL_BIND_GROUP}) @binding(101) +var particle_emissive_colors: array>; +#endif @fragment fn fragment( @@ -27,7 +34,17 @@ fn fragment( var pbr_input = pbr_input_from_standard_material(in, is_front); let tag = mesh_functions::get_tag(in.instance_index); + +#ifdef HAS_COLORS pbr_input.material.base_color = pbr_input.material.base_color * particle_colors[tag]; +#endif + +#ifdef HAS_EMISSIVE_COLORS + pbr_input.material.emissive = vec4( + pbr_input.material.emissive.rgb + particle_emissive_colors[tag].rgb, + pbr_input.material.emissive.a + ); +#endif pbr_input.material.base_color = alpha_discard(pbr_input.material, pbr_input.material.base_color); diff --git a/crates/processing_render/src/particles/scatter.rs b/crates/processing_render/src/particles/scatter.rs new file mode 100644 index 00000000..fef13f9c --- /dev/null +++ b/crates/processing_render/src/particles/scatter.rs @@ -0,0 +1,197 @@ +use bevy::mesh::{Indices, VertexAttributeValues}; +use bevy::prelude::*; + +use processing_core::app_mut; +use processing_core::error::{self, ProcessingError, Result}; + +use crate::geometry::Geometry; +use crate::shader_value::ShaderValue; +use crate::{ + buffer_create_with_data, compute_create, compute_set, geometry_attribute_position, shader_load, +}; + +pub fn particles_scatter_create(source_geometry: Entity) -> error::Result { + let (cdf_bytes, indices_bytes, face_count) = app_mut(|app| { + app.world_mut() + .run_system_cached_with(prepare_scatter_source, source_geometry) + .unwrap() + })?; + + let cdf_buf = buffer_create_with_data(cdf_bytes)?; + let idx_buf = buffer_create_with_data(indices_bytes)?; + + let shader = + shader_load("embedded://processing_render/particles/kernels/scatter_surface.wgsl")?; + let scatter = compute_create(shader)?; + + let position_attr = geometry_attribute_position(); + compute_set( + scatter, + "source_position", + ShaderValue::MeshAttribute(source_geometry, position_attr), + )?; + compute_set( + scatter, + "source_indices", + ShaderValue::Buffer(idx_buf), + )?; + compute_set(scatter, "cdf", ShaderValue::Buffer(cdf_buf))?; + compute_set(scatter, "face_count", ShaderValue::UInt(face_count))?; + compute_set(scatter, "seed", ShaderValue::UInt(0xc0ffeeu32))?; + + Ok(scatter) +} + +pub fn particles_scatter_volume_create(source_geometry: Entity) -> error::Result { + let (indices_bytes, aabb_min, aabb_max, face_count) = app_mut(|app| { + app.world_mut() + .run_system_cached_with(prepare_scatter_volume_source, source_geometry) + .unwrap() + })?; + + let idx_buf = buffer_create_with_data(indices_bytes)?; + + let shader = + shader_load("embedded://processing_render/particles/kernels/scatter_volume.wgsl")?; + let scatter = compute_create(shader)?; + + let position_attr = geometry_attribute_position(); + compute_set( + scatter, + "source_position", + ShaderValue::MeshAttribute(source_geometry, position_attr), + )?; + compute_set( + scatter, + "source_indices", + ShaderValue::Buffer(idx_buf), + )?; + compute_set( + scatter, + "aabb_min", + ShaderValue::Float4([aabb_min[0], aabb_min[1], aabb_min[2], 0.0]), + )?; + compute_set( + scatter, + "aabb_max", + ShaderValue::Float4([aabb_max[0], aabb_max[1], aabb_max[2], 0.0]), + )?; + compute_set(scatter, "face_count", ShaderValue::UInt(face_count))?; + compute_set(scatter, "max_attempts", ShaderValue::UInt(32))?; + compute_set(scatter, "seed", ShaderValue::UInt(0xc0ffeeu32))?; + + Ok(scatter) +} + +fn extract_scatter_geometry(mesh: &mut Mesh) -> Result<(Vec<[f32; 3]>, Vec)> { + mesh.deinterleave(); + + let positions = match mesh.attribute(Mesh::ATTRIBUTE_POSITION) { + Some(VertexAttributeValues::Float32x3(p)) => p.clone(), + _ => { + return Err(ProcessingError::InvalidArgument( + "scatter source mesh has no Float32x3 position attribute".to_string(), + )); + } + }; + + let dense_indices: Vec = match mesh.indices() { + Some(Indices::U16(v)) => v.iter().map(|&i| i as u32).collect(), + Some(Indices::U32(v)) => v.clone(), + None => { + if positions.len() % 3 != 0 { + return Err(ProcessingError::InvalidArgument( + "scatter source mesh has no indices and a vertex count that isn't a \ + multiple of 3" + .to_string(), + )); + } + (0..positions.len() as u32).collect() + } + }; + + if dense_indices.len() % 3 != 0 { + return Err(ProcessingError::InvalidArgument( + "scatter source mesh has a non-triangle index list".to_string(), + )); + } + if dense_indices.is_empty() { + return Err(ProcessingError::InvalidArgument( + "scatter source mesh has no triangles".to_string(), + )); + } + + Ok((positions, dense_indices)) +} + +pub fn prepare_scatter_source( + In(geom_entity): In, + geometries: Query<&Geometry>, + mut meshes: ResMut>, +) -> Result<(Vec, Vec, u32)> { + let geom = geometries + .get(geom_entity) + .map_err(|_| ProcessingError::GeometryNotFound)?; + let mesh = meshes + .get_mut(&geom.handle) + .ok_or(ProcessingError::GeometryNotFound)? + .into_inner(); + + let (positions, dense_indices) = extract_scatter_geometry(mesh)?; + let face_count = (dense_indices.len() / 3) as u32; + + let mut cum = Vec::with_capacity(face_count as usize); + let mut total = 0.0_f32; + for face in 0..face_count as usize { + let i0 = dense_indices[face * 3] as usize; + let i1 = dense_indices[face * 3 + 1] as usize; + let i2 = dense_indices[face * 3 + 2] as usize; + let p0 = Vec3::from_array(positions[i0]); + let p1 = Vec3::from_array(positions[i1]); + let p2 = Vec3::from_array(positions[i2]); + let area = 0.5 * (p1 - p0).cross(p2 - p0).length(); + total += area; + cum.push(total); + } + if total <= 0.0 { + return Err(ProcessingError::InvalidArgument( + "scatter source mesh has zero surface area".to_string(), + )); + } + let inv = 1.0 / total; + for v in &mut cum { + *v *= inv; + } + + let cdf_bytes: Vec = cum.iter().flat_map(|f| f.to_le_bytes()).collect(); + let indices_bytes: Vec = dense_indices.iter().flat_map(|i| i.to_le_bytes()).collect(); + Ok((cdf_bytes, indices_bytes, face_count)) +} + +pub fn prepare_scatter_volume_source( + In(geom_entity): In, + geometries: Query<&Geometry>, + mut meshes: ResMut>, +) -> Result<(Vec, [f32; 3], [f32; 3], u32)> { + let geom = geometries + .get(geom_entity) + .map_err(|_| ProcessingError::GeometryNotFound)?; + let mesh = meshes + .get_mut(&geom.handle) + .ok_or(ProcessingError::GeometryNotFound)? + .into_inner(); + + let (positions, dense_indices) = extract_scatter_geometry(mesh)?; + let face_count = (dense_indices.len() / 3) as u32; + + let mut min = Vec3::splat(f32::INFINITY); + let mut max = Vec3::splat(f32::NEG_INFINITY); + for p in &positions { + let v = Vec3::from_array(*p); + min = min.min(v); + max = max.max(v); + } + + let indices_bytes: Vec = dense_indices.iter().flat_map(|i| i.to_le_bytes()).collect(); + Ok((indices_bytes, min.to_array(), max.to_array(), face_count)) +} diff --git a/crates/processing_render/src/render/mod.rs b/crates/processing_render/src/render/mod.rs index c72a5494..75a57607 100644 --- a/crates/processing_render/src/render/mod.rs +++ b/crates/processing_render/src/render/mod.rs @@ -1419,7 +1419,8 @@ fn particles_fill_material( ..Default::default() }, extension: ParticlesExtension { - colors: buf.handle.clone(), + colors: Some(buf.handle.clone()), + emissive_colors: None, }, }); Some(handle.untyped()) diff --git a/crates/processing_render/src/render/primitive/shape3d.rs b/crates/processing_render/src/render/primitive/shape3d.rs index 11766c23..1e4889b3 100644 --- a/crates/processing_render/src/render/primitive/shape3d.rs +++ b/crates/processing_render/src/render/primitive/shape3d.rs @@ -90,7 +90,6 @@ pub fn tetrahedron_mesh(radius: f32) -> Mesh { mesh } -/// 3d lattice of `nx * ny * nz` points, centered at the origin with `spacing` between them. pub fn grid_mesh(nx: u32, ny: u32, nz: u32, spacing: f32) -> Mesh { let count = (nx as usize) * (ny as usize) * (nz as usize); let mut positions = Vec::with_capacity(count); diff --git a/crates/processing_render/src/shader_value.rs b/crates/processing_render/src/shader_value.rs index 9d2c77de..2ca8dce6 100644 --- a/crates/processing_render/src/shader_value.rs +++ b/crates/processing_render/src/shader_value.rs @@ -14,6 +14,8 @@ pub enum ShaderValue { Mat4([f32; 16]), Texture(Entity), Buffer(Entity), + MeshAttribute(Entity, Entity), + MeshIndex(Entity), } impl ShaderValue { @@ -29,7 +31,10 @@ impl ShaderValue { ShaderValue::Int4(v) => Some(v.iter().flat_map(|i| i.to_le_bytes()).collect()), ShaderValue::UInt(v) => Some(v.to_le_bytes().to_vec()), ShaderValue::Mat4(v) => Some(v.iter().flat_map(|f| f.to_le_bytes()).collect()), - ShaderValue::Texture(_) | ShaderValue::Buffer(_) => None, + ShaderValue::Texture(_) + | ShaderValue::Buffer(_) + | ShaderValue::MeshAttribute(..) + | ShaderValue::MeshIndex(_) => None, } } @@ -40,7 +45,10 @@ impl ShaderValue { ShaderValue::Float3(_) | ShaderValue::Int3(_) => Some(12), ShaderValue::Float4(_) | ShaderValue::Int4(_) => Some(16), ShaderValue::Mat4(_) => Some(64), - ShaderValue::Texture(_) | ShaderValue::Buffer(_) => None, + ShaderValue::Texture(_) + | ShaderValue::Buffer(_) + | ShaderValue::MeshAttribute(..) + | ShaderValue::MeshIndex(_) => None, } } @@ -76,7 +84,10 @@ impl ShaderValue { bytes[..4].try_into().ok()?, ))), ShaderValue::Mat4(_) => Some(ShaderValue::Mat4(f32s::<16>(bytes)?)), - ShaderValue::Texture(_) | ShaderValue::Buffer(_) => None, + ShaderValue::Texture(_) + | ShaderValue::Buffer(_) + | ShaderValue::MeshAttribute(..) + | ShaderValue::MeshIndex(_) => None, } } } diff --git a/docs/particles.md b/docs/particles.md deleted file mode 100644 index 06e61343..00000000 --- a/docs/particles.md +++ /dev/null @@ -1,23 +0,0 @@ -# Particles - -`Particles` are a collection of attribute buffers that can be used in order to sequence compute shaders. They are -isomorphic to `Mesh` in the sense that they contain attributes and sets of data. In this way, you can think of a -`Mesh` as the CPU representation of a `Particles` object, and the `Particles` object as the GPU representation of a -`Mesh`. This allows convenient initialization of particle simulations from existing meshes, or using a mesh as a -constraint for a particle simulation, like a volume or a surface. - -Another way to consider particles would be as the compute equivalent of `Graphics`. Where the `Grpahics` object -allows you to issue high level rasterization commands, the `Particles` object allows you to issue high level compute -commands. In this way, you can think of a `Particles` object as a compute shader that is executed on the GPU, and the -attributes as the inputs and outputs of the compute shader. In practice, a compute shader may also require additional -data, such as textures or bound vertex buffers, but the `Particles` object provides a high level abstraction for -sequencing compute shaders and managing their inputs and outputs. - -## Rasterization - -Of course, it's not very interesting just to have a collection of attribute buffers on the GPU. The real power of -`Particles` comes from the ability to rasterize them. This is done using the `particles` function, which takes a -`Particles` object and a `Geometry` object as input, and issues a draw call that rasterizes the particles using the -current drawing state. In other words, it instances the provided geometry over the particles, using the attributes as -to determine each instance's position, orientation, and other properties. This provides a powerful bridge between -`Particles` and the rest of the drawing API. \ No newline at end of file diff --git a/examples/compute_readback.rs b/examples/compute_readback.rs index f470274d..cc160de7 100644 --- a/examples/compute_readback.rs +++ b/examples/compute_readback.rs @@ -2,12 +2,9 @@ use processing::prelude::*; fn main() { match run() { - Ok(_) => { - eprintln!("Compute readback test passed!"); - exit(0).unwrap(); - } + Ok(_) => exit(0).unwrap(), Err(e) => { - eprintln!("Compute readback error: {:?}", e); + eprintln!("{e:?}"); exit(1).unwrap(); } } @@ -45,8 +42,7 @@ fn main() { .map(|c| u32::from_le_bytes([c[0], c[1], c[2], c[3]])) .collect(); - assert_eq!(values, vec![1, 2, 3, 4], "Compute readback mismatch!"); - eprintln!("PASS"); + assert_eq!(values, vec![1, 2, 3, 4]); let double_src = r#" @group(0) @binding(0) @@ -72,12 +68,7 @@ fn main(@builtin(global_invocation_id) id: vec3) { .chunks_exact(4) .map(|c| f32::from_le_bytes([c[0], c[1], c[2], c[3]])) .collect(); - assert_eq!( - floats, - vec![2.0, 4.0, 6.0, 8.0], - "In-place double mismatch!" - ); - eprintln!("PASS"); + assert_eq!(floats, vec![2.0, 4.0, 6.0, 8.0]); compute_destroy(compute)?; compute_destroy(compute2)?; diff --git a/examples/particles_emit.rs b/examples/particles_emit.rs index 0ce4ddd9..3ab58099 100644 --- a/examples/particles_emit.rs +++ b/examples/particles_emit.rs @@ -31,7 +31,6 @@ fn sketch() -> error::Result<()> { let color_buf = particles_buffer(p, color_attr)?.ok_or(error::ProcessingError::ParticlesNotFound)?; - // push unemitted slots off-screen so they don't render at the origin let init_positions: Vec = (0..capacity * 3).map(|_| 1.0e6).collect(); buffer_write( position_buf, diff --git a/examples/particles_emit_gpu.rs b/examples/particles_emit_gpu.rs index 39c2958f..3f0fc501 100644 --- a/examples/particles_emit_gpu.rs +++ b/examples/particles_emit_gpu.rs @@ -17,7 +17,7 @@ struct Spawn { @group(0) @binding(2) var color: array; @group(0) @binding(3) var scale: array; @group(0) @binding(4) var age: array; -@group(0) @binding(5) var dead: array; +@group(0) @binding(5) var life: array; @group(0) @binding(6) var spawn: Spawn; @group(0) @binding(7) var emit_range: vec4; @@ -70,7 +70,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { scale[slot * 3u + 2u] = 1.0; age[slot] = 0.0; - dead[slot] = 0.0; + life[slot] = 1.0; } "#; @@ -86,7 +86,7 @@ struct Params { @group(0) @binding(1) var velocity: array; @group(0) @binding(2) var scale: array; @group(0) @binding(3) var age: array; -@group(0) @binding(4) var dead: array; +@group(0) @binding(4) var life: array; @group(0) @binding(5) var params: Params; @compute @workgroup_size(64) @@ -94,7 +94,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { let i = gid.x; let count = arrayLength(&age); if i >= count { return; } - if dead[i] != 0.0 { return; } + if life[i] <= 0.0 { return; } age[i] = age[i] + params.dt; @@ -110,7 +110,7 @@ fn main(@builtin(global_invocation_id) gid: vec3) { scale[i * 3u + 1u] = s; scale[i * 3u + 2u] = s; - if age[i] > params.ttl { dead[i] = 1.0; } + if age[i] > params.ttl { life[i] = 0.0; } } "#; @@ -139,7 +139,7 @@ fn sketch() -> error::Result<()> { let position_attr = geometry_attribute_position(); let color_attr = geometry_attribute_color(); let scale_attr = geometry_attribute_scale(); - let dead_attr = geometry_attribute_dead(); + let life_attr = geometry_attribute_life(); let velocity_attr = geometry_attribute_create("velocity", AttributeFormat::Float3)?; let age_attr = geometry_attribute_create("age", AttributeFormat::Float)?; @@ -149,18 +149,12 @@ fn sketch() -> error::Result<()> { position_attr, color_attr, scale_attr, - dead_attr, + life_attr, velocity_attr, age_attr, ], )?; - // mark all unemitted slots dead so they don't render at origin. - let dead_buf = - particles_buffer(p, dead_attr)?.ok_or(error::ProcessingError::ParticlesNotFound)?; - let init_dead: Vec = (0..capacity).flat_map(|_| 1.0_f32.to_le_bytes()).collect(); - buffer_write(dead_buf, init_dead)?; - let color_buf = particles_buffer(p, color_attr)?.ok_or(error::ProcessingError::ParticlesNotFound)?; let mat = { diff --git a/examples/particles_lifecycle.rs b/examples/particles_lifecycle.rs index 17881a95..cb4d0371 100644 --- a/examples/particles_lifecycle.rs +++ b/examples/particles_lifecycle.rs @@ -7,7 +7,7 @@ use processing_render::render::command::DrawCommand; const AGING_SHADER: &str = r#" @group(0) @binding(0) var age: array; -@group(0) @binding(1) var dead: array; +@group(0) @binding(1) var life: array; @group(0) @binding(2) var position: array; @group(0) @binding(3) var scale: array; @group(0) @binding(4) var params: vec4; // x = dt, y = ttl @@ -22,22 +22,21 @@ fn main(@builtin(global_invocation_id) gid: vec3) { let dt = params.x; let ttl = params.y; - if dead[i] != 0.0 { + if life[i] <= 0.0 { return; } age[i] = age[i] + dt; position[i * 3u + 1u] = position[i * 3u + 1u] - dt * 1.5; - // shrink toward zero as age approaches ttl so dying is visible. - let life = clamp(1.0 - age[i] / ttl, 0.0, 1.0); - let s = life * life; // ease out + let remaining = clamp(1.0 - age[i] / ttl, 0.0, 1.0); + let s = remaining * remaining; scale[i * 3u + 0u] = s; scale[i * 3u + 1u] = s; scale[i * 3u + 2u] = s; if age[i] > ttl { - dead[i] = 1.0; + life[i] = 0.0; } } "#; @@ -64,22 +63,16 @@ fn sketch() -> error::Result<()> { let position_attr = geometry_attribute_position(); let color_attr = geometry_attribute_color(); let scale_attr = geometry_attribute_scale(); - let dead_attr = geometry_attribute_dead(); + let life_attr = geometry_attribute_life(); let age_attr = geometry_attribute_create("age", AttributeFormat::Float)?; let p = particles_create( capacity, - vec![position_attr, color_attr, scale_attr, dead_attr, age_attr], + vec![position_attr, color_attr, scale_attr, life_attr, age_attr], )?; - let dead_buf = - particles_buffer(p, dead_attr)?.ok_or(error::ProcessingError::ParticlesNotFound)?; let color_buf = particles_buffer(p, color_attr)?.ok_or(error::ProcessingError::ParticlesNotFound)?; - // mark all slots dead initially so the unemitted ring slots don't render. - let init_dead: Vec = (0..capacity).flat_map(|_| 1.0_f32.to_le_bytes()).collect(); - buffer_write(dead_buf, init_dead)?; - let mat = { let m = material_create_unlit()?; material_set_albedo_buffer(m, color_buf)?; @@ -130,7 +123,7 @@ fn sketch() -> error::Result<()> { let position_bytes: Vec = positions.iter().flat_map(|f| f.to_le_bytes()).collect(); let color_bytes: Vec = colors.iter().flat_map(|f| f.to_le_bytes()).collect(); let zero_floats: Vec = (0..burst).flat_map(|_| 0.0_f32.to_le_bytes()).collect(); - // init scale to 1; the aging shader shrinks it over time + let one_floats: Vec = (0..burst).flat_map(|_| 1.0_f32.to_le_bytes()).collect(); let one_scale: Vec = (0..burst) .flat_map(|_| { [1.0_f32, 1.0, 1.0] @@ -146,8 +139,8 @@ fn sketch() -> error::Result<()> { (position_attr, position_bytes), (color_attr, color_bytes), (scale_attr, one_scale), - (age_attr, zero_floats.clone()), - (dead_attr, zero_floats), + (age_attr, zero_floats), + (life_attr, one_floats), ], )?; diff --git a/examples/particles_oriented.rs b/examples/particles_oriented.rs index 29b45d43..4dcf443d 100644 --- a/examples/particles_oriented.rs +++ b/examples/particles_oriented.rs @@ -70,7 +70,6 @@ fn sketch() -> error::Result<()> { positions.push((x as f32 - 2.0) * 1.6); positions.push((y as f32 - 2.0) * 1.6); positions.push((z as f32 - 2.0) * 1.6); - // identity quat rotations.push(0.0); rotations.push(0.0); rotations.push(0.0); diff --git a/examples/particles_scatter.rs b/examples/particles_scatter.rs new file mode 100644 index 00000000..69a41173 --- /dev/null +++ b/examples/particles_scatter.rs @@ -0,0 +1,121 @@ +use processing_glfw::GlfwContext; +use std::time::Instant; + +use bevy::math::Vec3; +use processing::prelude::*; +use processing_render::geometry::AttributeFormat; +use processing_render::render::command::DrawCommand; + +const AGE_SHADER: &str = r#" +struct Params { dt: f32, ttl: f32, _pad0: f32, _pad1: f32 } + +@group(0) @binding(0) var scale: array; +@group(0) @binding(1) var age: array; +@group(0) @binding(2) var life: array; +@group(0) @binding(3) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&age); + if i >= count { return; } + if life[i] <= 0.0 { return; } + + age[i] = age[i] + params.dt; + let t = age[i] / params.ttl; + let rise = clamp(t / 0.05, 0.0, 1.0); + let fall = clamp((1.0 - t) / 0.6, 0.0, 1.0); + let s = rise * fall; + scale[i * 3u + 0u] = s; + scale[i * 3u + 1u] = s; + scale[i * 3u + 2u] = s; + + if age[i] > params.ttl { life[i] = 0.0; } +} +"#; + +fn main() { + sketch().unwrap(); + exit(0).unwrap(); +} + +fn sketch() -> error::Result<()> { + let mut glfw_ctx = GlfwContext::new(900, 700)?; + init(Config::default())?; + + let surface = glfw_ctx.create_surface(900, 700)?; + let graphics = graphics_create(surface, 900, 700, TextureFormat::Rgba16Float)?; + + graphics_mode_3d(graphics)?; + transform_set_position(graphics, Vec3::new(0.0, 0.4, 4.5))?; + transform_look_at(graphics, Vec3::ZERO)?; + + let _key = light_create_directional( + graphics, + bevy::color::Color::srgb(1.0, 0.95, 0.85), + 4500.0, + )?; + + let source = geometry_sphere(1.2, 96, 48)?; + let scatter = particles_scatter_create(source)?; + + let particle = geometry_sphere(0.005, 6, 4)?; + + let capacity: u32 = 40_000; + let position_attr = geometry_attribute_position(); + let scale_attr = geometry_attribute_scale(); + let life_attr = geometry_attribute_life(); + let age_attr = geometry_attribute_create("age", AttributeFormat::Float)?; + + let p = particles_create( + capacity, + vec![position_attr, scale_attr, life_attr, age_attr], + )?; + + let age_shader = shader_create(AGE_SHADER)?; + let aging = compute_create(age_shader)?; + + let mat = material_create_pbr()?; + material_set_albedo_color(mat, [0.9, 0.85, 1.0, 1.0])?; + + let burst: u32 = 600; + let dt: f32 = 1.0 / 60.0; + let ttl: f32 = 4.0; + let start = Instant::now(); + + while glfw_ctx.poll_events() { + graphics_begin_draw(graphics)?; + graphics_record_command( + graphics, + DrawCommand::BackgroundColor(bevy::color::Color::srgb(0.03, 0.03, 0.05)), + )?; + graphics_record_command(graphics, DrawCommand::Material(mat))?; + graphics_record_command( + graphics, + DrawCommand::Particles { + particles: p, + geometry: particle, + }, + )?; + graphics_end_draw(graphics)?; + + let t = start.elapsed().as_secs_f32(); + let cam_x = (t * 0.25).cos() * 4.5; + let cam_z = (t * 0.25).sin() * 4.5; + transform_set_position(graphics, Vec3::new(cam_x, 0.4, cam_z))?; + transform_look_at(graphics, Vec3::ZERO)?; + + compute_set( + scatter, + "seed", + shader_value::ShaderValue::UInt((t * 1000.0) as u32 ^ 0xc0ffeeu32), + )?; + particles_emit_gpu(p, burst, scatter)?; + + compute_set(aging, "dt", shader_value::ShaderValue::Float(dt))?; + compute_set(aging, "ttl", shader_value::ShaderValue::Float(ttl))?; + particles_apply(p, aging)?; + } + + Ok(()) +} diff --git a/examples/particles_scatter_volume.rs b/examples/particles_scatter_volume.rs new file mode 100644 index 00000000..df8636b9 --- /dev/null +++ b/examples/particles_scatter_volume.rs @@ -0,0 +1,113 @@ +use processing_glfw::GlfwContext; +use std::time::Instant; + +use bevy::math::Vec3; +use processing::prelude::*; +use processing_render::geometry::AttributeFormat; +use processing_render::render::command::DrawCommand; + +const AGE_SHADER: &str = r#" +struct Params { dt: f32, ttl: f32, _pad0: f32, _pad1: f32 } + +@group(0) @binding(0) var scale: array; +@group(0) @binding(1) var age: array; +@group(0) @binding(2) var life: array; +@group(0) @binding(3) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + let count = arrayLength(&age); + if i >= count { return; } + if life[i] <= 0.0 { return; } + + age[i] = age[i] + params.dt; + let t = age[i] / params.ttl; + let rise = clamp(t / 0.05, 0.0, 1.0); + let fall = clamp((1.0 - t) / 0.6, 0.0, 1.0); + let s = rise * fall; + scale[i * 3u + 0u] = s; + scale[i * 3u + 1u] = s; + scale[i * 3u + 2u] = s; + + if age[i] > params.ttl { life[i] = 0.0; } +} +"#; + +fn main() { + sketch().unwrap(); + exit(0).unwrap(); +} + +fn sketch() -> error::Result<()> { + let mut glfw_ctx = GlfwContext::new(600, 400)?; + init(Config::default())?; + + let surface = glfw_ctx.create_surface(900, 700)?; + let graphics = graphics_create(surface, 900, 700, TextureFormat::Rgba16Float)?; + + graphics_mode_3d(graphics)?; + transform_set_position(graphics, Vec3::new(0.0, 100.0, 400.0))?; + transform_look_at(graphics, Vec3::new(0.0, 80.0, 0.0))?; + graphics_orbit_camera(graphics)?; + + let gltf = gltf_load(graphics, "gltf/Duck.glb")?; + let duck = gltf_geometry(gltf, "LOD3spShape")?; + let scatter = particles_scatter_volume_create(duck)?; + + let particle = geometry_sphere(0.15, 4, 3)?; + + let capacity: u32 = 30_000; + let position_attr = geometry_attribute_position(); + let scale_attr = geometry_attribute_scale(); + let life_attr = geometry_attribute_life(); + let age_attr = geometry_attribute_create("age", AttributeFormat::Float)?; + + let p = particles_create( + capacity, + vec![position_attr, scale_attr, life_attr, age_attr], + )?; + + let age_shader = shader_create(AGE_SHADER)?; + let aging = compute_create(age_shader)?; + + let mat = material_create_unlit()?; + material_set_albedo_color(mat, [1.0, 1.0, 1.0, 1.0])?; + + let burst: u32 = 250; + let dt: f32 = 1.0 / 60.0; + let ttl: f32 = 5.0; + let start = Instant::now(); + + while glfw_ctx.poll_events() { + graphics_begin_draw(graphics)?; + graphics_record_command( + graphics, + DrawCommand::BackgroundColor(bevy::color::Color::srgb(0.03, 0.03, 0.05)), + )?; + graphics_record_command(graphics, DrawCommand::Material(mat))?; + graphics_record_command( + graphics, + DrawCommand::Particles { + particles: p, + geometry: particle, + }, + )?; + graphics_end_draw(graphics)?; + + let t = start.elapsed().as_secs_f32(); + + compute_set( + scatter, + "seed", + shader_value::ShaderValue::UInt((t * 1000.0) as u32 ^ 0xc0ffeeu32), + )?; + particles_emit_gpu(p, burst, scatter)?; + + compute_set(aging, "dt", shader_value::ShaderValue::Float(dt))?; + compute_set(aging, "ttl", shader_value::ShaderValue::Float(ttl))?; + particles_apply(p, aging)?; + } + + Ok(()) +} diff --git a/examples/particles_stress.rs b/examples/particles_stress.rs index e80827ae..281f4382 100644 --- a/examples/particles_stress.rs +++ b/examples/particles_stress.rs @@ -1,12 +1,10 @@ -//! stress test with `GRID^3` pbr-lit cubes rotating. tune `GRID` to scale. - use processing_glfw::GlfwContext; use bevy::math::Vec3; use processing::prelude::*; use processing_render::render::command::DrawCommand; -const GRID: u32 = 100; +const GRID: u32 = 100; // GRID^3 particles const SPACING: f32 = 1.0; const SPIN_SHADER: &str = r#" @@ -118,7 +116,7 @@ fn sketch() -> error::Result<()> { let spin_shader = shader_create(SPIN_SHADER)?; let spin = compute_create(spin_shader)?; - eprintln!("field_stress: {capacity} particles"); + eprintln!("{capacity} particles"); while glfw_ctx.poll_events() { graphics_begin_draw(graphics)?; diff --git a/examples/particles_text_whirl.rs b/examples/particles_text_whirl.rs new file mode 100644 index 00000000..1d98b802 --- /dev/null +++ b/examples/particles_text_whirl.rs @@ -0,0 +1,417 @@ +use processing_glfw::GlfwContext; +use std::time::Instant; + +use bevy::math::Vec3; +use processing::prelude::*; +use processing_render::render::command::{DrawCommand, TextStyle}; + +const BASE_COUNT: u32 = 6000; +const TRAIL_LEN_MIN: u32 = 8; +const TRAIL_LEN_MAX: u32 = 150; +const TRAIL_LEN: u32 = TRAIL_LEN_MAX; +const CAPACITY: u32 = BASE_COUNT * TRAIL_LEN; +const TRAIL_STRIDE: u32 = 1; + +const NOISE_STRENGTH: f32 = 50.5; +const SPHERE_RADIUS: f32 = 1.05; + +fn sim_shader() -> String { + r#" +struct Params { + base_count: u32, + trail_len_max: u32, + trail_len_min: u32, + seed: u32, + face_count: u32, + _pad0: u32, + _pad1: u32, + _pad2: u32, + noise_scale: f32, + noise_strength: f32, + time: f32, + step_speed: f32, +} + +@group(0) @binding(0) var source_position: array; +@group(0) @binding(1) var source_indices: array; +@group(0) @binding(2) var cdf: array; +@group(0) @binding(3) var head_pos: array; +@group(0) @binding(4) var anchor_normal: array; +@group(0) @binding(5) var trail_len_buf: array; +@group(0) @binding(6) var trail_head_buf: array; +@group(0) @binding(7) var position: array; +@group(0) @binding(8) var scale: array; +@group(0) @binding(9) var life: array; +@group(0) @binding(10) var params: Params; + +fn hash_u(n: u32) -> u32 { + var x = n; + x = (x ^ 61u) ^ (x >> 16u); + x = x + (x << 3u); + x = x ^ (x >> 4u); + x = x * 0x27d4eb2du; + x = x ^ (x >> 15u); + return x; +} +fn hash_unit(n: u32) -> f32 { return f32(hash_u(n)) / f32(0xffffffffu); } + +fn cdf_search(u: f32) -> u32 { + var lo: u32 = 0u; + var hi: u32 = params.face_count; + loop { + if lo >= hi { break; } + let mid = (lo + hi) >> 1u; + if cdf[mid] < u { lo = mid + 1u; } else { hi = mid; } + } + return min(lo, params.face_count - 1u); +} + +fn vhash(p: vec3) -> f32 { + let q = fract(p * 0.3183099) + vec3(0.1, 0.2, 0.3); + let r = q + dot(q, q.yzx + 19.19); + return fract(r.x * r.y * r.z); +} +fn value_noise(p: vec3) -> f32 { + let i = floor(p); + let f = fract(p); + let u = f * f * (3.0 - 2.0 * f); + return mix( + mix( + mix(vhash(i + vec3(0.0, 0.0, 0.0)), vhash(i + vec3(1.0, 0.0, 0.0)), u.x), + mix(vhash(i + vec3(0.0, 1.0, 0.0)), vhash(i + vec3(1.0, 1.0, 0.0)), u.x), + u.y), + mix( + mix(vhash(i + vec3(0.0, 0.0, 1.0)), vhash(i + vec3(1.0, 0.0, 1.0)), u.x), + mix(vhash(i + vec3(0.0, 1.0, 1.0)), vhash(i + vec3(1.0, 1.0, 1.0)), u.x), + u.y), + u.z); +} +fn noise3(p: vec3) -> vec3 { + return vec3( + value_noise(p), + value_noise(p + vec3(31.4, 0.0, 0.0)), + value_noise(p + vec3(0.0, 71.7, 0.0)), + ) * 2.0 - 1.0; +} +fn curl_noise(p: vec3) -> vec3 { + let eps = 0.01; + let dx = vec3(eps, 0.0, 0.0); + let dy = vec3(0.0, eps, 0.0); + let dz = vec3(0.0, 0.0, eps); + let n_xp = noise3(p + dx); let n_xm = noise3(p - dx); + let n_yp = noise3(p + dy); let n_ym = noise3(p - dy); + let n_zp = noise3(p + dz); let n_zm = noise3(p - dz); + let inv = 1.0 / (2.0 * eps); + let dn_dx = (n_xp - n_xm) * inv; + let dn_dy = (n_yp - n_ym) * inv; + let dn_dz = (n_zp - n_zm) * inv; + return vec3( + dn_dy.z - dn_dz.y, + dn_dz.x - dn_dx.z, + dn_dx.y - dn_dy.x, + ); +} + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let i = gid.x; + if i >= params.base_count { return; } + + let prev_len = trail_len_buf[i]; + let prev_head = trail_head_buf[i]; + let next_head_if_drift = + select((prev_head + 1u) % prev_len, 0u, prev_len == 0u); + let do_reset = (prev_len == 0u) || (next_head_if_drift == 0u); + + var hp: vec3; + var an: vec3; + var head_slot: u32 = 0u; + var len_i: u32 = prev_len; + + if do_reset { + let cycle_seed = params.seed ^ (i * 2654435761u + prev_head * 7919u + 1u); + let u01 = hash_unit(cycle_seed * 7u + 13u); + let face = cdf_search(u01); + let i0 = source_indices[face * 3u + 0u]; + let i1 = source_indices[face * 3u + 1u]; + let i2 = source_indices[face * 3u + 2u]; + let p0 = vec3(source_position[i0 * 3u + 0u], source_position[i0 * 3u + 1u], source_position[i0 * 3u + 2u]); + let p1 = vec3(source_position[i1 * 3u + 0u], source_position[i1 * 3u + 1u], source_position[i1 * 3u + 2u]); + let p2 = vec3(source_position[i2 * 3u + 0u], source_position[i2 * 3u + 1u], source_position[i2 * 3u + 2u]); + var u = hash_unit(cycle_seed * 31u + 23u); + var v = hash_unit(cycle_seed * 47u + 29u); + if u + v > 1.0 { u = 1.0 - u; v = 1.0 - v; } + hp = (1.0 - u - v) * p0 + u * p1 + v * p2; + an = normalize(cross(p1 - p0, p2 - p0)); + + let r = hash_unit(cycle_seed * 0x9E3779B9u + 17u); + let r_biased = pow(r, 5.0); + let span = f32(params.trail_len_max - params.trail_len_min); + len_i = params.trail_len_min + u32(r_biased * span); + if len_i < params.trail_len_min { len_i = params.trail_len_min; } + if len_i > params.trail_len_max { len_i = params.trail_len_max; } + head_slot = 0u; + + head_pos[i * 3u + 0u] = hp.x; + head_pos[i * 3u + 1u] = hp.y; + head_pos[i * 3u + 2u] = hp.z; + anchor_normal[i * 3u + 0u] = an.x; + anchor_normal[i * 3u + 1u] = an.y; + anchor_normal[i * 3u + 2u] = an.z; + trail_len_buf[i] = len_i; + trail_head_buf[i] = head_slot; + + let base = i * params.trail_len_max; + for (var k = 0u; k < params.trail_len_max; k = k + 1u) { + let s = base + k; + position[s * 3u + 0u] = hp.x; + position[s * 3u + 1u] = hp.y; + position[s * 3u + 2u] = hp.z; + scale[s * 3u + 0u] = 1.0; + scale[s * 3u + 1u] = 1.0; + scale[s * 3u + 2u] = 1.0; + if k < len_i { + life[s] = 1.0; + } else { + life[s] = 0.0; + } + } + } else { + hp = vec3( + head_pos[i * 3u + 0u], + head_pos[i * 3u + 1u], + head_pos[i * 3u + 2u], + ); + an = vec3( + anchor_normal[i * 3u + 0u], + anchor_normal[i * 3u + 1u], + anchor_normal[i * 3u + 2u], + ); + let sample = hp * params.noise_scale + + vec3(params.time, params.time * 0.7, params.time * 1.3); + let raw = curl_noise(sample); + let tangent = raw - dot(raw, an) * an; + hp = hp + tangent * params.noise_strength * params.step_speed; + head_pos[i * 3u + 0u] = hp.x; + head_pos[i * 3u + 1u] = hp.y; + head_pos[i * 3u + 2u] = hp.z; + head_slot = next_head_if_drift; + trail_head_buf[i] = head_slot; + + let trail_slot = i * params.trail_len_max + head_slot; + position[trail_slot * 3u + 0u] = hp.x; + position[trail_slot * 3u + 1u] = hp.y; + position[trail_slot * 3u + 2u] = hp.z; + scale[trail_slot * 3u + 0u] = 1.0; + scale[trail_slot * 3u + 1u] = 1.0; + scale[trail_slot * 3u + 2u] = 1.0; + life[trail_slot] = 1.0; + } +} +"# + .to_string() +} + +const COLOR_FADE: &str = r#" +struct Params { + base_count: u32, + trail_len_max: u32, +} +@group(0) @binding(0) var color: array; +@group(0) @binding(1) var trail_len_buf: array; +@group(0) @binding(2) var trail_head_buf: array; +@group(0) @binding(3) var params: Params; + +@compute @workgroup_size(64) +fn main(@builtin(global_invocation_id) gid: vec3) { + let s = gid.x; + let total = params.base_count * params.trail_len_max; + if s >= total { return; } + let i = s / params.trail_len_max; + let k = s % params.trail_len_max; + + let len_i = trail_len_buf[i]; + if len_i == 0u || k >= len_i { + return; + } + + let head = trail_head_buf[i]; + var slot_age: u32 = 0u; + if k <= head { + slot_age = head - k; + } else { + slot_age = head + len_i - k; + } + let denom = max(len_i - 1u, 1u); + let t = f32(slot_age) / f32(denom); + + let head_c = vec3(0.2, 0.8, 1.0) * 25.0; + let tail_c = vec3(1.0, 0.35, 0.45) * 55.0; + let c = mix(head_c, tail_c, t); + color[s * 4u + 0u] = c.x; + color[s * 4u + 1u] = c.y; + color[s * 4u + 2u] = c.z; + color[s * 4u + 3u] = 1.0; +} +"#; + +fn main() { + sketch().unwrap(); + exit(0).unwrap(); +} + +fn sketch() -> error::Result<()> { + let mut glfw_ctx = GlfwContext::new(1200, 800)?; + init(Config::default())?; + + let surface = glfw_ctx.create_surface(1200, 800)?; + let graphics = graphics_create(surface, 1200, 800, TextureFormat::Rgba16Float)?; + + graphics_mode_3d(graphics)?; + transform_set_position(graphics, Vec3::new(0.0, 0.0, 1700.0))?; + transform_look_at(graphics, Vec3::ZERO)?; + graphics_orbit_camera(graphics)?; + + const TEXT_PT: f32 = 700.0; + const EXTRUSION: f32 = 70.0; + processing_core::app_mut(|app| { + let mut state = app + .world_mut() + .get_mut::(graphics) + .ok_or(error::ProcessingError::GraphicsNotFound)?; + state.style.text_size = TEXT_PT; + Ok(()) + })?; + graphics_record_command(graphics, DrawCommand::TextStyle(TextStyle::Bold))?; + + let text = "processing5"; + let w = graphics_text_width(graphics, text)?; + println!("text width = {w} (size = {TEXT_PT}, extrusion = {EXTRUSION})"); + let mesh = graphics_text_to_model(graphics, text, -w / 2.0, -TEXT_PT * 0.4, EXTRUSION)?; + let source = geometry_create_from_mesh(mesh)?; + + let (cdf_bytes, indices_bytes, face_count) = processing_core::app_mut(|app| { + app.world_mut() + .run_system_cached_with( + processing_render::particles::prepare_scatter_source, + source, + ) + .unwrap() + })?; + let cdf_buf = buffer_create_with_data(cdf_bytes)?; + let idx_buf = buffer_create_with_data(indices_bytes)?; + + let head_pos_buf = buffer_create(BASE_COUNT as u64 * 3 * 4)?; + let anchor_normal_buf = buffer_create(BASE_COUNT as u64 * 3 * 4)?; + let trail_len_buf = buffer_create(BASE_COUNT as u64 * 4)?; + let trail_head_buf = buffer_create(BASE_COUNT as u64 * 4)?; + + let particle = geometry_sphere(SPHERE_RADIUS, 5, 4)?; + + let p = particles_create( + CAPACITY, + vec![ + geometry_attribute_position(), + geometry_attribute_scale(), + geometry_attribute_life(), + geometry_attribute_color(), + ], + )?; + + let sim_src = sim_shader(); + let sim_shader_e = shader_create(&sim_src)?; + let sim = compute_create(sim_shader_e)?; + compute_set( + sim, + "source_position", + shader_value::ShaderValue::MeshAttribute(source, geometry_attribute_position()), + )?; + compute_set(sim, "source_indices", shader_value::ShaderValue::Buffer(idx_buf))?; + compute_set(sim, "cdf", shader_value::ShaderValue::Buffer(cdf_buf))?; + compute_set(sim, "head_pos", shader_value::ShaderValue::Buffer(head_pos_buf))?; + compute_set( + sim, + "anchor_normal", + shader_value::ShaderValue::Buffer(anchor_normal_buf), + )?; + compute_set( + sim, + "trail_len_buf", + shader_value::ShaderValue::Buffer(trail_len_buf), + )?; + compute_set( + sim, + "trail_head_buf", + shader_value::ShaderValue::Buffer(trail_head_buf), + )?; + compute_set(sim, "base_count", shader_value::ShaderValue::UInt(BASE_COUNT))?; + compute_set(sim, "trail_len_max", shader_value::ShaderValue::UInt(TRAIL_LEN_MAX))?; + compute_set(sim, "trail_len_min", shader_value::ShaderValue::UInt(TRAIL_LEN_MIN))?; + compute_set(sim, "face_count", shader_value::ShaderValue::UInt(face_count))?; + compute_set(sim, "seed", shader_value::ShaderValue::UInt(0xc0ffeeu32))?; + compute_set(sim, "noise_scale", shader_value::ShaderValue::Float(0.005))?; + compute_set( + sim, + "noise_strength", + shader_value::ShaderValue::Float(NOISE_STRENGTH), + )?; + compute_set(sim, "step_speed", shader_value::ShaderValue::Float(0.05))?; + + compute_set(sim, "time", shader_value::ShaderValue::Float(0.0))?; + particles_apply(p, sim)?; + + let color_buf = particles_buffer(p, geometry_attribute_color())? + .ok_or(error::ProcessingError::ParticlesNotFound)?; + let mat = material_create_unlit()?; + material_set_albedo_buffer(mat, color_buf)?; + + let fade_shader_e = shader_create(COLOR_FADE)?; + let fade = compute_create(fade_shader_e)?; + compute_set(fade, "base_count", shader_value::ShaderValue::UInt(BASE_COUNT))?; + compute_set( + fade, + "trail_len_max", + shader_value::ShaderValue::UInt(TRAIL_LEN_MAX), + )?; + compute_set( + fade, + "trail_len_buf", + shader_value::ShaderValue::Buffer(trail_len_buf), + )?; + compute_set( + fade, + "trail_head_buf", + shader_value::ShaderValue::Buffer(trail_head_buf), + )?; + + particles_apply(p, fade)?; + + let start = Instant::now(); + let mut render_frame: u32 = 0; + while glfw_ctx.poll_events() { + if render_frame % TRAIL_STRIDE == 0 { + let t = start.elapsed().as_secs_f32(); + compute_set(sim, "time", shader_value::ShaderValue::Float(t * 0.015))?; + particles_apply(p, sim)?; + particles_apply(p, fade)?; + } + render_frame = render_frame.wrapping_add(1); + + graphics_begin_draw(graphics)?; + graphics_record_command( + graphics, + DrawCommand::BackgroundColor(bevy::color::Color::srgb(0.02, 0.02, 0.04)), + )?; + graphics_record_command(graphics, DrawCommand::Material(mat))?; + graphics_record_command( + graphics, + DrawCommand::Particles { + particles: p, + geometry: particle, + }, + )?; + graphics_end_draw(graphics)?; + } + + Ok(()) +} From 985f9590e93d47b51d0479a33a226a96d9e3942c Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?charlotte=20=F0=9F=8C=B8?= Date: Wed, 5 Aug 2026 22:34:16 -0700 Subject: [PATCH 2/2] Fmt. --- crates/processing_ffi/src/lib.rs | 115 +++++++++++------- crates/processing_pyo3/src/lib.rs | 12 +- crates/processing_pyo3/src/particles.rs | 63 +++++----- crates/processing_render/src/graphics.rs | 2 +- crates/processing_render/src/lib.rs | 32 ++--- .../src/particles/kernels/mod.rs | 33 ++--- crates/processing_render/src/particles/mod.rs | 29 ++--- .../src/particles/scatter.rs | 15 +-- examples/particles_scatter.rs | 7 +- examples/particles_text_whirl.rs | 47 +++++-- 10 files changed, 194 insertions(+), 161 deletions(-) diff --git a/crates/processing_ffi/src/lib.rs b/crates/processing_ffi/src/lib.rs index f520db75..722af2f4 100644 --- a/crates/processing_ffi/src/lib.rs +++ b/crates/processing_ffi/src/lib.rs @@ -3243,12 +3243,10 @@ pub extern "C" fn processing_particles_capacity(particles_id: u64) -> u32 { #[unsafe(no_mangle)] pub extern "C" fn processing_particles_buffer(particles_id: u64, attr_id: u64) -> u64 { error::clear_error(); - error::check(|| { - particles_buffer(Entity::from_bits(particles_id), Entity::from_bits(attr_id)) - }) - .flatten() - .map(|e| e.to_bits()) - .unwrap_or(0) + error::check(|| particles_buffer(Entity::from_bits(particles_id), Entity::from_bits(attr_id))) + .flatten() + .map(|e| e.to_bits()) + .unwrap_or(0) } /// # Safety @@ -3284,11 +3282,7 @@ pub unsafe extern "C" fn processing_particles_emit( } #[unsafe(no_mangle)] -pub extern "C" fn processing_particles_emit_gpu( - particles_id: u64, - n: u32, - compute_id: u64, -) { +pub extern "C" fn processing_particles_emit_gpu(particles_id: u64, n: u32, compute_id: u64) { error::clear_error(); error::check(|| { particles_emit_gpu( @@ -3302,38 +3296,45 @@ pub extern "C" fn processing_particles_emit_gpu( #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_noise() -> u64 { error::clear_error(); - error::check(particles_kernel_noise).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_noise) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_transform() -> u64 { error::clear_error(); - error::check(particles_kernel_transform).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_transform) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_attract() -> u64 { error::clear_error(); - error::check(particles_kernel_attract).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_attract) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_drag() -> u64 { error::clear_error(); - error::check(particles_kernel_drag).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_drag) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_vortex() -> u64 { error::clear_error(); - error::check(particles_kernel_vortex).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_vortex) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] -pub extern "C" fn processing_particles_attribute_add( - particles_id: u64, - attribute_id: u64, -) -> i32 { +pub extern "C" fn processing_particles_attribute_add(particles_id: u64, attribute_id: u64) -> i32 { error::clear_error(); error::check(|| { particles_attribute_add( @@ -3349,31 +3350,41 @@ pub extern "C" fn processing_particles_attribute_add( #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_force() -> u64 { error::clear_error(); - error::check(particles_kernel_force).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_force) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_integrate() -> u64 { error::clear_error(); - error::check(particles_kernel_integrate).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_integrate) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_age() -> u64 { error::clear_error(); - error::check(particles_kernel_age).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_age) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_bounds_sphere() -> u64 { error::clear_error(); - error::check(particles_kernel_bounds_sphere).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_bounds_sphere) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_bounds_box() -> u64 { error::clear_error(); - error::check(particles_kernel_bounds_box).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_bounds_box) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] @@ -3387,55 +3398,73 @@ pub extern "C" fn processing_particles_kernel_bounds_geometry(geometry_entity: u #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_impulse() -> u64 { error::clear_error(); - error::check(particles_kernel_impulse).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_impulse) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_flock() -> u64 { error::clear_error(); - error::check(particles_kernel_flock).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_flock) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_orient() -> u64 { error::clear_error(); - error::check(particles_kernel_orient).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_orient) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_field() -> u64 { error::clear_error(); - error::check(particles_kernel_field).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_field) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_attr_linear() -> u64 { error::clear_error(); - error::check(particles_kernel_attr_linear).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_attr_linear) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_attr_combine() -> u64 { error::clear_error(); - error::check(particles_kernel_attr_combine).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_attr_combine) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_attr_mix() -> u64 { error::clear_error(); - error::check(particles_kernel_attr_mix).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_attr_mix) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_attr_lookup1d() -> u64 { error::clear_error(); - error::check(particles_kernel_attr_lookup1d).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_attr_lookup1d) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] pub extern "C" fn processing_particles_kernel_attr_lookup2d() -> u64 { error::clear_error(); - error::check(particles_kernel_attr_lookup2d).map(|e| e.to_bits()).unwrap_or(0) + error::check(particles_kernel_attr_lookup2d) + .map(|e| e.to_bits()) + .unwrap_or(0) } #[unsafe(no_mangle)] @@ -3528,11 +3557,7 @@ pub extern "C" fn processing_particles_apply(particles_id: u64, compute_id: u64) } #[unsafe(no_mangle)] -pub extern "C" fn processing_particles_draw( - graphics_id: u64, - particles_id: u64, - geometry_id: u64, -) { +pub extern "C" fn processing_particles_draw(graphics_id: u64, particles_id: u64, geometry_id: u64) { error::clear_error(); let graphics_entity = Entity::from_bits(graphics_id); error::check(|| { @@ -3599,9 +3624,9 @@ pub unsafe extern "C" fn processing_graphics_world_from_screen( out_z: *mut f32, ) { error::clear_error(); - if let Some(world) = error::check(|| { - graphics_world_from_screen(Entity::from_bits(graphics_id), sx, sy, depth) - }) { + if let Some(world) = + error::check(|| graphics_world_from_screen(Entity::from_bits(graphics_id), sx, sy, depth)) + { unsafe { *out_x = world.x; *out_y = world.y; @@ -3611,11 +3636,7 @@ pub unsafe extern "C" fn processing_graphics_world_from_screen( } #[unsafe(no_mangle)] -pub extern "C" fn processing_graphics_set_bloom( - graphics_id: u64, - intensity: f32, - threshold: f32, -) { +pub extern "C" fn processing_graphics_set_bloom(graphics_id: u64, intensity: f32, threshold: f32) { error::clear_error(); error::check(|| graphics_set_bloom(Entity::from_bits(graphics_id), intensity, threshold)); } diff --git a/crates/processing_pyo3/src/lib.rs b/crates/processing_pyo3/src/lib.rs index ba08e6f2..d5dd6c31 100644 --- a/crates/processing_pyo3/src/lib.rs +++ b/crates/processing_pyo3/src/lib.rs @@ -334,12 +334,6 @@ mod mewnala { #[pymodule_export] use super::Compute; #[pymodule_export] - use super::particles::Attribute; - #[pymodule_export] - use super::particles::AttributeFormat; - #[pymodule_export] - use super::particles::Particles; - #[pymodule_export] use super::Font; #[pymodule_export] use super::Geometry; @@ -375,6 +369,12 @@ mod mewnala { #[pymodule_export] use super::monitor::Monitor; #[pymodule_export] + use super::particles::Attribute; + #[pymodule_export] + use super::particles::AttributeFormat; + #[pymodule_export] + use super::particles::Particles; + #[pymodule_export] use super::surface::Surface; #[pymodule_init] diff --git a/crates/processing_pyo3/src/particles.rs b/crates/processing_pyo3/src/particles.rs index afb7e699..e2a6c2f4 100644 --- a/crates/processing_pyo3/src/particles.rs +++ b/crates/processing_pyo3/src/particles.rs @@ -262,57 +262,56 @@ impl Particles { #[staticmethod] pub fn noise() -> PyResult { - let entity = particles_kernel_noise() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_noise().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn transform() -> PyResult { - let entity = particles_kernel_transform() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_transform().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn attract() -> PyResult { - let entity = particles_kernel_attract() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_attract().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn drag() -> PyResult { - let entity = particles_kernel_drag() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_drag().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn vortex() -> PyResult { - let entity = particles_kernel_vortex() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_vortex().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn force() -> PyResult { - let entity = particles_kernel_force() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_force().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn integrate() -> PyResult { - let entity = particles_kernel_integrate() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_integrate().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn age() -> PyResult { - let entity = particles_kernel_age() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = particles_kernel_age().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } @@ -325,8 +324,8 @@ impl Particles { #[staticmethod] pub fn bounds_box() -> PyResult { - let entity = particles_kernel_bounds_box() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_bounds_box().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } @@ -339,50 +338,50 @@ impl Particles { #[staticmethod] pub fn impulse() -> PyResult { - let entity = particles_kernel_impulse() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_impulse().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn flock() -> PyResult { - let entity = particles_kernel_flock() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_flock().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn orient() -> PyResult { - let entity = particles_kernel_orient() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_orient().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn field() -> PyResult { - let entity = particles_kernel_field() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_field().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn attr_linear() -> PyResult { - let entity = particles_kernel_attr_linear() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_attr_linear().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn attr_combine() -> PyResult { - let entity = particles_kernel_attr_combine() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_attr_combine().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } #[staticmethod] pub fn attr_mix() -> PyResult { - let entity = particles_kernel_attr_mix() - .map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; + let entity = + particles_kernel_attr_mix().map_err(|e| PyRuntimeError::new_err(format!("{e}")))?; Ok(Compute::from_entity(entity)) } diff --git a/crates/processing_render/src/graphics.rs b/crates/processing_render/src/graphics.rs index 9cb39ee0..4fc9071b 100644 --- a/crates/processing_render/src/graphics.rs +++ b/crates/processing_render/src/graphics.rs @@ -9,9 +9,9 @@ use bevy::{ ImageRenderTarget, MsaaWriteback, Projection, RenderTarget, visibility::RenderLayers, }, core_pipeline::tonemapping::Tonemapping, - post_process::bloom::Bloom, ecs::query::QueryEntityError, math::{Mat4, Vec3A}, + post_process::bloom::Bloom, prelude::*, render::{ RenderApp, diff --git a/crates/processing_render/src/lib.rs b/crates/processing_render/src/lib.rs index b7ef15ce..6f4cc72e 100644 --- a/crates/processing_render/src/lib.rs +++ b/crates/processing_render/src/lib.rs @@ -26,14 +26,14 @@ pub use particles::{ FALLOFF_INVERSE, FALLOFF_LINEAR, FALLOFF_QUADRATIC, FALLOFF_SMOOTHSTEP, particles_apply, particles_attribute_add, particles_buffer, particles_capacity, particles_create, particles_create_from_geometry, particles_destroy, particles_emit, particles_emit_gpu, - particles_kernel_age, - particles_kernel_attr_combine, particles_kernel_attr_linear, particles_kernel_attr_lookup1d, - particles_kernel_attr_lookup2d, particles_kernel_attr_mix, particles_kernel_attract, - particles_kernel_bounds_box, particles_kernel_bounds_geometry, particles_kernel_bounds_sphere, - particles_kernel_drag, particles_kernel_field, particles_kernel_flock, particles_kernel_force, - particles_kernel_impulse, particles_kernel_integrate, particles_kernel_noise, - particles_kernel_orient, particles_kernel_transform, particles_kernel_vortex, - particles_scatter_create, particles_scatter_volume_create, + particles_kernel_age, particles_kernel_attr_combine, particles_kernel_attr_linear, + particles_kernel_attr_lookup1d, particles_kernel_attr_lookup2d, particles_kernel_attr_mix, + particles_kernel_attract, particles_kernel_bounds_box, particles_kernel_bounds_geometry, + particles_kernel_bounds_sphere, particles_kernel_drag, particles_kernel_field, + particles_kernel_flock, particles_kernel_force, particles_kernel_impulse, + particles_kernel_integrate, particles_kernel_noise, particles_kernel_orient, + particles_kernel_transform, particles_kernel_vortex, particles_scatter_create, + particles_scatter_volume_create, }; use std::path::PathBuf; @@ -875,7 +875,10 @@ pub fn graphics_world_from_screen( ) -> error::Result { app_mut(|app| { app.world_mut() - .run_system_cached_with(graphics::world_from_screen, (graphics_entity, sx, sy, depth)) + .run_system_cached_with( + graphics::world_from_screen, + (graphics_entity, sx, sy, depth), + ) .unwrap() }) } @@ -887,10 +890,7 @@ pub fn graphics_set_bloom( ) -> error::Result<()> { app_mut(|app| { app.world_mut() - .run_system_cached_with( - graphics::set_bloom, - (graphics_entity, intensity, threshold), - ) + .run_system_cached_with(graphics::set_bloom, (graphics_entity, intensity, threshold)) .unwrap() }) } @@ -1874,7 +1874,11 @@ pub fn material_set_emissive_buffer( entity: Entity, emissive_buffer_entity: Entity, ) -> error::Result<()> { - material_set_particles_buffer(entity, emissive_buffer_entity, ParticlesBufferSlot::Emissive) + material_set_particles_buffer( + entity, + emissive_buffer_entity, + ParticlesBufferSlot::Emissive, + ) } pub fn material_set( diff --git a/crates/processing_render/src/particles/kernels/mod.rs b/crates/processing_render/src/particles/kernels/mod.rs index c0068373..357461e6 100644 --- a/crates/processing_render/src/particles/kernels/mod.rs +++ b/crates/processing_render/src/particles/kernels/mod.rs @@ -128,8 +128,7 @@ pub fn particles_kernel_force() -> error::Result { } pub fn particles_kernel_integrate() -> error::Result { - let shader = - shader_load("embedded://processing_render/particles/kernels/integrate.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/integrate.wgsl")?; let entity = compute_create(shader)?; set_requires(entity, &["position", "velocity"])?; compute_set(entity, "dt", ShaderValue::Float(1.0))?; @@ -157,8 +156,7 @@ pub fn particles_kernel_vortex() -> error::Result { } pub fn particles_kernel_bounds_sphere() -> error::Result { - let shader = - shader_load("embedded://processing_render/particles/kernels/bounds_sphere.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/bounds_sphere.wgsl")?; let entity = compute_create(shader)?; set_requires(entity, &["position", "velocity"])?; compute_set(entity, "center", ShaderValue::Float3([0.0; 3]))?; @@ -170,8 +168,7 @@ pub fn particles_kernel_bounds_sphere() -> error::Result { } pub fn particles_kernel_bounds_box() -> error::Result { - let shader = - shader_load("embedded://processing_render/particles/kernels/bounds_box.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/bounds_box.wgsl")?; let entity = compute_create(shader)?; set_requires(entity, &["position", "velocity"])?; compute_set(entity, "aabb_min", ShaderValue::Float3([-1.0, -1.0, -1.0]))?; @@ -236,7 +233,11 @@ pub fn particles_kernel_impulse() -> error::Result { compute_set(entity, "radius", ShaderValue::Float(1.0))?; compute_set(entity, "position_kick", ShaderValue::Float(0.0))?; compute_set(entity, "velocity_kick", ShaderValue::Float(0.0))?; - compute_set(entity, "falloff_mode", ShaderValue::UInt(FALLOFF_SMOOTHSTEP))?; + compute_set( + entity, + "falloff_mode", + ShaderValue::UInt(FALLOFF_SMOOTHSTEP), + )?; Ok(entity) } @@ -270,13 +271,16 @@ pub fn particles_kernel_field() -> error::Result { set_requires(entity, &["position"])?; compute_set(entity, "center", ShaderValue::Float3([0.0; 3]))?; compute_set(entity, "radius", ShaderValue::Float(1.0))?; - compute_set(entity, "falloff_mode", ShaderValue::UInt(FALLOFF_SMOOTHSTEP))?; + compute_set( + entity, + "falloff_mode", + ShaderValue::UInt(FALLOFF_SMOOTHSTEP), + )?; Ok(entity) } pub fn particles_kernel_attr_linear() -> error::Result { - let shader = - shader_load("embedded://processing_render/particles/kernels/attr_linear.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/attr_linear.wgsl")?; let entity = compute_create(shader)?; compute_set(entity, "scale", ShaderValue::Float(1.0))?; compute_set(entity, "offset", ShaderValue::Float(0.0))?; @@ -284,8 +288,7 @@ pub fn particles_kernel_attr_linear() -> error::Result { } pub fn particles_kernel_attr_combine() -> error::Result { - let shader = - shader_load("embedded://processing_render/particles/kernels/attr_combine.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/attr_combine.wgsl")?; let entity = compute_create(shader)?; compute_set(entity, "op", ShaderValue::UInt(COMBINE_ADD))?; compute_set(entity, "b_scale", ShaderValue::Float(1.0))?; @@ -303,8 +306,7 @@ pub fn particles_kernel_attr_mix() -> error::Result { } pub fn particles_kernel_attr_lookup1d() -> error::Result { - let shader = - shader_load("embedded://processing_render/particles/kernels/attr_lookup1d.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/attr_lookup1d.wgsl")?; let entity = compute_create(shader)?; compute_set(entity, "scale", ShaderValue::Float(1.0))?; compute_set(entity, "offset", ShaderValue::Float(0.0))?; @@ -312,8 +314,7 @@ pub fn particles_kernel_attr_lookup1d() -> error::Result { } pub fn particles_kernel_attr_lookup2d() -> error::Result { - let shader = - shader_load("embedded://processing_render/particles/kernels/attr_lookup2d.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/attr_lookup2d.wgsl")?; let entity = compute_create(shader)?; compute_set(entity, "u_scale", ShaderValue::Float(1.0))?; compute_set(entity, "u_offset", ShaderValue::Float(0.0))?; diff --git a/crates/processing_render/src/particles/mod.rs b/crates/processing_render/src/particles/mod.rs index 28efe687..577f2261 100644 --- a/crates/processing_render/src/particles/mod.rs +++ b/crates/processing_render/src/particles/mod.rs @@ -10,14 +10,13 @@ pub use emit::{particles_apply, particles_emit, particles_emit_gpu}; pub use kernels::{ BOUNDS_CLAMP, BOUNDS_REFLECT, BOUNDS_SOFT, BOUNDS_WRAP, COMBINE_ADD, COMBINE_DIV, COMBINE_MAX, COMBINE_MIN, COMBINE_MUL, COMBINE_POW, COMBINE_SUB, FALLOFF_CONST, FALLOFF_CUBIC, - FALLOFF_INVERSE, FALLOFF_LINEAR, FALLOFF_QUADRATIC, FALLOFF_SMOOTHSTEP, - particles_kernel_age, particles_kernel_attr_combine, particles_kernel_attr_linear, - particles_kernel_attr_lookup1d, particles_kernel_attr_lookup2d, particles_kernel_attr_mix, - particles_kernel_attract, particles_kernel_bounds_box, particles_kernel_bounds_geometry, - particles_kernel_bounds_sphere, particles_kernel_drag, particles_kernel_field, - particles_kernel_flock, particles_kernel_force, particles_kernel_impulse, - particles_kernel_integrate, particles_kernel_noise, particles_kernel_orient, - particles_kernel_transform, particles_kernel_vortex, + FALLOFF_INVERSE, FALLOFF_LINEAR, FALLOFF_QUADRATIC, FALLOFF_SMOOTHSTEP, particles_kernel_age, + particles_kernel_attr_combine, particles_kernel_attr_linear, particles_kernel_attr_lookup1d, + particles_kernel_attr_lookup2d, particles_kernel_attr_mix, particles_kernel_attract, + particles_kernel_bounds_box, particles_kernel_bounds_geometry, particles_kernel_bounds_sphere, + particles_kernel_drag, particles_kernel_field, particles_kernel_flock, particles_kernel_force, + particles_kernel_impulse, particles_kernel_integrate, particles_kernel_noise, + particles_kernel_orient, particles_kernel_transform, particles_kernel_vortex, }; pub use scatter::{ particles_scatter_create, particles_scatter_volume_create, prepare_scatter_source, @@ -278,7 +277,9 @@ pub fn materialize_attribute( None => { let mut hit = None; for (&e, &buf) in &particles.buffers { - let Ok(other) = attributes.get(e) else { continue }; + let Ok(other) = attributes.get(e) else { + continue; + }; if other.name == attr.name { if other.format != attr.format { return Err(ProcessingError::InvalidArgument(format!( @@ -341,10 +342,7 @@ pub fn materialize_attribute( Ok(buffer_entity) } -pub fn particles_create( - capacity: u32, - attribute_entities: Vec, -) -> error::Result { +pub fn particles_create(capacity: u32, attribute_entities: Vec) -> error::Result { app_mut(|app| { app.world_mut() .run_system_cached_with(create, (capacity, attribute_entities)) @@ -381,10 +379,7 @@ pub fn particles_capacity(entity: Entity) -> error::Result { }) } -pub fn particles_buffer( - entity: Entity, - attribute_entity: Entity, -) -> error::Result> { +pub fn particles_buffer(entity: Entity, attribute_entity: Entity) -> error::Result> { app_mut(|app| { Ok(app .world() diff --git a/crates/processing_render/src/particles/scatter.rs b/crates/processing_render/src/particles/scatter.rs index fef13f9c..f42c5b64 100644 --- a/crates/processing_render/src/particles/scatter.rs +++ b/crates/processing_render/src/particles/scatter.rs @@ -30,11 +30,7 @@ pub fn particles_scatter_create(source_geometry: Entity) -> error::Result error::Result let idx_buf = buffer_create_with_data(indices_bytes)?; - let shader = - shader_load("embedded://processing_render/particles/kernels/scatter_volume.wgsl")?; + let shader = shader_load("embedded://processing_render/particles/kernels/scatter_volume.wgsl")?; let scatter = compute_create(shader)?; let position_attr = geometry_attribute_position(); @@ -61,11 +56,7 @@ pub fn particles_scatter_volume_create(source_geometry: Entity) -> error::Result "source_position", ShaderValue::MeshAttribute(source_geometry, position_attr), )?; - compute_set( - scatter, - "source_indices", - ShaderValue::Buffer(idx_buf), - )?; + compute_set(scatter, "source_indices", ShaderValue::Buffer(idx_buf))?; compute_set( scatter, "aabb_min", diff --git a/examples/particles_scatter.rs b/examples/particles_scatter.rs index 69a41173..dd046618 100644 --- a/examples/particles_scatter.rs +++ b/examples/particles_scatter.rs @@ -50,11 +50,8 @@ fn sketch() -> error::Result<()> { transform_set_position(graphics, Vec3::new(0.0, 0.4, 4.5))?; transform_look_at(graphics, Vec3::ZERO)?; - let _key = light_create_directional( - graphics, - bevy::color::Color::srgb(1.0, 0.95, 0.85), - 4500.0, - )?; + let _key = + light_create_directional(graphics, bevy::color::Color::srgb(1.0, 0.95, 0.85), 4500.0)?; let source = geometry_sphere(1.2, 96, 48)?; let scatter = particles_scatter_create(source)?; diff --git a/examples/particles_text_whirl.rs b/examples/particles_text_whirl.rs index 1d98b802..b9b2a778 100644 --- a/examples/particles_text_whirl.rs +++ b/examples/particles_text_whirl.rs @@ -292,10 +292,7 @@ fn sketch() -> error::Result<()> { let (cdf_bytes, indices_bytes, face_count) = processing_core::app_mut(|app| { app.world_mut() - .run_system_cached_with( - processing_render::particles::prepare_scatter_source, - source, - ) + .run_system_cached_with(processing_render::particles::prepare_scatter_source, source) .unwrap() })?; let cdf_buf = buffer_create_with_data(cdf_bytes)?; @@ -326,9 +323,17 @@ fn sketch() -> error::Result<()> { "source_position", shader_value::ShaderValue::MeshAttribute(source, geometry_attribute_position()), )?; - compute_set(sim, "source_indices", shader_value::ShaderValue::Buffer(idx_buf))?; + compute_set( + sim, + "source_indices", + shader_value::ShaderValue::Buffer(idx_buf), + )?; compute_set(sim, "cdf", shader_value::ShaderValue::Buffer(cdf_buf))?; - compute_set(sim, "head_pos", shader_value::ShaderValue::Buffer(head_pos_buf))?; + compute_set( + sim, + "head_pos", + shader_value::ShaderValue::Buffer(head_pos_buf), + )?; compute_set( sim, "anchor_normal", @@ -344,10 +349,26 @@ fn sketch() -> error::Result<()> { "trail_head_buf", shader_value::ShaderValue::Buffer(trail_head_buf), )?; - compute_set(sim, "base_count", shader_value::ShaderValue::UInt(BASE_COUNT))?; - compute_set(sim, "trail_len_max", shader_value::ShaderValue::UInt(TRAIL_LEN_MAX))?; - compute_set(sim, "trail_len_min", shader_value::ShaderValue::UInt(TRAIL_LEN_MIN))?; - compute_set(sim, "face_count", shader_value::ShaderValue::UInt(face_count))?; + compute_set( + sim, + "base_count", + shader_value::ShaderValue::UInt(BASE_COUNT), + )?; + compute_set( + sim, + "trail_len_max", + shader_value::ShaderValue::UInt(TRAIL_LEN_MAX), + )?; + compute_set( + sim, + "trail_len_min", + shader_value::ShaderValue::UInt(TRAIL_LEN_MIN), + )?; + compute_set( + sim, + "face_count", + shader_value::ShaderValue::UInt(face_count), + )?; compute_set(sim, "seed", shader_value::ShaderValue::UInt(0xc0ffeeu32))?; compute_set(sim, "noise_scale", shader_value::ShaderValue::Float(0.005))?; compute_set( @@ -367,7 +388,11 @@ fn sketch() -> error::Result<()> { let fade_shader_e = shader_create(COLOR_FADE)?; let fade = compute_create(fade_shader_e)?; - compute_set(fade, "base_count", shader_value::ShaderValue::UInt(BASE_COUNT))?; + compute_set( + fade, + "base_count", + shader_value::ShaderValue::UInt(BASE_COUNT), + )?; compute_set( fade, "trail_len_max",